06.CUDA编程模型概述（二）

不要那么孤单

已于 2024-08-25 16:05:00 修改

阅读量967

点赞数 25

分类专栏：从零开始学cuda 文章标签：语言模型机器学习算法 c语言

于 2024-07-16 23:51:31 首次发布

本文链接：https://blog.csdn.net/sunhjb/article/details/140480083

版权

从零开始学cuda 专栏收录该内容

15 篇文章 0 订阅

订阅专栏

4. CUDA核函数

CUDA核函数：在GPU上并行执行的函数称为CUDA核函数（Kernel Function)，它属于CUDA编程中最为重要且核心的一个环节，也是我们重点要写的代码部分。我们先从启动一个核函数开始！

5. 启动一个CUDA核函数

C语言中函数调用语句：

function_name(argument list);

CUDA内核调用是对C语言函数调用语句的扩展，<<<>>>运算符内是核函数的执行配置，即需要指定网格和块的维度。

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

执行配置的第一个值是网格维度，也就是启动块的数目。第二个值是块维度，也就是每个块中线程的数目。通过指定网格和块的维度，你可以进行以下配置：

内核中线程的数目
内核中使用的线程布局

同一个块中的线程之间可以相互协作，不同块内的线程不能协作。对于一个给定的问题，可以使用不同的网格和块布局来组织你的线程。例如，假设你有32个数据元素用于计算，每8个元素一个块，需要启动4个块。我们可以使用dim3类型的grid维度和block维度配置内核，也可以使用int类型的变量，或者常量直接初始化，如下：

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

由于数据在全局内存中是线性存储的，因此可以用变量blockIdx.x和threadId.x来进行以下操作。

在网格中标识一个唯一的线程
建立线程和数据元素之间的映射关系

如果把所有32个元素放到一个块里，那么只会得到一个块：

kernel_name<<<1,32>>>(argument list);

如果每个块只含有一个元素，那么会有32个块：

kernel_name<<<32,1>>>(argument list);

核函数的调用与主机线程是异步的。核函数调用结束后，控制权立刻返回给主机端。你可以调用以下函数来强制主机端程序等待所有的核函数执行结束：

cudaError_t cudaDeviceSynchronize(void);

这是一个显示的方法，对应的也有隐式方法，隐式方法就是不明确说明主机要等待设备端，而是设备端不执行完，主机没办法进行，比如内存拷贝函数：

cudaError_t cudaMemcpy(void* dst,const void * src,size_t count,cudaMemcpyKind kind);

6. 编写核函数

核函数用__global__符号声明，在devie（GPU）上执行，在host（CPU）上调用，返回类型必须时void，不支持可变参数，不支持静态变量，不支持函数指针，核函数相对于CPU是异步的，在核函数执行完之前就会返回，这样CPU可以不用等待核函数的完成，继续执行后续代码。在host端核函数的调用方式为：

__global__ void kernel_name(argument list);

我们之前讲过CPU-GPU的异构模型，所以需要区分host和device上的代码，在CUDA中是通过函数类型限定词来区分host和device上的函数，主要的三个函数类型限定词如下：

限定符	执行	调用	备注
__ global__	设备端执行	可以从主机调用，也可以从计算能力3以上的设备调用	必须有一个void的返回类型
__ device__	设备端执行	设备端调用
__ host __	主机端执行	主机端调用	一般忽略不写

__device__和__host__限定符可以一齐使用，这样函数可以同时在主机和设备端进行
编译。

6.1 CUDA核函数的限制

核函数编写有以下限制：

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

考虑一个简单的例子：循环将两个大小为N的向量A和B相加，主机端代码：

void sumArraysOnHost(float *A, float *B, float *C, const int N) {
  for (int i = 0; i < N; i++)
    C[i] = A[i] + B[i];
}

设备端代码：

__global__ void sumArraysOnGPU(float *A, float *B, float *C) {
  int i = threadIdx.x;
  C[i] = A[i] + B[i];
}

可以看到设备端代码中的循环体消失了，内置的线程坐标变量替换了数组索引。

7. 验证核函数

开发过程中，验证自己的代码是必须的且高效的。这里的验证方法是为了验证我们最开始写的两个数组相加的程序，通过比较两个数组在主机端和设备端分别执行相加得出的结果，来进行验证核函数即设备端的代码是否正确。

#include <stdio.h>
#include <math.h>
void checkResult(float *hostRef, float *gpuRef, const int N){
    double epsilon = 1.0E-8 ;
    int match = 1;
    for (int i=0; i<N; i++){
        if (fabsf(hostRef[i] - gpuRef[i]) > epsilon){
            match = 0;
            printf("Arrays do not match!\n");
            printf("host %5.2f gpu %5.2f at current %d\n", hostRef[i], gpuRef[i], i);
            break;
        }
    }
    if (match){
        printf("Arrays match.\n\n");
        return;
    }
}

math.h：提供了一组用于数学计算的函数，这些函数涵盖了基本的算术运算、三角函数、指数函数、对数函数、幂函数和其他一些常用的数学操作。

abs、fabs 和 fabsf 是C/C++标准库中用于取绝对值的函数，它们的主要区别在于处理的数据类型和返回值类型。

abs：

用于整数数据类型（如 int、long）。
返回一个整数类型的绝对值。
不需要包含额外的头文件，通常在或 <stdlib.h> 中定义。

fabs：

用于双精度浮点数数据类型（double）。
返回一个双精度浮点数的绝对值。
需要包含或 <math.h> 头文件。

fabsf：

用于单精度浮点数数据类型（float）。
返回一个单精度浮点数的绝对值。
需要包含或 <math.h> 头文件。

两个非常简单实用的方法可以验证核函数。

可以在Fermi及更高版本的设备端的核函数中使用printf函数。
可以将执行参数设置为<<<1，1>>>，因此强制用一个块和一个线程执行核函数，这模拟了串行执行程序。这对于调试和验证结果是否正确是非常有用的，而且，如果你遇到了运算次序的问题，这有助于你对比验证数值结果是否是按位精确的。

8. 处理错误

由于许多CUDA调用是异步的，所以有时可能很难确定某个错误是由哪一步程序引起的。定义一个错误处理宏封装所有的CUDA API调用，这简化了错误检查过程：

#define CHECK(call)                                                             \
{                                                                               \
    const cudaError_t error = call;                                             \
    if (error != cudaSuccess)                                                   \
    {                                                                           \
        printf("Error: %s:%d, ", __FILE__, __LINE__);                           \
        printf("code: %d, reason: %s\n", error, cudaGetErrorString(error));     \
        exit(1);                                                                \
    }                                                                           \
}

宏定义：是C提供的三种预处理功能的其中一种。其主要目的是为程序员在编程时提供一定的方便，并能在一定程度上提高程序的运行效率。

#define命令是C语言中的一个宏定义命令，它用来讲一个标识符定义为一个字符串，该标识符被称为宏名，被定义的字符串称为替换文本。该命令有两种格式：一种是简单的宏定义（不带参数的宏定义），另一种是带参数的宏定义。
例：#define PI 3.1415926
例：#define S(a,b) a*b

注意：

宏名一般用大写
宏定义末尾不加分号 ;
可以用#undef命令终止宏定义的作用域
宏定义可以嵌套
字符串“”中永远不包含宏
宏名和参数的括号间不能有空格

你可以在以下代码中使用宏：

 CHECK(cudaMemcpy(d_c, gpuRef, nBytes, cudaMemcpyHostToDevice));

如果内存拷贝或之前的异步操作产生了错误，这个宏会报告错误代码，并输出一个可读信息，然后停止程序。也可以用下述方法，在核函数调用后检查核函数错误：

kernel_name<<<grid,block>>>(argument list);
CHECK(cudaDeviceSynchronize());

CHECK（cudaDeviceSynchronize（））会阻塞主机端线程的运行直到设备端所有的请求任务都结束，并确保最后的核函数启动部分不会出错。当然在release版本中可以去除这部分，但是开发的时候一定要有的。

9. 编译和执行

OK，现在把所有的代码放在一个文件名为sumArraysOnGPU-small-case.cu的文件中，如下：

完整代码：https://github.com/dive-into-cuda

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

#define CHECK(call)                                                             \
{                                                                               \
    const cudaError_t error = call;                                             \
    if (error != cudaSuccess)                                                   \
    {                                                                           \
        printf("Error: %s:%d, ", __FILE__, __LINE__);                           \
        printf("code: %d, reason: %s\n", error, cudaGetErrorString(error));     \
        exit(1);                                                                \
    }                                                                           \
} 

void checkResult(float *hostRef, float *gpuRef, const int N){
    double epsilon = 1.0E-8 ;
    int match = 1;
    for (int i=0; i<N; i++){
        if (fabsf(hostRef[i] - gpuRef[i]) > epsilon){
            match = 0;
            printf("Arrays do not match!\n");
            printf("host %5.2f gpu %5.2f at current %d\n", hostRef[i], gpuRef[i], i);
            break;
        }
    }
    if (match){
        printf("Arrays match.\n\n");
        return;
    }
}

void initialData(float *ip, int size){
    // 为随机数生成不同的种子
    // time_t是一个数据类型，用于表示时间
    time_t t;
    // &t获取变量t的地址
    srand((unsigned int) time(&t));

    for (int i=0; i<size; i++){
        ip[i] = (float)(rand() & 0xFF) / 10.0f;
    }
}

// 当a是一个指针的时候，*a就是这个指针指向的内存的值
// const含义：只要一个变量前用const来修饰，就意味着该变量里的数据只能被访问，而不能被修改，也就是意味着“只读”（readonly）
void sumArraysOnHost(float *A, float *B, float *C, const int N){
    for (int idx=0; idx<N; idx++){
        C[idx] = A[idx] + B[idx];
    }
}

__global__ void sumArraysOnGPU(float*a,float*b,float*res)
{
  int i=threadIdx.x;
  res[i]=a[i]+b[i];
}

int main(int argc, char **argv){
    printf("%s 开始...\n", argv[0]);

    // 设置设备
    int dev = 0;
    cudaSetDevice(dev);

    // 设置向量数据
    int nElem = 32;
    printf("向量大小为 %d\n", nElem);

    // 主机申请内存
    size_t nBytes = nElem * sizeof(float);

    float *h_A, *h_B, *hostRef, *gpuRef;
    h_A     = (float *)malloc(nBytes);
    h_B     = (float *)malloc(nBytes);
    hostRef = (float *)malloc(nBytes);
    gpuRef  = (float *)malloc(nBytes);

    // 主机端初始化数据
    initialData(h_A, nElem);
    initialData(h_B, nElem);

    memset(hostRef, 0, nBytes);
    memset(gpuRef,  0, nBytes);

    // 设备端申请全局内存
    float *d_A, *d_B, *d_C;
    CHECK(cudaMalloc((float**)&d_A, nBytes));
    CHECK(cudaMalloc((float**)&d_B, nBytes));
    CHECK(cudaMalloc((float**)&d_C, nBytes));

    // 将主机数据传到设备端
    CHECK(cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice));

    // 在主机端设置线程块，线程格
    dim3 block  (nElem);
    dim3 grid   (nElem/block.x);

    sumArraysOnGPU<<< grid, block >>>(d_A, d_B, d_C);
    printf("线程设置：<<<%d, %d>>>\n", grid.x, block.x);

    // 复制设备端结果到主机
    CHECK(cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost));

    // 主机端计算结果
    sumArraysOnHost(h_A, h_B, hostRef, nElem);

    // 对比设备端和主机端计算结果
    checkResult(hostRef, gpuRef, nElem);

    printf("打印数组A，数组B，主机结果，设备结果前5个数：\n");
    for(int i=0; i < 5; i++){
        printf("h_A[%d]=%f h_B[%d]=%f hostRef[%d]=%f gpuRef[%d]=%f\n", i, h_A[i], i, h_B[i], i, *hostRef[i], i, *gpuRef[i]);
    }

    // 释放设备端内存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    // 释放主机端内存
    free(h_A);
    free(h_B);
    free(hostRef);
    free(gpuRef);

    return 0;
}

size_t：是个无符号整型，它并不是一个全新的数据类型，更不是一个关键字。size_t是由typedef定义而来的，我们在很多标准库头文件中都能发现。size_t的含义是size type，是一种计数类型。取值范围与机器架构与操作系统相关。32 位机器一般是unsigned int，占 4 字节；而 64 位机器一般是unsigned long，占 8 字节。
size_t类型常被用作计数用途，例如：sizeof运算符得到对象所占的字节数；字符串函数strlen返回字符串的长度等等，其返回值都为size_t类型。

C语言的内存必须初始化。‌在C语言中，‌不仅新定义的变量需要初始化，‌新分配的内存空间也需要初始化。‌

memset：是计算机中C/C++语言初始化函数。作用是将某一块内存中的内容全部设置为指定的值，这个函数通常为新申请的内存做初始化工作，它是直接操作内存空间。memset 一般使用“0”初始化内存单元，而且通常是给数组或结构体进行初始化。一般的变量如 char、int、float、double 等类型的变量直接初始化即可，没有必要用 memset。如果用 memset 的话反而显得麻烦。
函数原型：void *memset(void *s, int c, unsigned long n);
将s中当前位置后面的n个字节（typedef unsigned int size_t ）用 c 替换并返回 s。

执行结果如下：
在这里插入图片描述