CUDA

基础知识

__global__用于定义核函数，没有返回值，定义的时候需要<<<block数量，每个block的线程数量>>>

__device__用定义设备函数，在GPU上执行，也只能在GPU中被调用，global函数可以调用device.

__host__device__定义的函数可以在GPU和CPU中通用

__global__ void kernel()
{
    printf("Thread %d of %d \n",threadIdx.x,blockDim.x); 
    //一个grid有多少 线程编号和线程数量


    printf("Block %d of %d,Thread %d of %d\n",
        blockIdx.x,gridDim.x,threadIdx.x,blockDim.x)
        /* block的编号 ：block 1 of 2, block 0 of 2
        thread的编号：thread 0 of 3, thread 1 of 3,thread 2 of 3
        */ 
}

__global__ void kernel()
{
    int tid = blockDim.x * blockIdx.x + treadIdx.x;  //线程标记
    int tnum  = blockDim.x * gridDim.x; //总线程数
}

int main()
{
    kernel<<<2,3>>>();  //几个block，每个block有几个线程
    cudaDeviceSynchronize();
    /* 等待GPU处理结束，不然会直接return */
    return 0;
}

内存管理

cuda是利用的是异步机制，简答来说，就是GPU和CPU两个设备各算各的，而且两个设备之间的内存是相互隔离的，但是需要将GPU的计算结果返回到CPU中，这是就需要对内存进行管理，比如可以先用cudaMalloc分配好GPU中的内存，在GPU中计算完成后，再利用cudaMemcpy将结果拷回到CPU内存中。当然，还有一种办法就是利用cudaMallocManaged的统一寻址功能，让内存在GPU和CPU中进行统一。

int main()
{
    int *pret = (int *)malloc(sizeof(int));
    kernel<<<1,1>>>(pret);
    checkCudaErrors(cudaDeviceSynchronize());//同步错误

/*先申请一块GPU内存，再将计算后的拷贝到CPU中，最后在main中用*/
    int *npret;
    checkCudaErrors(cudaMalloc(&npret,sizeof(int))); 
    kernel<<<1,1>>>(npret);
    int ret;
    cudaMemcpy(&ret,npret,sizeof(int),cudaMemcpyDeviceToHost);
    printf(ret);

/*利用统一寻址的功能在GPU和CPU中使用内存*/
    int *pret;
    cudaMallocManaged(&pret,sizeof(int));
    kernel<<<1,1>>>(pret);
    cudaDeviceSynchronize();
    printf("result:%d\n",*pret);
    return 0;
    
}

Thrust库

Thrust库提供一些算法啊和容器，和C++中的STL相似，比如vector,Thrust提供了统一寻址的universal_vector，这样可以在GPU和CPU中直接进行修改，也可以用host_vector创建CPU的vector，用device_vector创建GPU的vector，比较好的就是，host给device进行赋值操作，会自动调用cudaMemcpy。

int main()
{
    int n = 31312;
    thrust::universal_vector<float> x(n);
    thrust::host_vector<int> x_host(5);
    thrust::device_vector<int> x_device = x_host;
}

原子操作

由于线程的并行操作，会导致对某一资源的修改发生冲突，利用原子操作进行控制，但是，如果并行线程都是修改同一个共享资源，那么原子操作可能会让其变成顺序执行。

__global__ void paralel_sum(int *sum,int const *arr,int n)
{
    for(int i =blockDim.x *blockIdx.x+threadIdx.x；
        i<n；i+=blockDim.x * gridDim.x)
    {
        /*sum[0]+= arr[i]; //错误，GPU并行修改sum[0]的寄存器值*/
        atomicAdd(&sum[0],arr[i]); 
        /* 其他原子操作
         -atomicMax(dst, src)
         -atomicMin(dst, src)
         -atomicSub(dst, src)
         -atomicAdd(dst, src)
        */
    }
}