cuda-2-程序框架

cuda 程序框架

头文件包含
常量定义（或者宏定义）
C++ 自定义函数和CUDA核函数的声明（原型）
int main(void)
{
    分配主机与设备内存
    初始化主机中的数据
    将某些数据从主机复制到设备
    调用核函数在设备中进行计算
    将某些数据从设备复制到主机
    释放主机与设备内存
}
C++自定义函数与CUDA核函数的定义（实现）

一些常用函数

限定符的使用

限定符	执行	调用	备注
`__global__`	设备端执行	可以从主机调用也可以从计算能力3以上的设备调用	必须有一个void的返回类型
`__device__`	设备端执行	设备端调用
`__host__`	主机端执行	主机调用	可以省略

获取GPU设备数量

1	`__host__ __device__ cudaError_t cudaGetDeviceCount(int* devices);`

功能描述：

输入参数 int* count 是一个指向整型变量的指针，用于接收当前系统的CUDA设备数量
函数执行成功时，将填充 count 变量
如果没有可用的CUDA设备，那么 count 将被设置为0，并返回相应的错误状态

1 2	`int iDeviceCount = 0; cudaGetDeviceCount(&iDeviceCount);`

设置GPU执行时使用的设备

1 2	`int iDev = 0; cudaSetDevice(iDev);`

cudaSetDevice只能在主机上运行（不然已经到GPU上哪还需要设置）

内存管理（操作 global memory）

CUDA通过内存分配，数据传递，内存初始化，内存释放进行内存管理

C语言	CUDA语言	说明
malloc	cudaMalloc	内存分配
memcpy	cudaMemcpy	内存复制
memset	cudaMemset	内存设置
free	cudaFree	释放内存

数据分配：

1	`__host__ __device__ cudaFrror_t cudaMalloc(void **devPtr, size_t size)`

devPtr：用于接受分配内存的地址（是一个指向指针的指针，类型为 void）
size：用于指定内存分配的大小

// 主机分配内存：
float *fpHost_A;
fpHost_A = (float*)malloc(nBtytes);
// 设备分配内存：
float *fpDevice_A;
cudaMalloc((float**)&fpDevice_A, nBytes);

数据拷贝：

1	`__host__ cudaFrror_t cudaMemcpy(void dest, const void src, size_t count, cudaMemcpyKind kind)`

dest：用于接受拷贝的目的地址
src：用于拷贝的源地址
count：用于指定拷贝的长度
kind：用于指定拷贝从哪里到哪里，有四种可能：

Kind	拷贝方向
cudaMemcpyHostToHost	主机->主机
cudaMemcpyHostToDevice	主机->设备
cudaMemcpyDeviceToHost	设备->主机
cudaMemcpyDeviceToDevice	设备->设备
cudaMemcpyDefault	默认

默认方式只允许在支持同意虚拟寻址的系统上使用（自动判断拷贝方向）

内存初始化：

1	`__host__ cudaFrror_t cudaMemset(void *devPtr, int value, size_t count)`

devPtr：用于接受内存初始化的地址
value：用于指定内存初始化的值
count：用于指定内存初始化的长度

内存释放：

1	`__host__ __device__ cudaFrror_t cudaFree(void *devPtr)`

devPtr：用于接受内存释放的地址

一个经常会发生的错误就是混用设备和主机的内存地址！

为了避免这个情况，建议每个地址后面接上一个 _d 或者 _h 表示是主机端（host）还是设备端（device）

核函数

核函数就是在CUDA模型上诸多线程中运行的那段串行代码，这段代码在设备上运行，用NVCC编译，产生的机器码是GPU的机器码，所以我们写CUDA程序就是写核函数，第一步我们要确保核函数能正确的运行产生正确的结果，第二优化CUDA程序的部分，无论是优化算法，还是调整内存结构，线程结构都是要调整核函数内的代码来完成这些优化的

Kernel核函数编写有以下限制

只能访问设备内存
必须为 void 返回类型
不支持静态变量
显示异步行为

启动核函数

启动核函数，通过的以下的 ANSI C 扩展出的 CUDA C 指令：

1	`kernel_name<<<grid,block>>>(argument list);`

其标准C的原型就是C语言函数调用

1	`function_name(argument list);`

<<<grid,block>>> 内是对设备代码执行的线程结构的配置（或者简称为对内核进行配置），也就是我们上一篇中提到的线程结构中的 grid，block。我们通过 CUDA C 内置的数据类型 dim3 类型的变量来配置 grid 和 block

线程管理

当核函数开始执行，组织 GPU 的线程就变成了最主要的问题了，我们必须明确，一个核函数只能有一个 grid，一个 grid 可以有很多个块，每个块可以有很多的线程，这种分层的组织结构使得我们的并行过程更加自如灵活：
kernel

一个线程块block中的线程可以完成下述协作：

同步（相互影响，注意竞争冒险的问题）
共享内存

不同块内线程不能相互影响！他们是物理隔离的！

1	`kernel_name<<<4,8>>>(argument list);`

线程布局是：

为了给每个线程一个编号了，我们知道每个线程都执行同样的一段串行代码，那么怎么让这段相同的代码对应不同的数据呢？首先第一步就是让这些线程彼此区分开，才能对应到相应从线程，使得这些线程也能区分自己的数据。如果线程本身没有任何标记，那么没办法确认其行为。
依靠下面两个内置结构体确定线程标号：

blockIdx（线程块在线程网格内的位置索引）
threadIdx（线程在线程块内的位置索引）

这两个内置结构体基于 uint3 定义，包含三个无符号整数的结构，通过三个字段来指定：

blockIdx.x，blockIdx.y，blockIdx.z
threadIdx.x，threadIdx.y，threadIdx.z

上面这两个是坐标，当然我们要有同样对应的两个结构体来保存其范围，也就是 blockIdx 中三个字段的范围 threadIdx 中三个字段的范围：

blockDim，gridDim

他们是 dim3 类型(基于uint3定义的数据结构)的变量，也包含三个字段 x，y，z

#include <stdio.h>
#include <cuda_runtime.h>

__global__ void checkIndex(void){
    printf("threadIdx:(%d,%d,%d) blockIdx:(%d,%d,%d) blockDim:(%d,%d,%d) \
    gridDim:(%d,%d,%d)\n", threadIdx.x, threadIdx.y, threadIdx.z, blockIdx.x, blockIdx.y, blockIdx.z,\
     blockDim.x, blockDim.y, blockDim.z, gridDim.x, gridDim.y, gridDim.z);
     // there is no threadDim.z, threadDim.y, threadDim.z, gridIdx.x and gridIdx.y ...  
}

int main(){
    int nElem = 6;  // number elements of thread block
    dim3 block(3); // number of threads in a block, x=3; y=1, z=1 in default
    dim3 grid = (nElem + block.x - 1) / block.x; // number of blocks in grid
    printf("grid.x %d grid.y %d grid.z %d\n", grid.x, grid.y, grid.z);


    // CUDA中的线程和块是并行执行的，但它们的输出是通过标准输出（stdout）串行化的
    // 输出的顺序可能会表现出某种顺序性，但这并不意味着线程和块是按照这种顺序执行的
    checkIndex<<<grid, block>>>();
    cudaDeviceReset();
}

dim3 数据类型是通过括号进行初始化（顺序为 x,y,z ），通过结构体访问属性的方式修改的

1	`dim3 block(16, 8, 4); // x 维度为 16，y 维度为 8，z 维度为 4`

应用实例：cuda实现矩阵加法

#include <stdio.h>

__global__ void addFromGPU(float *A, float *B, float *C, const int N)
{
    const int bid = blockIdx.x;
    const int tid = threadIdx.x;
    const int id = tid + bid * blockDim.x; 

    C[id] = A[id] + B[id];
    
}

void initialData(float *addr, int elemCount);
void setGPU();

int main(void)
{
    // 1、设置GPU设备
    setGPU();

    // 2、分配主机内存和设备内存，并初始化
    int iElemCount = 512;                               // 设置元素数量
    size_t stBytesCount = iElemCount * sizeof(float);   // 字节数
    
    // （1）分配主机内存，并初始化
    float *fpHost_A, *fpHost_B, *fpHost_C;
    fpHost_A = (float *)malloc(stBytesCount);
    fpHost_B = (float *)malloc(stBytesCount);
    fpHost_C = (float *)malloc(stBytesCount);
    if (fpHost_A != NULL && fpHost_B != NULL && fpHost_C != NULL) {
        memset(fpHost_A, 0, stBytesCount);  // 主机内存初始化为0
        memset(fpHost_B, 0, stBytesCount);
        memset(fpHost_C, 0, stBytesCount);
    }
    else {
        printf("Fail to allocate host memory!\n");
        exit(-1);
    }

    // （2）分配设备内存，并初始化
    float *fpDevice_A, *fpDevice_B, *fpDevice_C;
    cudaMalloc((float**)&fpDevice_A, stBytesCount);
    cudaMalloc((float**)&fpDevice_B, stBytesCount);
    cudaMalloc((float**)&fpDevice_C, stBytesCount);
    if (fpDevice_A != NULL && fpDevice_B != NULL && fpDevice_C != NULL) {
        cudaMemset(fpDevice_A, 0, stBytesCount);  // 设备内存初始化为0
        cudaMemset(fpDevice_B, 0, stBytesCount);
        cudaMemset(fpDevice_C, 0, stBytesCount);
    }
    else {
        printf("fail to allocate memory\n");
        free(fpHost_A);
        free(fpHost_B);
        free(fpHost_C);
        exit(-1);
    }

    // 3、初始化主机中数据
    srand(666); // 设置随机种子
    initialData(fpHost_A, iElemCount);
    initialData(fpHost_B, iElemCount);
    
    // 4、数据从主机复制到设备
    cudaMemcpy(fpDevice_A, fpHost_A, stBytesCount, cudaMemcpyHostToDevice); 
    cudaMemcpy(fpDevice_B, fpHost_B, stBytesCount, cudaMemcpyHostToDevice); 
    cudaMemcpy(fpDevice_C, fpHost_C, stBytesCount, cudaMemcpyHostToDevice);


    // 5、调用核函数在设备中进行计算
    dim3 block(32);
    dim3 grid(iElemCount / 32);

    addFromGPU<<<grid, block>>>(fpDevice_A, fpDevice_B, fpDevice_C, iElemCount);    // 调用核函数
    // cudaDeviceSynchronize();

    // 6、将计算得到的数据从设备传给主机
    cudaMemcpy(fpHost_C, fpDevice_C, stBytesCount, cudaMemcpyDeviceToHost);


    for (int i = 0; i < 10; i++) {
        printf("idx=%2d\tmatrix_A:%.2f\tmatrix_B:%.2f\tresult=%.2f\n", i+1, fpHost_A[i], fpHost_B[i], fpHost_C[i]);
    }

    // 7、释放主机与设备内存
    free(fpHost_A);
    free(fpHost_B);
    free(fpHost_C);
    cudaFree(fpDevice_A);
    cudaFree(fpDevice_B);
    cudaFree(fpDevice_C);

    cudaDeviceReset();
    return 0;
}


void setGPU() {
    int iDeviceCount = 0;
    cudaError_t error = cudaGetDeviceCount(&iDeviceCount);

    if(error != cudaSuccess || iDeviceCount == 0) {
        printf("no compatible GPU found\n");
        exit(-1);
    }
    else {
        printf("the count of GPU id %d\n", iDeviceCount);
    }

    int iDev = 0;
    error = cudaSetDevice(iDev);
    if (error != cudaSuccess) {
        printf("fail to set GPU 0 for computing\n");
        exit(-1);
    }
    else {
        printf("set GPU 0 for computing");
    }
}


void initialData(float *addr, int elemCount) {
    for (int i = 0; i < elemCount; i++) {
        addr[i] = (float)(rand() & 0xFF) / 10.f;
    }
    return;
}

在开发阶段，每一步都进行验证是绝对高效的，比把所有功能都写好，然后进行测试这种过程效率高很多，同样写CUDA也是这样的每个代码小块都进行测试，看起来慢，实际会提高很多效率
CUDA小技巧，当我们进行调试的时候可以把核函数配置成单线程的：

1	`kernel_name<<<1,1>>>(argument list)`

概率论熵上一篇

cuda-1(兼容性) 下一篇