cuda c 第三章 第四章 第五章

  • 定义
    • CPU以及系统的内存称为主机。
    • GPU及其内存称为设备
    • 核函数为在设备上运行的函数，使用__global__修饰
  • 内存分配与释放
    1
    2
    3
    4
    int c;
    int * dev_c;
    cudaMalloc((void**)&dev_c,sizeof(int));
    cudaFree(dev_c);
  • 内存拷贝
    1
    2
    3
    4
    5
    6
    int c;
    int * dev_c;
    cudaMemcpy(&c,dev_c,sizeof(int),cudaMemcpyDeviceToHost);
    \\cudaMemcpyDeviceToHost
    \\cudaMemcpyHostToDevice
    \\cudaMemcpyDeviceToDevice
  • 获取设备数量和属性
    1
    2
    3
    4
    int count;
    cudaDeviceProp prop;
    cudaGetDeviceCount(&count);
    cudaGetDeviceProperties(&prop,0);

prop

  • 并行编程
    1
    2
    3
    4
    5
    __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];
    }

注意 __global__和__device__函数的区别是：前者说明是由主机调用的在GPU上执行的函数，而后者只能由在GPU上的函数调用，即只能由其他__global__和__device__的函数调用。

  • 函数调用
    1
    kernel<<<m,n>>>()

其中，m指使用的并行线程块的数量，这个线程块集合成为一个gird。blockIdx.x当前位于哪一个线程块。cuda支持二维并行线程块数组。当使用二维的grid时，使用dim3类型设置。如：

1
2
dim3 grid(DIM,DIM);
kernel<<<grid,1>>>();

此时得到当前block的id涉及

1
2
3
int x = blockIdx.x;
int y = blockIdx.y;
int offset = x + y*gridDim.x;

n表示在每个block创建按的线程数目。

  • 单线程块内多线程
    1
    kernel<<<1,n>>>

使用线程号索引

1
int tid = blockIdx.x;

索引可能用到的

1
2
3
4
blockIdx.x,blockIdx.y : grid内的block索引
gridDim.x,gridDim.y : grid的大小，即内部的block排列维度
threadIdx.x : block内的线程数
blockDim.x : 线程数

其中gridDim是二维的，blockDim是三维的，也就是说允许在2维的grid中每个block中有多个线程。此时的索引方式：

1
int offset = x + y * DIM;

此时DIM表示线程块中线程的数量，y为线程块索引，x为线程块内的线程索引。
例子：

1
2
3
4
5
6
7
8
9
10
//任意长向量相加
__global__ void add( int *a, int *b, int *c ) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    while (tid < N) {
        c[tid] = a[tid] + b[tid];
        tid += blockDim.x * gridDim.x;
    }
}
//调用
add<<<128,128>>>( dev_a, dev_b, dev_c );

说明：总共N=(33 * 1024)，总线程数128*128，因此，如果add中没有循，则只处理前128×128个元素。blockDim为单block内的总共线程数目，gridDim为总共block数目。是指为处理下一个128*128的元素。

  • 共享内存
    关键字：
    1
    __share__

线程同步使用函数

1
__synchreads();

线程块中所有线程都执行完__synctheads()前面的代码后才会执行下一条语句。例如：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
__global__ void dot( float *a, float *b, float *c ) 
{
    __shared__ float cache[threadsPerBlock];
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int cacheIndex = threadIdx.x;

    float   temp = 0;
    while (tid < N) {
        temp += a[tid] * b[tid];
        tid += blockDim.x * gridDim.x;
    }

    // set the cache values
    cache[cacheIndex] = temp;

    // synchronize threads in this block
    __syncthreads();

    // for reductions, threadsPerBlock must be a power of 2
    // because of the following code
    int i = blockDim.x/2;
    while (i != 0) 
    {
        if (cacheIndex < i)
        {
            cache[cacheIndex] += cache[cacheIndex + i];
        }
        __syncthreads();
        i /= 2;
    }

    if (cacheIndex == 0)
        c[blockIdx.x] = cache[0];
}

以上为求向量点积代码，后段求和使用两两向前求和的方式，称之为归约。另需特别注意同步函数的位置。