【CUDA】《CUDA编程:基础与实践》CUDA程序基本框架
创始人
2024-06-02 15:37:17
0

介绍

以向量加法为例介绍了CUDA程序得基本结构,使用xmake进行项目组织及编译。
相关代码:https://github.com/JilinLi4/LearningCUDA

一、 相关API

1. 设备CUDARuntime 初始化

在CUDA运行时API中,没有显示地初始化设备的函数。在第一次调用一个和设备管理及版本查询无关的运行时API函数时,设备将自动地初始化。

2. 设备内存的分配及释放

申请显存

cudaError_t cudaMalloc(void **address, size_t size);
  • address 待分配GPU内存的指针。
  • size 分配内存的字节数
  • 返回错误码,如果成功,返回 cudaSuccess

释放显存

cudaError_t cudaFree(void* address);

3. 主机与设备之间数据的传递

cudaError_t cudaMemcpy(void                *dst,const void          *src,size_t              count,enum cudaMemcpyKind kind
)
  • dst 是目标地址
  • src 是源地址
  • count 复制数据的字节数
  • kind 一个枚举类型的变量,取如下几个值:
    • cudaMemcpyHostToHost 从主机复制到主机
    • cudaMemcpyHostToDevice 从主机复制到设备
    • cudaMemcpyDevicetoHost 从设备复制到主机
    • cudaMemcpyDefault, 根据指针dst和src所指地址自行判断数据传输方向。

二、 核函数中数据与线程的对应

1. 使用多个线程的核函数

核函数中允许指派很多线程。因为一个GPU中有几千个计算核心,而总的线程数必须至少等于计算核心数才能充分利用GPU中的全部计算资源。
所以根据需要,在调用核函数时可以指定多个线程。比如:

func<<<2, 4>>>();

func指定了2个线程块,每个线程块中包含4个线程。总的线程数是 2×4=82 \times 4 = 82×4=8。也就是说,该程序中的核函数中代码执行方式是:“单指令-多线程”,即每一个线程都执行同一串指令。

2 使用线程索引

线程组织结构是由调用函数配置的。如:

func<<>>();

一般来说:

  • gridSize≤231−1\text{gridSize} \le 2^{31}-1gridSize≤231−1
  • blockSize≤1024\text{blockSize} \le 1024blockSize≤1024
    每个线程在核函数中都有一个唯一的身份表示。由于我们用两个参数指定线程数目。

3 核函数中的数据与线程对应

将有关数据从主机传至设备之后,就可以调用核函数在设备中进行计算了。使用以下线程布局:

  • blockSize=128\text{blockSize} = 128blockSize=128
  • gridSize=108/128\text{gridSize} = 10^8 / 128gridSize=108/128
    在这种情况下,核函数的数据与线程对应关系为:
const int n = blockDim.x * blockIdx.x + threadIdx.x;

4 是否需要使用 if 判断 n是否越界?

核函数没有使用参数N,当N是blockDim.x的整数倍时,不会引起问题,因为核函数的线程数刚好等于数组元素个数。然而,当N不是 blockDim.x的整数倍时,就有可能发生错误。例如将 N改为: 108+110^8 + 1108+1 而 block_size依然是128,如果griadSize的计算方式不变,那么必将有一个元素无法处理。所以我们应该将原来的计算方式:
gridSize=N/blockSize\text{gridSize} = N / \text{blockSize}gridSize=N/blockSize 改为
gridSize=(N+blockSize−1)/blockSize\text{gridSize} = (N + blockSize - 1) / \text{blockSize}gridSize=(N+blockSize−1)/blockSize

这样一来也会引入越界的风险,所以需要将线程索引的计算方式改为:

const int n = blockDim.x * blockIdx.x + threadIdx.x;
if(n >= N) return;
z[n] = x[n] + y[n];

三、 自定义核函数

1. 核函数的要求

  • 返回值必须是void,可以使用return 关键字,但是不可返回任何值
  • 必须使用限定符 global 也可以加上一些其他C++的限定符,如:static 限定符的次序任意
  • 函数名支持重载
  • 不可以向核函数传递非指针变量
  • 除非使用统一内存编程机制,否则传给核函数的指针必须只想设备内存
  • 核函数不可成为一个类的成员。通常做法是用一个包装函数调用核函数,将包装函数定义为一个类。
  • 核函数之间在算力 > 3.5 开始,可以调用其他核函数,也可以调用自己。
  • 核函数必须指定线程布局配置。

2. 函数执行空间标识符

  • __ global__ 称为核函数,一般由主机调用,也可以被核函数包括自己调用。
  • __ device__ 称为设备函数,只能被核函数或其他设备函数调用,在设备中执行。
  • __host __ 主机段普通的C++函数,在主机中被调用,在主机执行。可省略,存在的意义在于,可以使用**host**和 **device**修饰同一个函数,表示即可在设备端调用也可以在主机端调用,减少代码冗余。
  • 不能使用__global__ 和 __device__修饰统一个函数
  • 不能使用__host__和 __device__修饰统一个函数
  • 编译器决定把设备函数当做内联函数或非内联函数,但可以使用 __noinline__建议一个设备函数为非内联函数,也可以用 __forceinline__建议一个设备为内联函数。

3. 实例代码

3.1 code1 返回值

double __device__ add1_device(const double x, const double y) {return x + y;
}void __global__ add1(const double *x, const double *y, double *z, const int N) {const int n = blockDim.x * blockIdx.x + threadIdx.x;if(n < N) {z[n] = add1_device(x[n], y[n]);}
}

3.2 code2 传指针

void __device__ add2_device(const double x, const double y, double *z) {*z = *x + *y;
}void __global__ add1(const double *x, const double *y, double *z, const int N) {const int n = blockDim.x * blockIdx.x + threadIdx.x;if(n < N) {add1_device(x[n], y[n], &z[n]);}
}

3.3 code3 传引用

void __device__ add3_device(const double x, const double y, double &z) {z = x + y;
}void __global__ add1(const double *x, const double *y, double *z, const int N) {const int n = blockDim.x * blockIdx.x + threadIdx.x;if(n < N) {add1_device(x[n], y[n], z[n]);}
}

CUDA 错误检查

bool checkRuntime(cudaError_t e, const char* call, int line, const char *file){if (e != cudaSuccess) {INFOE("CUDA Runtime error %s # %s, code = %s [ %d ] in file %s:%d", call, cudaGetErrorString(e), cudaGetErrorName(e), e, file, line);return false;}return true;}

错误信息实例:

CUDA Runtime error cudaMemcpy(d_x, h_x, M, cudaMemcpyDeviceToDevice) # invalid argument, code = cudaErrorInvalidValue [ 1 ] in file D:\work_dir\experiment\LearningCUDA\src\add.cu:29

相关内容

热门资讯

激发暑期文旅消费新动能   内蒙古阿尔山,漫步国家森林公园,感受清凉夏日;甘肃敦煌,坐在鸣沙山上,听一场星空演唱会;江苏连云...
日股收跌0.19% 格隆汇7月11日|日经225指数收盘下跌76.68点,跌幅0.19%,报39569.68点。
遵义小伙跳水救人牺牲,年仅31... (转自:遵义晚报)7月8日下午6时,遵义市播州区出现感人一幕——尚嵇镇大坝社区全军坝组28岁的村民陈...
“小记者”走进河北省邯郸市食品... 中国质量新闻网讯 7月7日下午,河北省邯郸市食品药品检验中心(市食品药品安全科普基地)联合邯郸新闻传...
局地超42℃!下周河南持续高温... 大范围高温闷热天气强势来袭!12日起,西太平洋副热带高压将再次控制我省,受其影响,13日至20日我省...
毕市监处罚[2025]1021... 毕市监处罚[2025]1021号浏览:1次基本信息:处罚机关毕节市市场监督管理局处罚日期2025年0...
达势股份-达美乐中国以4D战略...   达势股份-达美乐中国("达美乐中国"或"达势股份"或"公司")(1405.HK)是达美乐比萨在中...
“向上取整”:变相“明抢” 转自:中国质量报□ 胡立彪近日,有媒体对快递行业称重计费问题进行测评,结果发现,被测评的8家快递企业...
以改革破局 盐城经开区多点突破... 中新网江苏新闻7月11日电(倪玲)改革是发展的根本动力,开放是进步的必由之路。近年来,盐城经开区将推...
中国代表强调国际刑事法院应依法... 转自:千龙网新华社联合国7月10日电 中国常驻联合国副代表孙磊10日在安理会审议国际刑事法院涉苏丹问...