【C++】CUDA期末复习指南下（详细）
温馨提示：这篇文章已超过411天没有更新，请注意相关的内容是否还可用！
🍎 博客主页：🌙@披星戴月的贾维斯
🍎 欢迎关注：👍点赞🍃收藏🔥留言
🍇系列专栏：🌙 C/C++专栏
🌙请不要相信胜利就像山坡上的蒲公英一样唾手可得，但是请相信，世界上总有一些美好值得我们全力以赴，哪怕粉身碎骨！🌙
🍉一起加油，去追寻、去成为更好的自己！
文章目录
- 前言
- 🍎1、cuda常考函数
- 🍎2、CUDA编程
- 🍇一个典型的CUDA程序的基本框架
- 🍇简单的CUDA加法
- 🍇获取计算机线程块的分配
- 🍇在GPU任意长度的矢量求和
- 🍇点积运算
- 🍇常量内存光线跟踪(使用共享内存)
- 🍇GPU使用一维纹理内存的热传导模拟计算
- 🍇统计直方图（普通版本）
- 🍇GPU原子递增操作统计直方图
- 🍎总结
提示：以下是本篇文章正文内容，下面案例可供参考
前言
上期CUDA期末复习指南我们主要讲了GPU的串行/并行以及一些背诵的知识点，这篇博客我们继续介绍cuda的函数以及cuda编程，常考的CUDA函数和编程题博主在这里为大家总结一下，希望对大家有所帮助。
🍎1、cuda常考函数
在上一节我们已经讲到了cudaFreeHost这个函数，是用来释放固定主机内存的。我们直接承接上一篇的内容
8. cudaSetDevice
cudaSetDevice 函数用于设置当前线程要使用的 CUDA 设备。它的函数原型为
cudaError_t cudaSetDevice ( int device );
其中， device 参数表示要使用的 CUDA 设备的编号，编号从 0 开始。
使用方法如下代码：
int main() { int device = 0; cudaSetDevice(device); return 0; }- cudaDeviceReset（这个函数用得较少）
cudaDeviceReset 是 CUDA 运行时 API 中的一个函数，用于重置当前设备的状态。该函数执行后，会清除当前设备上的所有运行时状态，并释放所有与该设备有关的资源。它的函数原型为：
cudaError_t cudaDeviceReset(void); // 该函数不接受任何参数
- cudaHostAlloc
cudaHostAlloc 是 CUDA 中用于在主机（Host）上分配内存的函数。与普通的 malloc 或 new不同，通过 cudaHostAlloc 分配的主机内存可以与设备（Device）之间进行零拷贝（Zerocopy）操作，这意味着可以在主机和设备之间非常高效地共享数据。它的函数原型为：
cudaError_t cudaHostAlloc( void **ptr, size_t size, unsigned int flags );
其中， ptr 为分配的内存的指针所存放的地址； size ：分配的内存的大小，以字节为单位；flags ：分配内存的标志，可以是以下任意组合：
cudaHostAllocDefault ：默认标志，表示按照系统的最佳准则分配主机内存。
cudaHostAllocPortable ：可以在不同 CUDA 设备之间交换的可移植内存。
cudaHostAllocMapped ：将主机内存映射到设备内存，以实现零拷贝操作。
cudaHostAllocWriteCombined ：对写缓冲区进行优化的内存，用于使用
cudaMemcpyAsync 在主机和设备之间高效地复制内存。
使用方法如下代码：
int main() { int n = 1024; float *data; HANDLE_ERROR( cudaHostAlloc((void**)&data, n * sizeof(float), cudaHostAllocMapped) ); // ... 使用 data 在主机和设备之间进行零拷贝操作 ... cudaFreeHost(data); return 0; }- cudaHostGetDevicePointer
cudaHostGetDevicePointer 函数是 CUDA 主机内存和设备内存之间数据传输的重要函数之一。
该函数的作用是将主机内存指针映射到设备内存地址空间中，并返回设备内存的指针。这样，就可
以在主机内存和设备内存之间直接传递数据，从而避免了在主机内存和设备内存之间复制数据的开
销，提高了程序运行效率。它的函数原型为：
cudaError_t cudaHostGetDevicePointer( void **pDevice, void *pHost, unsigned int flags // 此处需要被置为0 );
其中， pDevice 是指向设备内存指针的指针； pHost 是主机内存指针； flags 是标志位，用于控制映射的行为。
与 cudaHostAlloc 搭配使用方法如下：
int main() { float *h_ptr, *d_ptr; size_t size = N * sizeof(float); // 分配 CUDA 主机内存 HANDLE_ERROR( cudaHostAlloc(&h_ptr, size, cudaHostAllocDefault) ); // 在设备上分配内存 HANDLE_ERROR( cudaMalloc((void**)&d_ptr, size) ); // 映射设备内存 HANDLE_ERROR( cudaHostGetDevicePointer(&d_ptr, h_ptr, 0) ); // 将数据从主机内存复制到设备内存 HANDLE_ERROR( cudaMemcpy(d_ptr, h_ptr, size, cudaMemcpyHostToDevice) ); // 对设备上的数据进行操作 SomeKernel(d_ptr, N); // 将数据从设备内存复制到主机内存 HANDLE_ERROR( cudaMemcpy(h_ptr, d_ptr, size, cudaMemcpyDeviceToHost) ); // 释放内存 HANDLE_ERROR( cudaFree(d_ptr) ); HANDLE_ERROR( cudaFreeHost(h_ptr) ); return 0; }- cudaGetErrorString
cudaGetErrorString 是一个 CUDA 函数，它可以将 CUDA 错误码转换为可读字符串，方便开发者调试和查错。它的函数原型为：
cudaError_t cudaGetErrorString( cudaError_t error, const char **pStr);
其中， error 为要转换为字符串的 CUDA 错误码； **pStr 为指向指针的指针，用于存储转换后的字符串。用法如下：
使用方法如下：
cudaError_t err = cudaMalloc(&devPtr, size); if (err != cudaSuccess) { printf("CUDA error: %s\n", cudaGetErrorString(err)); }- cudaMemcpyToSymbol
cudaMemcpyToSymbol 将 count 个字节从 src 指向的内存复制到 symbol 指向的内存中，这个变量存放在设备的 全局内存或常量内存中。
使用方法如下：
__constant__ Sphere s[SPHERES]; // 在全局处定义 int main() { Sphere *temp_s = (Sphere*)malloc( sizeof(Sphere) * SPHERES); // 在局部动态拷贝 HANDLE_ERROR( cudaMemcpyToSymbol( s, temp_s, sizeof(Sphere) * SPHERES)); return 0; }- cudaEventCreate(重点)
cudaEventCreate 是一个CUDA运行时API函数，用于创建CUDA事件。CUDA事件可以被用于测量GPU程序的执行时间、在异步操作间同步数据等。它的函数原型为：
cudaError_t cudaEventCreate ( cudaEvent_t* event );
其中， event 是一个指向 cudaEvent_t 类型对象的指针。
使用方法如下：
int main() { cudaEvent_t start, stop; cudaEventcreate(&start); cudaEventcreate(&stop); }- cudaEventRecord
cudaEventRecord 是一个 CUDA 运行时函数，用于记录一个调用 cudaEventCreate 创建的CUDA 事件的时间点。通常情况下，它用作测量时间间隔或异步操作的同步。
- cudaEventSynchronize、
cudaEventSynchronize 函数用于等待一个事件完成，并阻塞当前主机线程，直到该事件完成为止。它接收一个事件作为参数，并会一直等待该事件完成，直到可以安全地使用该事件依赖的所有CUDA 内存和其他资源。
- cudaEventElapsedTime
cudaEventElapsedTime 函数用于计算两个事件记录的时间间隔，该函数返回两个事件之间的时间间隔（以毫秒为单位）。
- cudaEventDestroy
cudaEventDestroy 函数用于销毁一个 CUDA 事件对象，释放与之关联的资源。
使用方法如下：
int main() { cudaEvent_t start, stop; HANDLE_ERROR( cudaEventCreate(&start) ); HANDLE_ERROR( cudaEventCreate(&stop) ); // 记录开始时间点 HANDLE_ERROR( cudaEventRecord(start, 0) ); // 执行核函数 myKernel >(...); // 记录结束时间点 HANDLE_ERROR( cudaEventRecord(stop, 0) ); // 等待GPU操作完成 HANDLE_ERROR( cudaEventSynchronize(stop) ); float elapsedTime; // 计算两个时间点之间经过的时间间隔 HANDLE_ERROR( cudaEventElapsedTime(&elapsedTime, start, stop) ); printf("Execution time: %f ms\n", elapsedTime); // 销毁已经创建的事件对象 HANDLE_ERROR( cudaEventDestroy(start) ); HANDLE_ERROR( cudaEventDestroy(stop) ); return 0; }- atomicAdd
atomicAdd 是 CUDA 提供的一种原子操作函数，用于在 GPU 全局内存中进行原子加操作。它的、作用是确保在并发情况下，对该内存地址执行原子加操作，以避免数据争用的问题，使得并发访问的结果能够正确累加。
🍎2、CUDA编程
🍇一个典型的CUDA程序的基本框架
1、头文件包含 2、常量定义 3、C++自定义函数和CUDA核函数声明 4、int main (void) { 定义主机和设备变量/数组。 分配主机和设备内存 初始化主机中的数据 将某些数据从主机复制到设备 调用核函数在设备中计算 将计算结果从设备中拷贝回主机变量/数组。 释放主机/设备内存 } C++自定义函数和CUDA核函数实现（定义）🍇简单的CUDA加法
#include #include #include //定义global类型的函数 __global__ void add(int *const c, int const a, int const b) { *c = a + b; } int main () { int c; int *dev_c; //申请显存 cudaMalloc((void**)&dev_c, sizeof(int)); //调用核函数 addadd(dev_c, 2, 7); //把在dev_c的结果拷贝回cpu变量c中 cudaMemcpy(&c, dev_c, sizeof(int), cudaMemcpyDeviceToHost); printf("%d\n", c); //释放显存 cudaFree(dev_c); return 0; }🍇获取计算机线程块的分配
一个简短内核程序输出线程块、线程、线程束和线程全局标号到屏幕上，了解线程块的分配。
#include #include __global__ void what_is_my_id(unsigned int* const block, unsigned int* const thread, unsigned int* const warp, unsigned int* const calc_thread){ const unsigned int thread_idx = blockIdx.x * blockDim.x + threadIdx.x; //其中这个thread_idx是当前线程的坐标 block[thread_idx] = blockIdx.x;//当前线程位于的线程块 thread[thread_idx] = threadIdx.x;//当前线程的索引 warp[thread_idx] = threadIdx.x / warpSize;//当前线程的宽度 calc_thread[thread_idx] = thread_idx;//当前线程在整个网格中的索引 } #define ARRAY_SIZE 128 #define ARRAY_SIZE_IN_BYTES (sizeof(unsigned int) * (ARRAY_SIZE)) unsigned int cpu_block[ARRAY_SIZE]; unsigned int cpu_thread[ARRAY_SIZE]; unsigned int cpu_warp[ARRAY_SIZE]; unsigned int cpu_calc_thread[ARRAY_SIZE]; int main () { //确定调用核函数的多少 const unsigned int num_blocks = 2; const unsigned int num_threads = 64; //定义gpu线程块数组 unsigned int* gpu_block; unsigned int* gpu_thread; unsigned int* gpu_warp; unsigned int* gpu_calc_thread; //申请显存 cudaMalloc((void **)&gpu_block, ARRAY_SIZE_IN_BYTES); cudaMalloc((void **)&gpu_thread, ARRAY_SIZE_IN_BYTES); cudaMalloc((void **)&gpu_warp, ARRAY_SIZE_IN_BYTES); cudaMalloc((void **)&gpu_calc_thread, ARRAY_SIZE_IN_BYTES); //调用核函数 what_is_my_id(gpu_block, gpu_thread,gpu_warp,gpu_calc_thread); //再把显存内容拷贝回cpu的数组 cudaMemcpy(cpu_block, gpu_block,ARRAY_SIZE_IN_BYTES,cudaMemcpyDeviceToHost); cudaMemcpy(cpu_thread, gpu_thread,ARRAY_SIZE_IN_BYTES,cudaMemcpyDeviceToHost); cudaMemcpy(cpu_warp, gpu_warp,ARRAY_SIZE_IN_BYTES,cudaMemcpyDeviceToHost); cudaMemcpy(cpu_calc_thread, gpu_calc_thread,ARRAY_SIZE_IN_BYTES,cudaMemcpyDeviceToHost); //释放指针 cudaFree(gpu_block); cudaFree(gpu_thread); cudaFree(gpu_warp); cudaFree(gpu_calc_thread); //打印数据 for(int i = 0; u🍇在GPU任意长度的矢量求和
#include #include #include #define N (65536 * 128 * 10) __global__ void add(int* a, int *b, int *c) { int tid = blockIdx.x * blockDim.x + threadIdx.x; while(tid🍇点积运算
矢量的点积运算(Dot Product，也称为内积)。这个计算包含两个步骤。首先把两个输入矢量中相应的元素相乘，类似矩阵加法，不过使用的乘法操作。其次把计算出来的值相加起来得到一个标量输出值（单个数值）。
(x1,x2,x3,x4)·(y1,y2,y3,y4) = x1y1+x2y2+x3y3+x4y4
先像矢量加法那样，每个线程将两个相应的元素相乘，每个线程保存计算的乘积的求和。
代码示例：
#include"../common/book.h" #define imin (a,b) (a
🍇常量内存光线跟踪(使用共享内存)
以光线跟踪为例，从三维对象场景中生产二维图像的一种方式。需要判断图像中的每个像素的颜色和强度。
一组包含球状物体的场景，从每个像素发射一道光线，跟踪这些光线会命中哪些球面，当一道光线穿过多个球面时，只有最接近的球面才有机会被看到。
对球面进行数据结构建模，包含了球面中心点(x,y,z)，半径radius，颜色值(r,g,b)
代码示例：
#include "cuda_runtime.h" #include "device_launch_parameters.h" #include "H:\cuda_by_example\common\book.h" #include "H:\cuda_by_example\common\cpu_bitmap.h" #include "device_functions.h" #include #define DIM 1024 #define rnd( x ) (x * rand() / RAND_MAX) #define INF 2e10f struct Sphere { float r,b,g; float radius; float x,y,z; __device__ float hit( float ox, float oy, float *n ) { float dx = ox - x; float dy = oy - y; if (dx*dx + dy*dy bitmap; //由于tex是全局并且是有界的，因此我们必须通过一个标志来来选择每次迭代中哪个是输入/输出 volatile bool dstOut = true; for (int i = 0; i dev_inSrc; out = d->dev_outSrc; } else { in = d->dev_outSrc; out = d->dev_inSrc; } copy_const_kernal > (in); blend_kernal > (out,dstOut); dstOut = !dstOut; } //将float数值映射成颜色，以便用图像显示它 float_to_color >(d->output_bitmap, d->dev_inSrc); cudaMemcpy(bitmap->get_ptr(), d->output_bitmap, bitmap->image_size(), cudaMemcpyDeviceToHost); cudaEventRecord(d->stop, 0); cudaEventSynchronize(d->stop); float elapsedTime; cudaEventElapsedTime(&elapsedTime, d->start, d->stop); d->totalTime += elapsedTime; ++d->frames; printf("Average Time per frame:%3.1fms\n", d->totalTime / d->frames); } void anim_exit(DataBlock *d) { cudaUnbindTexture(texIn); cudaUnbindTexture(texOut); cudaUnbindTexture(texConstSrc); cudaFree(d->dev_constSrc); cudaFree(d->dev_inSrc); cudaFree(d->dev_outSrc); cudaEventDestroy(d->start); cudaEventDestroy(d->stop); }🍇统计直方图（普通版本）
#include "../common/book.h" #define Size (100*1024*1024) int main() { unsigned char *buffer = (unsigned char*)big_random_block(Size); clock_t start, stop; start = clock(); unsigned int histo[256]; for(int i = 0; i🍇GPU原子递增操作统计直方图
代码示例：
#include "../common/book.h" #define Size (100*1024*1024) __global__ void histo_kernel(unsigned char *buffer,long size,unsigned int *histo){ int i = threadIdx.x + blockIdx.x * blockDim.x; int stride = blockDim.x * gridDim.x; while(i🍎总结
本文总共总结多个常考的函数和编程题，希望对大家学习和复习CUDA有所帮助，感谢大家的支持，一起进步！
- cudaDeviceReset（这个函数用得较少）
