CUDA编程(二) —— CUDA编程模型

CUDA编程模型

常见的CUDA术语

      我们怎么写一个能在GPU跑的程序或函数呢？

      通过关键字就可以表示某个程序在CPU上跑还是在GPU上跑！如下表所示，比如我们用__global__定义一个kernel函数，就是CPU上调用，GPU上执行，注意__global__函数的返回值必须设置为void。

CPU和GPU间的数据传输怎么写？

GPU内存分配回收内存的函数接口：

• cudaMalloc(): 在设备端分配global memory
• cudaFree(): 释放存储空间

CPU的数据和GPU端数据做数据传输的函数接口是一样的，他们通过传递的函数实参（枚举类型）来表示传输方向：

cudaMemcpy(void *dst, void *src, size_t nbytes, enum cudaMemcpyKind direction)

enum cudaMemcpyKind:

• cudaMemcpyHostToDevice（CPU到GPU）
• cudaMemcpyDeviceToHost（GPU到CPU）
• cudaMemcpyDeviceToDevice（GPU到GPU）

怎么用代码表示线程组织模型？

我们可以用dim3类来表示网格和线程块的组织方式，网格grid可以表示为一维和二维格式，线程块block可以表示为一维、二维和三维的数据格式。

dim3 DimGrid(100, 50);  //5000个线程块，维度是100*50
dim3 DimBlock(4, 8, 8);  //每个线层块内包含256个线程，线程块内的维度是4*8*8

怎样计算线程号

①如果使用N个线程块，每一个线程块只有一个线程

即

dim3 dimGrid(N);
dim3 dimBlock(1);

此时的线程号的计算方式就是

threadId = blockIdx.x;

其中threadId的取值范围为0到N-1。对于这种情况，我们可以将其看作是一个列向量，列向量中的每一行对应一个线程块。列向量中每一行只有1个元素，对应一个线程。

②使用M×N个线程块，每个线程块1个线程

由于线程块是2维的，故可以看做是一个M*N的2维矩阵，其线程号有两个维度，即：

dim3 dimGrid(M,N);
dim3 dimBlock(1);

其中

blockIdx.x 取值0到M-1

blcokIdx.y 取值0到N-1

这种情况一般用于处理2维数据结构，比如2维图像。每一个像素用一个线程来处理，此时需要线程号来映射图像像素的对应位置，如

pos = blockIdx.y * blcokDim.x + blockIdx.x; //其中gridDim.x等于M

③使用一个线程块，该线程具有N个线程

即

dim3 dimGrid(1);
dim3 dimBlock(N);

此时线程号的计算方式为

threadId = threadIdx.x;

其中threadId的范围是0到N-1，对于这种情况，可以看做是一个行向量，行向量中的每一个元素的每一个元素对应着一个线程。

④使用M个线程块，每个线程块内含有N个线程

即

dim3 dimGrid(M);
dim3 dimBlock(N);

这种情况，可以把它想象成二维矩阵，矩阵的行与线程块对应，矩阵的列与线程编号对应，那线程号的计算方式为

threadId = threadIdx.x + blcokIdx*blockDim.x;

上面其实就是把二维的索引空间转换为一维索引空间的过程。

⑤使用M×N的二维线程块，每一个线程块具有P×Q个线程

即

dim3 dimGrid(M, N);
dim3 dimBlock(P, Q);

这种情况其实是我们遇到的最多情况，特别适用于处理具有二维数据结构的算法，比如图像处理领域。

其索引有两个维度

threadId.x = blockIdx.x*blockDim.x+threadIdx.x;
threadId.y = blockIdx.y*blockDim.y+threadIdx.y;

上述公式就是把线程和线程块的索引映射为图像像素坐标的计算方法。

CUDA代码示例

test.cu

#include "device_launch_parameters.h"
#include <iostream>

int main()
{
    int deviceCount;
    cudaGetDeviceCount(&deviceCount);
    for(int i=0;i<deviceCount;i++)
    {
        cudaDeviceProp devProp;
        cudaGetDeviceProperties(&devProp, i);
        std::cout << "使用GPU device " << i << ": " << devProp.name << std::endl;
        std::cout << "设备全局内存总量： " << devProp.totalGlobalMem / 1024 / 1024 << "MB" << std::endl;
        std::cout << "SM的数量：" << devProp.multiProcessorCount << std::endl;
        std::cout << "每个线程块的共享内存大小：" << devProp.sharedMemPerBlock / 1024.0 << " KB" << std::endl;
        std::cout << "每个线程块的最大线程数：" << devProp.maxThreadsPerBlock << std::endl;
        std::cout << "设备上一个线程块（Block）种可用的32位寄存器数量： " << devProp.regsPerBlock << std::endl;
        std::cout << "每个EM的最大线程数：" << devProp.maxThreadsPerMultiProcessor << std::endl;
        std::cout << "每个EM的最大线程束数：" << devProp.maxThreadsPerMultiProcessor / 32 << std::endl;
        std::cout << "设备上多处理器的数量： " << devProp.multiProcessorCount << std::endl;
        std::cout << "======================================================" << std::endl;     
        
    }
    return 0;
}

所以.cu其实是C++代码

编译

nvcc test.cu -o test

输出结果：

使用GPU device 0: GeForce RTX 2080 Ti
设备全局内存总量： 11019MB
SM的数量：68
每个线程块的共享内存大小：48 KB
每个线程块的最大线程数：1024
设备上一个线程块（Block）种可用的32位寄存器数量： 65536
每个EM的最大线程数：1024
每个EM的最大线程束数：32
设备上多处理器的数量： 68
======================================================
使用GPU device 1: GeForce RTX 2080 Ti
设备全局内存总量： 11019MB
SM的数量：68
每个线程块的共享内存大小：48 KB
每个线程块的最大线程数：1024
设备上一个线程块（Block）种可用的32位寄存器数量： 65536
每个EM的最大线程数：1024
每个EM的最大线程束数：32
设备上多处理器的数量： 68
======================================================
......

10块GPU一共输出10组