CUDA编程：Hello world

第三章 CUDA C简介

1. 输出hello world

   
   #include<stdio.h>
   
   __global__ void kernel() {
     printf("hello world");
   }
   
   int main() {
     kernel<<<1, 1>>>();
     return 0;
   }
   

   这个程序和普通的C程序的区别值得注意

   • 函数的定义带有了__global__这个标签，表示这个函数是在GPU上运行
   • 函数的调用除了常规的参数之外，还增加了<<<>>>修饰。而其中的数字将传递个CUDA的运行时系统，至于能干啥，下一章会讲。

2. 进阶版

   
   #include<stdio.h>
   
   __global__ void add(int a,int b,int *c){
     *c = a + b;
   }
   int main(){
     int c;
     int *dev_c;
     cudaMalloc((void**)&dev_c,sizeof(int));
     add<<<1,1>>>(2,7,dev_c);
     cudaMemcpy(&c,dev_c,sizeof(int),cudaMemcpyDeviceToHost);
     printf("2 + 7 = %d",c);
     return 0;
   }

   这里就涉及了GPU和主机之间的内存交换了，cudaMalloc是在GPU的内存里开辟一片空间，然后通过操作之后，这个内存里有了计算出来内容，再通过cudaMemcpy这个函数把内容从GPU复制出来。就是这么简单。

第四章 CUDA C并行编程

这一章开始体现CUDA并行编程的魅力。 
以下是一个数组求和的代码

#include<stdio.h>

#define N   10

__global__ void add( int *a, int *b, int *c ) {
    int tid = blockIdx.x;    // this thread handles the data at its thread id
    if (tid < N)
        c[tid] = a[tid] + b[tid];
}

int main( void ) {
    int a[N], b[N], c[N];
    int *dev_a, *dev_b, *dev_c;

    // allocate the memory on the GPU
    cudaMalloc( (void**)&dev_a, N * sizeof(int) );
    cudaMalloc( (void**)&dev_b, N * sizeof(int) );
    cudaMalloc( (void**)&dev_c, N * sizeof(int) );

    // fill the arrays 'a' and 'b' on the CPU
    for (int i=0; i<N; i++) {
        a[i] = -i;
        b[i] = i * i;
    }

    // copy the arrays 'a' and 'b' to the GPU
    cudaMemcpy( dev_a, a, N * sizeof(int),
                              cudaMemcpyHostToDevice );
    cudaMemcpy( dev_b, b, N * sizeof(int),
                              cudaMemcpyHostToDevice );

    add<<<N,1>>>( dev_a, dev_b, dev_c );

    // copy the array 'c' back from the GPU to the CPU
    cudaMemcpy( c, dev_c, N * sizeof(int),
                              cudaMemcpyDeviceToHost );

    // display the results
    for (int i=0; i<N; i++) {
        printf( "%d + %d = %d\n", a[i], b[i], c[i] );
    }

    // free the memory allocated on the GPU
    cudaFree( dev_a );
    cudaFree( dev_b );
    cudaFree( dev_c );
    return 0;
}

重点也是对于初学者最难理解的就是kernel函数了：

 __global__ void add( int *a, int *b, int *c ) {
    int tid = blockIdx.x;
    if (tid < N)
        c[tid] = a[tid] + b[tid];
}

GPU编程和CPU编程的最大区别也就在这里体现出来了，就是数组求和竟然不要循环！为什么不要循环，就是因为这里的tid可以把整个循环的工作做了。这里的tid也就是thread的id，每个thread负责数组一个数的操作，所以将10个循环操作拆分成了十个线程同时搞定。这里的kernel函数也就是可以同时并发执行，而里面的tid的数值是不一样的。