CUDA02_03CUDA编程入门与GPU模式

CUDA的核心就是扩展了C/C++语法，提出了核函数的语法，使得单一在CPU上运算的函数，可以指定在GPU上计算。同时提供辅助的API完成一些计算相关的操作。
CUDA的扩展语法还是采用PRO*C/C++等类似的思想，就是预编译，CUDA提供了一个nvcc的预编译工具，该工具可以自动调用本地编译器，实现完整的编译过程。工具根据扩展名来识别，cuda的扩展语法源代码扩展名是.cu。
因为Visual Studio貌似采用UTF-8BOM编码，使得微软cl编译器在编译CUDA的头文件的时候，出现各种警告。本主题在附录简单介绍了一下编译器中源代码编码的设置选项。
测试代码证明，GPU的运算快的一米。

CUAD设备管理

cuda驱动管理的参考文档
- https://docs.nvidia.com/cuda/cuda-driver-api/index.html
cuda驱动管理的库是：
- cuda.lib
这里主要介绍设备管理，其他后面使用的时候，再根据需要介绍
- 初始化
- 版本管理
- 设备管理
- 基本上下文管理
- 上下文管理
- 错误处理
根据C/C++的兼容习惯，一般会定义一套跨平台的类型。
- 使用的时候略加注意即可。

检查CUDA版本

初始化驱动APU
- CUresult cuInit ( unsigned int Flags )
  - 参数必须是0
  - 返回的类型是枚举类型：CUresult
    - CUDA_SUCCESS,
    - CUDA_ERROR_INVALID_VALUE,
    - CUDA_ERROR_INVALID_DEVICE,
    - CUDA_ERROR_SYSTEM_DRIVER_MISMATCH,
    - CUDA_ERROR_COMPAT_NOT_SUPPORTED_ON_DEVICE
- 该函数必须在调用任何驱动相关API之前调用。
获取CUDA的版本
- CUresult cuDriverGetVersion ( int* driverVersion )
  - 参数返回版本号，格式是(1000 major + 10 minor)，比如我的就是10010
  - 返回的也是枚举类型，错误返回
    - CUDA_ERROR_INVALID_VALUE
- 相关的两个API：
  - cudaDriverGetVersion
    - __host__ cudaError_t cudaDriverGetVersion ( int* driverVersion )
  - cudaRuntimeGetVersion
    - __host__ __device__ cudaError_t cudaRuntimeGetVersion ( int* runtimeVersion )
纯C/C++例子代码

    #include <stdio.h>
    #include <stdlib.h>
    #include <cuda.h>

    int main(int argc, char const *argv[])
    {
        // 1. 任何驱动API调用之前必须先初始化
        CUresult re = cuInit(0);
        if (re != CUDA_SUCCESS){
            printf("cuda设备初始化失败！\n");
            exit(EXIT_FAILURE);
        }
        printf("cuda设备初始化成功！\n");

        // 2. 获取版本
        int cu_version;
        re = cuDriverGetVersion (&cu_version);
        if(re == CUDA_SUCCESS){
            printf("获取版本是：%d\n", cu_version);
        }else{
            printf("获取版本失败！\n");
        }
        return 0;
    }

编译脚本


main: c01_version.c
    @nvcc -o main.exe  -l cuda  c01_version.c
clean:
    @del  *.exe *.exp  *.lib  2>nul

运行结果：

    C:\01works\02cuda\c03cuda_kernel>main
    cuda设备初始化成功！
    获取版本是：10010

CUDA扩展语法的例子代码
- CUDA的C/C++扩展语法主要提供了CPU运行与GPU设备运行代码的区分。
- __host__ cudaError_t cudaDriverGetVersion ( int* driverVersion )
- __host__ __device__ cudaError_t cudaRuntimeGetVersion ( int* runtimeVersion )
  - 返回类型：cudaError_t：typedef enumcudaError cudaError_t:也是枚举类型，这个函数返回下面两个值之一（这个枚举的值很多）
    - cudaSuccess
    - cudaErrorInvalidValue

这个函数的两个扩展字:__host__与__device__用来标识该函数能在什么设备上运行
- __host__：CPU
- __device__：GPU

    #include <stdio.h>
    #include <stdlib.h>
    #include <cuda.h>

    int main(int argc, char const *argv[])
    {
        int cu_version;
        cudaError_t  cu_re;

        // 1. 获取cuda版本
        cu_re = cudaDriverGetVersion(&cu_version);
        if(cu_re == cudaSuccess){
            printf("获取驱动版本是：%d\n", cu_version);
        }else{
            printf("获取驱动版本失败！\n");
        }

        // 2. 获取cuda运行时版本
        cu_re = cudaRuntimeGetVersion(&cu_version);
        if(cu_re == cudaSuccess){
            printf("获取运行时版本是：%d\n", cu_version);
        }else{
            printf("获取运行时版本失败！\n");
        }

        return 0;
    }

编译脚本

    main: c01_version.c
    @ # nvcc -o main.exe  -l cuda  c01_version.c
    @ nvcc -o main.exe  c02_version.cu
    clean:
    @del  *.exe *.exp  *.lib  2>nul

运行结果

    C:\01works\02cuda\c03cuda_kernel>main
    获取驱动版本是：10010
    获取运行时版本是：10010

设备信息检测

API
1. 获取设备列表与数量
  - CUresult cuDeviceGetCount ( int* count )
  - CUresult cuDeviceGet ( CUdevice* device, int ordinal )
2. 获取指定设备的name，id，属性，内存大小
  - CUresult cuDeviceGetAttribute ( int* pi, CUdevice_attribute attrib, CUdevice dev )
  - CUresult cuDeviceGetLuid ( char* luid, unsigned int* deviceNodeMask, CUdevice dev )
  - CUresult cuDeviceGetUuid ( CUuuid* uuid, CUdevice dev )
  - CUresult cuDeviceGetName ( char* name, int len, CUdevice dev )
  - CUresult cuDeviceTotalMem ( size_t* bytes, CUdevice dev )

CUuuid类型的定义

    typedef struct CUuuid_st {                                /**< CUDA definition of UUID */
        char bytes[16];
    } CUuuid;
    #endif

参数比较简答，返回值尽管很多，我们一般情况区分成功与失败即可。
- 其中属性比较繁琐一点，下面单独说明。

API的使用例子代码
- 编译脚本与运行效果，就不再罗列。

    #include <stdio.h>
    #include <stdlib.h>
    #include <cuda.h>

    int main(int argc, char const *argv[]){
        cuInit(0);   // 不做错误处理了！

        // 1. 获取设备的数量
        CUresult re;
        int count;
        re = cuDeviceGetCount(&count);
        if(re == CUDA_SUCCESS){
            printf("设备数量为：%d\n", count);
        }else{
            printf("设备数量获取失败！\n");
        }


        for(int i =0; i < count; i++){
            CUdevice device;
            re = cuDeviceGet (&device, i);
            if(re != CUDA_SUCCESS){
                printf("获取设备%d失败！\n", i);
                continue;
            }
            // 2. 获取设备的信息-uid
            CUuuid uid;
            cuDeviceGetUuid ( &uid,  device);
            printf("设备%d的uid：", i);
            for(int j = 0; j < 16; j++){
                printf("%02hhX", uid.bytes[j]);
            }
            printf("\n");
            // UUID：一共32为16字节，字符串表示如右：1859B803-EA0D-5A37-74E4-73F6F9B208FF
            // UUID格式：8-4-4-4-12
            // 3. 获取设备的信息-lid
            unsigned int lid;
            unsigned int mask;
            cuDeviceGetLuid ((char*)&lid, &mask, device);
            printf("设备%d的mask：%u\n", i, mask);
            printf("设备%d的uid：%u", i, lid);
            printf("设备节点 ：%u\n", lid & mask);

            // 4. 获取设备的信息-name
            char name[256] = {0};
            cuDeviceGetName (name, sizeof(name) - 1, device);
            printf("GPU设备名：%s\n", name);
            // 5. 获取设备的信息-mem
            size_t mem;
            cuDeviceTotalMem (&mem, device);
            printf("GPU内存大小：%zd Bytes\n", mem);
            printf("GPU内存大小：%zd MBytes\n", mem/1024/1024);
        }
        return 0;
    }

nvcc -o main.exe -l cuda c03_device.c

内存信息的cuda函数
- __host__cudaError_t cudaMemGetInfo ( size_t* free, size_t* total )
  - 返回总内存与剩余内存。

#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>

int main(int argc, char const *argv[]){

    size_t  mem_free, mem_total;
    cudaMemGetInfo (&mem_free, &mem_total);
    printf("总内存：%zd Mbytes\n", mem_total/1024/1024);
    printf("空余内存：%zd MBytes\n", mem_free/1024/1024);
    return 0;
}

编译指令：nvcc -o main.exe c04_device_mem.cu

设备属性检测

cuDeviceGetAttribute函数
- CUresult cuDeviceGetAttribute ( int* pi, CUdevice_attribute attrib, CUdevice dev )
  1. pi是返回的属性值
  2. attrib指定需要查询的属性；
  3. 指定设备

可以查询的属性，使用枚举类型预先定义，属性的枚举值太多，这儿不一一；列出，使用例子代码说明几个常用的.

#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>


int main(int argc, char const *argv[]){
    // 为了代码的简洁性，省略了异常处理
    cuInit(0);   
    CUdevice device;
    cuDeviceGet (&device, 0);   // 本机只有一块GPU

    // 属性1-兼容的最低版本
    int major, minor;
    cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, device);
    cuDeviceGetAttribute(&minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, device);
    printf("兼容版本：%d.%d\n", major, minor);

    // 属性2-CPU数量
    int  core_count;
    cuDeviceGetAttribute(&core_count, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, device);
    printf("GPU设备上的核数量：%d\n", core_count);

    // 属性3-时钟频率
    int kernel_rate;
    cuDeviceGetAttribute(&kernel_rate, CU_DEVICE_ATTRIBUTE_CLOCK_RATE, device);
    printf("GPU时钟频率：%0.2fGHz\n", kernel_rate*1e-6f);

    // 属性4-内存频率与带宽
    int mem_rate;
    cuDeviceGetAttribute(&mem_rate, CU_DEVICE_ATTRIBUTE_MEMORY_CLOCK_RATE, device);
    int bus_width;
    cuDeviceGetAttribute(&bus_width, CU_DEVICE_ATTRIBUTE_GLOBAL_MEMORY_BUS_WIDTH, device);
    printf("GPU内存时钟频率：%0.2fGHz，内存带宽：%d比特\n", mem_rate*1e-6f, bus_width);
    // 属性5-L2缓存大小
    int l2_size;
    cuDeviceGetAttribute(&l2_size, CU_DEVICE_ATTRIBUTE_L2_CACHE_SIZE, device);
    printf("L2缓存大小：%d MB\n", l2_size/1024/1024);
    return 0;
}

编译命令：nvcc -o main.exe -l cuda c05_attribute.c
运行结果：


C:\01works\02cuda\c03cuda_kernel>main
兼容版本：6.6
GPU设备上的核数量：10
GPU时钟频率：1.67GHz
GPU内存时钟频率：4.00GHz，内存带宽：192比特
L2缓存大小：1 MB

cudaDeviceGetAttribute函数
- 这个函数是可以在GPU上运行的。
- 参数与上面一样，只是枚举类型的拼写不同。
- __host__ __device__ cudaError_t cudaDeviceGetAttribute ( int* value, cudaDeviceAttr attr, int device )

#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>


int main(int argc, char const *argv[]){
    // 为了代码的简洁性，省略了异常处理
    int kernel_count;
    cudaDeviceGetAttribute (&kernel_count, cudaDevAttrMultiProcessorCount, 0);
    printf("GPU设备上的核数量：%d\n", kernel_count);
    return 0;
}

// nvcc -o main.exe  c06_attribute.cu

cudaGetDeviceProperties函数
- __host__ cudaError_t cudaGetDeviceProperties ( cudaDeviceProp* prop, int device )
  - 返回结构，结构包含基本上所有属性。

#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>


int main(int argc, char const *argv[]){
    // 为了代码的简洁性，省略了异常处理
    cudaDeviceProp prop;
    cudaGetDeviceProperties (&prop, 0);

    printf("GPU设备上的核数量：%d\n", prop.multiProcessorCount);
    return 0;
}

// nvcc -o main.exe c07_properties.cu

CUDA基本编程模式

CUDA扩展C语法

cuda扩展了一种可以在核上调用的函数：核函数，这个函数的语法扩展体现在两个方面：
1. 定义需要指定该函数在主机上运行，还是在GPU设备上运行：
  - __device__:GPU上执行，GPU上调用；
  - __host__：CPU上执行，CPU上调用；
  - __global__：CPU上调用，GPU上执行；
2. 调用函数需要指定分配的块数，以及块上的线程数。
  - MyKernel<<<blocksPerGrid, THREADS_PER_BLOCK>>>(函数的参数);
    - 如果是在GPU上运行，则需要传递GPU内存地址。

GPU调用函数定义

其中几个变量blockDim，blockIdx，threadIdx后面解释，是扩展C/C++语法的一部分。

#include <stdio.h>
#include <cuda.h>

/*
  1. 定义核函数：
*/
__global__ void kernel_add(float* a, float *b, float *r){

    int i = blockDim.x * blockIdx.x + threadIdx.x;
    r[i] = a[i] + b[i];
}

GPU内存申请与拷贝

函数定义
- 内存分配：
  - __host__ __device__ cudaError_t cudaMalloc ( void** devPtr, size_t size )
- 内存拷贝：
  - __host__ cudaError_t cudaMemcpy ( void* dst, const void* src, size_t count, cudaMemcpyKind kind )
代码实现

    // 2. GPU内存分配与初始化
    // 2.1. CPU数据
    int n = 1024 * 1024;
    int size = n * sizeof(float);

    float *cpu_a, *cpu_b, *cpu_r;
    cpu_a = (float*)malloc(size);
    cpu_b = (float*)malloc(size);
    cpu_r = (float*)malloc(size);

    // 初始化数cpu内存
    for(int i=0; i < n; i++){
        cpu_a[i] = 90.0;
        cpu_b[i] = 10.0;
    }
    
    // 2.2. GPU内存
    float *gpu_a, *gpu_b, *gpu_r;
    cudaMalloc((void**)&gpu_a, size);
    cudaMalloc((void**)&gpu_b, size);
    cudaMalloc((void**)&gpu_r, size);
    
    // 2.3. GPU数据
    cudaMemcpy(gpu_a, cpu_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(gpu_b, cpu_b, size, cudaMemcpyHostToDevice);

在GPU上调用函数

需要指定函数调用需要的GPU线程资源。
代码实现：
- 我们分配了512块
- 每个块分配的线程是1024 * 1024 / 512个线程。

    // 3. 分配GPU
    dim3 dimBlock(512);       // 块数 blockDim封装了块的描述
    dim3 dimGrid(n/512);      // 每块的线程数，blockIdx封装了块中kernel的索引
                              // kernel被多个线程调用，线程的描述封装为threadIdx

    // 4. 调用kernel
    kernel_add<<<dimGrid, dimBlock>>>(gpu_a, gpu_b, gpu_r);

从GPU拷贝运算结果

运算结束，需要从GPU把数据拷贝到CPU内存，才能访问。

    // 5. GPU返回数据
    cudaMemcpy(cpu_r, gpu_r, size, cudaMemcpyDeviceToHost);

    printf("%f\n", cpu_r[1]);

计算结束，需要释内存资源
- 包含CPU与GPU上资源

    // 6. GPU内存释放
    cudaFree(gpu_a);
    cudaFree(gpu_b);
    cudaFree(gpu_r);

    free(cpu_a);
    free(cpu_b);
    free(cpu_r);

完整的程序代码

包含结尾附加了编译命令。

#include <stdio.h>
#include <cuda.h>

/*
  1. 定义核函数：
*/
__global__ void kernel_add(float* a, float *b, float *r){

    int i = blockDim.x * blockIdx.x + threadIdx.x;
    r[i] = a[i] + b[i];
}

int main(int argc, char const *argv[]){
    
    // 2. GPU内存分配与初始化
    // 2.1. CPU数据
    int n = 1024 * 1024;
    int size = n * sizeof(float);

    float *cpu_a, *cpu_b, *cpu_r;
    cpu_a = (float*)malloc(size);
    cpu_b = (float*)malloc(size);
    cpu_r = (float*)malloc(size);

    // 初始化数cpu内存
    for(int i=0; i < n; i++){
        cpu_a[i] = 90.0;
        cpu_b[i] = 10.0;
    }
    
    // 2.2. GPU内存
    float *gpu_a, *gpu_b, *gpu_r;
    cudaMalloc((void**)&gpu_a, size);
    cudaMalloc((void**)&gpu_b, size);
    cudaMalloc((void**)&gpu_r, size);
    
    // 2.3. GPU数据
    cudaMemcpy(gpu_a, cpu_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(gpu_b, cpu_b, size, cudaMemcpyHostToDevice);

    // 3. 分配GPU
    dim3 dimBlock(512);       // 块数 blockDim封装了块的描述
    dim3 dimGrid(n/512);      // 每块的线程数，blockIdx封装了块中kernel的索引
                              // kernel被多个线程调用，线程的描述封装为threadIdx

    // 4. 调用kernel
    kernel_add<<<dimGrid, dimBlock>>>(gpu_a, gpu_b, gpu_r);

    // 5. GPU返回数据
    cudaMemcpy(cpu_r, gpu_r, size, cudaMemcpyDeviceToHost);

    printf("%f\n", cpu_r[1]);

    // 6. GPU内存释放
    cudaFree(gpu_a);
    cudaFree(gpu_b);
    cudaFree(gpu_r);

    free(cpu_a);
    free(cpu_b);
    free(cpu_r);

    return 0;
}

// nvcc -o main.exe  -Xcompiler /source-charset:utf-8  c08kernel_function.cu

运行结果

GPU运算结果

观察GPU运算过程中的性能

下面是循环10000此调用产生的一个峰值。后面还会使用更加复杂的计算来评估这个性能优化的效果。

GPU的利用率

附录

1. 关于编译器对源代码编码的识别

因为VS2019默认支持的UTF-8编码支持带BOM的UTF-8，使用nvcc编译的时候，cl会报一大堆与编码UTF-8有关的警告。解决办法有两种：
- 关闭警告。这个只是回避问题，并没有解决问题；
  - 关闭警告选项：-w
- 指定编码器使用的编码；
指定对源代码处理的编码
1. 使用cl的编译选项（对所有源代码有效）：
  - cl /Femain.exe /source-charset:utf-8 /link /execution-charset:utf-8 c01_readbmp.c
2. 使用#pragma预处理扩展指令（只对当前文件的代码有效）：
  - #pragma source_character_set("utf-8")
指定连接器使用的编码
1. 使用link的链接选项
  - cl /Femain.exe /source-charset:utf-8 /link /execution-charset:utf-8 c01_readbmp.c
2. 使用#pragma`预处理扩展指令
  - #pragma execution_character_set("utf-8")
在nvcc中指定
- nvcc -o main.exe -Xcompiler /source-charset:utf-8 c07_properties.cu
cuda的h文件的编码问题：
- 这个建议在编译命令行中指定代码的编码处理，如上。
- 下面是编译器对cuda的头文件的编码不一致导致的。

默认编译情况下产生的编码警告

使用选项指定编码可以解决

指定编码后，警告消失

人面猴
序言：七十年代末，一起剥皮案震惊了整个滨河市，随后出现的几起案子，更是在滨河造成了极大的恐慌，老刑警刘岩，带你破解...
沈念sama阅读 216,544评论 6赞 501
死咒
序言：滨河连续发生了三起死亡事件，死亡现场离奇诡异，居然都是意外死亡，警方通过查阅死者的电脑和手机，发现死者居然都...
沈念sama阅读 92,430评论 3赞 392
救了他两次的神仙让他今天三更去死
文/潘晓璐我一进店门，熙熙楼的掌柜王于贵愁眉苦脸地迎上来，“玉大人，你说我怎么就摊上这事。” “怎么了？”我有些...
开封第一讲书人阅读 162,764评论 0赞 353
道士缉凶录：失踪的卖姜人
文/不坏的土叔我叫张陵，是天一观的道长。经常有香客问我，道长，这世上最难降的妖魔是什么？我笑而不...
开封第一讲书人阅读 58,193评论 1赞 292
港岛之恋（遗憾婚礼）
正文为了忘掉前任，我火速办了婚礼，结果婚礼上，老公的妹妹穿的比我还像新娘。我一直安慰自己，他们只是感情好，可当我...
茶点故事阅读 67,216评论 6赞 388
恶毒庶女顶嫁案：这布局不是一般人想出来的
文/花漫我一把揭开白布。她就那样静静地躺着，像睡着了一般。火红的嫁衣衬着肌肤如雪。梳的纹丝不乱的头发上，一...
开封第一讲书人阅读 51,182评论 1赞 299
城市分裂传说
那天，我揣着相机与录音，去河边找鬼。笑死，一个胖子当着我的面吹牛，可吹牛的内容都是我干的。我是一名探鬼主播，决...
沈念sama阅读 40,063评论 3赞 418
双鸳鸯连环套：你想象不到人心有多黑
文/苍兰香墨我猛地睁开眼，长吁一口气：“原来是场噩梦啊……” “哼！你这毒妇竟也来了？” 一声冷哼从身侧响起，我...
开封第一讲书人阅读 38,917评论 0赞 274
万荣杀人案实录
序言：老挝万荣一对情侣失踪，失踪者是张志新（化名）和其女友刘颖，没想到半个月后，有当地人在树林里发现了一具尸体，经...
沈念sama阅读 45,329评论 1赞 310
护林员之死
正文独居荒郊野岭守林人离奇死亡，尸身上长有42处带血的脓包…… 初始之章·张勋以下内容为张勋视角年9月15日...
茶点故事阅读 37,543评论 2赞 332
白月光启示录
正文我和宋清朗相恋三年，在试婚纱的时候发现自己被绿了。大学时的朋友给我发了我未婚夫和他白月光在一起吃饭的照片。...
茶点故事阅读 39,722评论 1赞 348
活死人
序言：一个原本活蹦乱跳的男人离奇死亡，死状恐怖，灵堂内的尸体忽然破棺而出，到底是诈尸还是另有隐情，我是刑警宁泽，带...
沈念sama阅读 35,425评论 5赞 343
日本核电站爆炸内幕
正文年R本政府宣布，位于F岛的核电站，受9级特大地震影响，放射性物质发生泄漏。R本人自食恶果不足惜，却给世界环境...
茶点故事阅读 41,019评论 3赞 326
男人毒药：我在死后第九天来索命
文/蒙蒙一、第九天我趴在偏房一处隐蔽的房顶上张望。院中可真热闹，春花似锦、人声如沸。这庄子的主人今日做“春日...
开封第一讲书人阅读 31,671评论 0赞 22
一桩弑父案，背后竟有这般阴谋
文/苍兰香墨我抬头看了看天上的太阳。三九已至，却和暖如春，着一层夹袄步出监牢的瞬间，已是汗流浃背。一阵脚步声响...
开封第一讲书人阅读 32,825评论 1赞 269
情欲美人皮
我被黑心中介骗来泰国打工，没想到刚下飞机就差点儿被人妖公主榨干…… 1. 我叫王不留，地道东北人。一个月前我还...
沈念sama阅读 47,729评论 2赞 368
代替公主和亲
正文我出身青楼，却偏偏与公主长得像，于是被迫代替她去往敌国和亲。传闻我的和亲对象是个残疾皇子，可洞房花烛夜当晚...
茶点故事阅读 44,614评论 2赞 353

CUDA02_03CUDA编程入门与GPU模式

CUAD设备管理

检查CUDA版本

设备信息检测

设备属性检测

CUDA基本编程模式

CUDA扩展C语法

GPU调用函数定义

GPU内存申请与拷贝

在GPU上调用函数

从GPU拷贝运算结果

完整的程序代码

运行结果

观察GPU运算过程中的性能

附录

1. 关于编译器对源代码编码的识别

推荐阅读更多精彩内容