CUDA学习系列之cuda-runtime-api(二)

AI_央木

已于 2022-08-25 20:07:05 修改

阅读量631

点赞数 1

文章标签：学习 linux c++

于 2022-08-25 19:48:21 首次发布

本文链接：https://blog.csdn.net/m0_46464037/article/details/126530932

版权

1. 认识gridDim，blockIdx，blockDim，threadIdx

1.1 CUDA从逻辑上将GPU线程分成了三个层次——线程格（grid）、线程块（block）和线程（thread）。GPU硬件架构

每个核函数对应一个线程格，一个线程格中有一个或多个线程块，一个线程块中有一个或多个线程。

我们可以将Grid想象为一栋楼，将Block想象为楼里面的房间，而Thread就是房间里面的工作人员。这样，启动一个核函数就像将一项任务交给一栋楼来完成，楼将任务分解给各个房间，房间再将任务分解给各个工作人员。

1.2 线程全局索引

int idx = threadIdx.x + blockIdx.x * blockDim.x

这里有通用计算方法左乘右加

CUDA规定一个Block内最多包含1024个线程，Block每个维度上的最大数为(1024, 1024, 64），而Grid每个维度上的最大数为(2147483647, 65535 , 65535）

int main(){

    cudaDeviceProp prop;
    checkRuntime(cudaGetDeviceProperties(&prop, 0));

    // 通过查询maxGridSize和maxThreadsDim参数，得知能够设计的gridDims、blockDims的最大值
    // warpSize则是线程束的线程数量
    // maxThreadsPerBlock则是一个block中能够容纳的最大线程数，也就是说blockDims[0] * blockDims[1] * blockDims[2] <= maxThreadsPerBlock
    printf("prop.maxGridSize = %d, %d, %d\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
    printf("prop.maxThreadsDim = %d, %d, %d\n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);
    printf("prop.warpSize = %d\n", prop.warpSize);
    printf("prop.maxThreadsPerBlock = %d\n", prop.maxThreadsPerBlock);

    int grids[] = {1, 2, 3};     // gridDim.x  gridDim.y  gridDim.z 
    int blocks[] = {1024, 1, 1}; // blockDim.x blockDim.y blockDim.z 
    launch(grids, blocks);       // grids表示的是有几个大格子，blocks表示的是每个大格子里面有多少个小格子
    checkRuntime(cudaPeekAtLastError());   // 获取错误 code 但不清楚error
    checkRuntime(cudaDeviceSynchronize()); // 进行同步，这句话以上的代码全部可以异步操作
    printf("done\n");
    return 0;
}

prop.maxGridSize = 2147483647, 65535, 65535
prop.maxThreadsDim = 1024, 1024, 64
prop.warpSize = 32
prop.maxThreadsPerBlock = 1024
Run kernel. blockIdx = 0,0,0 threadIdx = 0,0,0
Run kernel. blockIdx = 0,1,2 threadIdx = 0,0,0
Run kernel. blockIdx = 0,0,2 threadIdx = 0,0,0
Run kernel. blockIdx = 0,1,1 threadIdx = 0,0,0
Run kernel. blockIdx = 0,0,1 threadIdx = 0,0,0
Run kernel. blockIdx = 0,1,0 threadIdx = 0,0,0

2. 核函数

2.1 编写核函数必须遵循CUDA规范，有哪些规范？
1）必须写在*.cu文件中
2）必须以__global__限定符声明定义；
3）返回类型必须是void；
4）不支持可变数量的参数；
5）核函数内部只能访问设备内存；
6）核函数内部不能使用静态变量。

2.2 函数声明中，__global__、__device__、__host__三者区别是什么？
1） __global__修饰的函数是核函数，在设备端执行，可以从主机端调用，也可以在sm3以上的设备端调用（比如动态并行）；
2） __device__修饰的函数是设备函数，在设备端执行，只能从设备端调用；
3） __host__修饰的函数是主机函数，在主机端执行，只能从主机端调用；
4） __device__和__host__可以一起使用，来表示该函数可以同时在主机端和设备端执行；
5） nvcc编译选项中添加-dc（相当于--relocatable-device-code=true --compile）时，__global__函数可以调用其它文件中的__device__函数，否则只能调用同文件中的__device__函数。

2.3 如何启动核函数？

启动CUDA核函数与启动C/C++函数很相似，只是额外添加了<<<>>>尖括号配置信息，尖括号内的配置信息并不是传递给核函数的，而是传递给CUDA运行时系统，告诉运行时系统如何启动核函数。

尖括号中包括四种信息，<<<块个数，线程个数，动态分配共享内存，流>>>，其中动态分配共享内存和流不是必填项。确定块个数和线程个数的一般步骤为：
1）先根据GPU设备的硬件资源确定一个块内的线程个数；
2）再根据数据大小和每个线程处理数据个数确定块个数。

__global__ void test(const float* pdata, int ndata){

    int idx = threadIdx.x + blockIdx.x * blockDim.x

    printf("Element[%d] = %f, threadIdx.x=%d, blockIdx.x=%d, blockDim.x=%d\n", \
            idx, pdata[idx], threadIdx.x, blockIdx.x, blockDim.x);
}

void test_print(condt float* pdata, int ndata){

    test<<<1, 10, 0, nullptr>>>(pdata, 10);
    
}

ref1.2.CUDA核函数 - 知乎