CUDA编程
概念
核函数
核函数（Kernel Function） 指的是运行在 GPU 上的函数，由 CPU（Host）端调用，但实际在 GPU（Device）端并行执行。
核函数的声明需要限定符 __global__ ，例如：
__global__ void add(int *a, int *b, int *c, int N) {int idx = blockIdx.x * blockDim.x + threadIdx.x; // 全局线程索引if (idx < N) {c[idx] = a[idx] + b[idx];}
}
核函数的调用需要用<<...>>运算符来指定线程的层次结构：
dim3 gridDim(x, y, z);//目前，grid还不支持三维，z始终等于 1。
dim3 blockDim(x, y, z);
add<<<gridDim, blockDim>>>(d_a, d_b, d_c, N);
其中，第一个参数是Grid的维度，即Grid由多少个Block组成；第二个参数是Block的维度，即每个Block里的线程数。注意，这里的gridDim,blockDim是最多可以支持三维数据的。
在核函数调用的尖括号中还可以带有一个流参数，此时核函数调用将是异步的。
kernel<<<N/256,256,0, stream>>>( dev_a, dev_b, dev_c );
Block中的线程是有最大限制的
blockDim.x × blockDim.y × blockDim.z ≤ 1024（cudaDeviceProp中的maxThreadsPerBlock是有描述的）
- blockDim.x ≤ 1024
- blockDim.y ≤ 1024
- blockDim.z ≤ 64
Grid中的Block也是有限制的
- gridDim.x ≤ 2³¹-1
- gridDim.y ≤ 65535
- gridDim.z = 1 （不支持超过二维）
限定符
对于限定符，有以下常用的。在声明函数时，他们被放在函数返回类型的前面。对于修饰变量的限定符，他们被放在类型前。
global
- 声明核函数（kernel function）
- 调用端：Host（CPU）
- 执行端：Device（GPU）
- 返回值：必须是 void
device
- 声明设备函数（只能在 GPU 上运行）
- 调用端：只能由 GPU（Device）里的函数调用
- 执行端：Device（GPU）
- 返回值：任意
host
- 声明主机函数（只能在 CPU 上运行）
- 调用端：Host（CPU）
- 执行端：Host（CPU）
- 默认情况下，所有普通函数就是 host，写不写效果一样。
__host__device__
- 声明主机函数（只能在 CPU 上运行）
- 调用端：Host（CPU）
- 执行端：Host（CPU）
- 默认情况下，所有普通函数就是 host，写不写效果一样。
constant（用于修饰变量）
- 表示定义在 GPU 常量内存里的变量，所有线程可共享，且只读。
- 注意！这是用来修饰变量的！
shared（用于修饰变量）
- 表示定义在 GPU 共享内存（SM 内部共享内存）里的变量，线程块内的线程可共享读写。
Grid&Block&Thread与关系
Kernel
- 写在 GPU 上运行的函数，用 global 声明。
- 由 CPU 启动，但实际在 GPU 上并行执行。
- 一次 调用核函数 就会生成大量线程并行运行同一段代码。但是一次调用只启动一个Grid
Grid
- Grid 是由很多 Block 组成的整体。
- Grid 也可以是一维、二维或三维的结构。
- 不同 Block 之间一般不能直接通信（除非通过全局内存）。
Block
- 一个 Block 里包含很多 Thread。
- Block 内部的线程可以通过 共享内存（shared） 互相通信，也可以通过 同步（__syncthreads()） 协调执行。
- Block 可以是一维、二维或三维的结构（最多 3D）。
Thread
- GPU 上最基本的执行单元。
- 每个线程都有自己独立的 寄存器、局部变量。
- 所有线程都会运行核函数里的代码，但通常处理不同的数据。
共享内存
Block内共享，Block之间不共享
编写代码时，你可以将 _ _s h a r e_ _ 添加到变量声明中，这将使这个变量驻留在共享内存中。
对于在GPU上启动的每个Block，CUDAC编译器都将创建该变量的一个副本。线程块中的每个线程都共享这块内存，但线程却无法看到也不能修改其他线程块的变量副本。
而且，共享内存缓冲区驻留在物理GPU上，而不是驻留在GPU之外的系统内存中。因此，在访问共享内存时的延迟要远远低于访问普通缓冲区的延迟，使得共享内存像每个线程块的高速缓存或者中间结果暂存器那样高效。
常量内存
常量内存 是 GPU 上的一块只读存储区域，专门用于存储程序运行过程中不变的数据。是 GPU 全局内存的一部分，但有专门的 缓存（constant cache），访问速度很快。所有线程共享同一份常量内存。一般只有 64KB。
使用 __constant__ 修饰。常量内存只能由 Host（CPU）写入，不能由 Device（GPU）写入。
使用cudaMemcpyToSymbol( )写入常量内存。例如
float h_array[256];
// 初始化
for (int i = 0; i < 256; i++) h_array[i] = i * 1.0f;
// 拷贝到常量内存
cudaMemcpyToSymbol(constArray, h_array, 256 * sizeof(float));
为什么常量内存快？
- 对常量内存的单次读操作可以广播到其他的“邻近(Nearby)”线程，这将节约读取操作。
- 这里所谓的邻近，其实是半个线程束(warp)。一个warp有 32 个thread，所以一次广播节约 16次读取操作
- 常量内存的数据将缓存起来，因此对相同地址的连续读操作将不会产生额外的内存通信量。
当处理常量内存时,NVIDIA硬件将把单次内存读取操作广播到每个半线程束(Half-Warp)在半线程束中包含了16个线程，即线程束中线程数量的一半。如果在半线程束中的每个线程都从常量内存的相同地址上读取数据，那么GPU只会产生一次读取请求并在随后将数据广播到每个线程。如果从常量内存中读取大量的数据，那么这种方式产生的内存流量只是使用全局内存时的1/16(大约6%)。
但在读取常量内存时，所节约的并不仅限于减少了94%的带宽。由于这块内存的内容是不会发生变化的，因此硬件将主动把这个常量数据缓存在GPU上。在第一次从常量内存的某个地址上读取后，当其他半线程束请求同一个地址时，那么将命中缓存，这同样减少了额外的内存流量。
为什么常量内存不是一定会快？
当使用常量内存时，也可能对性能产生负面影响。半线程束广播功能实际上是一把双刃剑。虽然当所有16个线程都读取相同地址时，这个功能可以极大地提升性能，但当所有16个线程分别读取不同的地址时，它实际上会降低性能。
只有当16个线程每次都只需要相同的读取请求时，才值得将这个读取操作广播到16个线程。然而，如果半线程束中的所有16个线程需要访问常量内存中不同的数据，那么这个16次不同的读取操作会被串行化，从而需要16倍的时间来发出请求。但如果从全局内存中读取，那么这些请求会同时发出。在这种情况中，从常量内存读取就慢于从全局内存中读取。
流
流（cudaStream_t） 是 CUDA 的执行队列。核函数（kernel）和内存拷贝等操作都会在流里排队执行。默认情况下，0 表示 默认流。在默认流上发出的操作会顺序执行，且会与其他流的操作有一些同步规则。
流的并行与其他并行不太一样，它是如同CPU多线程应用程序中的任务并行性。任务并行性是指并行执行两个或多个不同的任务，而不是在大量数据上执行同一个任务。
创建&销毁流
cudaStream_t stream;//声明
cudaStreamCreate(&stream);//注册
cudaStreamSynchronize(stream);//等待流完成
cudaStreamDestory(stream)//销毁流
设备重叠
就是 数据传输（CPU ↔ GPU）和 GPU 上的 kernel 计算可以并行执行
要让数据传输和 kernel 真正并行，需要满足以下条件：
- 硬件支持：deviceOverlap == 1。
- 使用流 (CUDA Streams)：
- 默认流（stream 0）是串行的。
- 必须创建多个流，才能并行拷贝和计算。
- 异步拷贝：
- 用 cudaMemcpyAsync（异步拷贝），而不是 cudaMemcpy（同步）。
- 而且异步拷贝的 host 内存必须是 pinned memory（通过 cudaHostAlloc 分配），否则也会退化成同步。
现在几乎所有设备都支持设备重叠了，甚至有大部分已经有 2 个copy engine了。可以通过以下代码检查
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, deviceID);
printf("deviceOverlap = %d\n", prop.deviceOverlap);
printf("asyncEngineCount = %d\n", prop.asyncEngineCount);
//deviceOverlap = 1 → 支持拷贝和计算重叠。
//asyncEngineCount = 2 → H→D、D→H、Compute 可以三向并行。
硬件上的调度机制
硬件的执行模式和直观上的“流”模式是不同的，硬件是将核函数、数据复制分成了两个队列单元来进行的。
可以想象成，硬件上对内存复制引擎、核函数执行引擎各自做了一个队列。两个队列将会把流上的操作各自复制进来进行排队，并进行相应的阻塞。
通过这三张图片，我们可以更直观的感受到硬件工作模式上的差别。
因此，在使用流的时候，将操作放入流队列的时候应该遵守宽度优先模式，而不是深度优先模式。
页锁定内存
C库函数malloc()将分配标准的，可分页的(Pagable)主机内存 ，而cudaHostAlloc()将分配页锁定的主机内存。页锁定内存也称为固定内存(Pinned Memory)或者不可分页内存，它有一个重要的属性:操作系统将不会对这块内存分页并交换到磁盘上，从而确保了该内存始终驻留在物理内存中。因此，操作系统能够安全地使某个应用程序访问该内存的物理地址，因为这块内存将不会被破坏或者重新定位。
为什么使用页锁定内存？DMA中，复制时无需CPU介入，但CPU很可能在DMA的执行过程中将目标内存交换到磁盘上，或者通过更新操作系统的可分页表来重新定位目标内存的物理地址。CPU可能会移动可分页的数据，这就可能对DMA操作造成延迟。当使用可分页内存进行复制时，CUDA驱动程序仍然会通过DAM把数据传输给GPU。因此，复制操作将执行两遍，第一遍从可分页内存复制到一块“临时的”页锁定内存，然后再从这个页锁定内存复制到GPU上。因此，每当从可分页内存中执行复制操作时，复制速度将受限于PCIE传输速度和系统前端总线速度相对较低的一方。
但同时需要注意到，并不应该将所有的内存都声明为页锁定内存。使用固定内存时，会失去虚拟内存的所有功能，系统将更快的耗尽内存。
建议，仅对cudaMemcpy()调用中的源内存或者目标内存，才使用页锁定内存，并且在不再需要使用它们时立即释放，而不是等到应用程序关闭时才释放。
纹理内存
纹理内存（Texture Memory） 是一种只读的、经过特殊优化的内存，主要用于处理有空间局部性的数据访问。纹理缓存是专门为那些在内存访问模式中存在大量空间局部性(Spatial Locality)的图形应用程序而设计的。
纹理内存其实还是存放在 显存（device memory） 里，不是单独的硬件。不过，访问它时会经过 纹理缓存（texture cache），而这个 cache 对二维/三维访问模式优化比较好。
特点
-
只读
-
内核（kernel）中只能读取纹理内存，不能写。
-
主机端（CPU）可以把数据拷贝到绑定的显存区域。
-
-
空间局部性优化
-
如果相邻线程访问相邻的数据点（尤其是二维/三维数据），纹理缓存会大大减少显存带宽消耗。
-
这点和普通的 shared 或 L1 cache 不同，它专门为这种「采样」型访问优化。
-
-
插值与边界处理（图形学特性）
-
在图形渲染中，纹理内存支持线性插值（linear interpolation）、边界模式（wrap/clamp） 等功能。
-
在通用 CUDA 里，如果启用这些特性，可以在硬件层面获得插值结果。
-
cudaArray
cudaArray 是 CUDA 里 专门给纹理（texture）和表面（surface）内存用的数据存储格式。是一种 特殊的存储对象，通常存放在显存里。它的存储方式和普通线性内存不一样，设计上是 针对图像/纹理访问优化的
必须要用 cudaArray 的场景：
- 使用 二维/三维纹理内存（如 tex2D, tex3D）。
- 使用 线性插值（filterMode=Linear）采样。
- 使用 表面内存（surface memory） 进行写操作。
线性插值：线性插值采样就是 CUDA 纹理单元在采样非整数坐标时，会根据相邻像素/体素的值按比例平滑计算出中间值，而不是简单取整。
创建与释放
分配 cudaArray 要用 cudaMallocArray（或 cudaMalloc3DArray）：
cudaArray* cuArray;
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
cudaMallocArray(&cuArray, &channelDesc, width, height);
释放时用
cudaFreeArray(cuArray);
数据拷贝
因为 cudaArray 不是普通的指针内存，不能直接 cudaMemcpy，而是要用专门的函数：
cudaMemcpy2DToArray(cuArray, 0, 0, h_data,pitch, width*sizeof(float), height,cudaMemcpyHostToDevice);
使用方法
三种资源描述符
cudaChannelFormatDesc 数据格式描述符
它描述了「存储单元里一条数据」的格式，比如一个像素里有几个分量、每个分量多少位、是 int 还是 float。
struct cudaChannelFormatDesc {int x; // 通道 X 的位宽（例如 R 通道）int y; // 通道 Y 的位宽（例如 G 通道）int z; // 通道 Z 的位宽（例如 B 通道）int w; // 通道 W 的位宽（例如 A 通道）enum cudaChannelFormatKind f; // 数据类型（float / signed / unsigned 等）
};
enum cudaChannelFormatKind {cudaChannelFormatKindSigned,cudaChannelFormatKindUnsigned,cudaChannelFormatKindFloat,cudaChannelFormatKindNone
};cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>(); // 单通道 float
cudaResourceDesc资源描述符
告诉 CUDA 数据存在哪里，比如在 cudaArray、线性内存、还是 MipMap。
struct cudaResourceDesc {cudaResourceType resType; // 资源类型：Array / Linear / Pitch2D / MipmappedArrayunion {struct {cudaArray_t array; // 如果是 cudaArray} array;struct {void* devPtr; // 线性内存指针struct cudaChannelFormatDesc desc; // 数据格式size_t sizeInBytes; // 总字节数} linear;struct {void* devPtr; // pitch 内存指针struct cudaChannelFormatDesc desc; // 数据格式size_t width;size_t height;size_t pitch; // 每行字节数} pitch2D;struct {cudaMipmappedArray_t mipmap; // MipMap 数据} mipmap;};
};
enum cudaResourceType {cudaResourceTypeArray = 0,cudaResourceTypeMipmappedArray,cudaResourceTypeLinear,cudaResourceTypePitch2D
};
cudaTextureDesc采样方式描述符
告诉 CUDA 怎么采样这个资源，包括：
- 地址模式（超出边界怎么办？Clamp/Wrap/Border）
- 过滤模式（最近邻 / 线性插值）
- 坐标归一化（是用 0~1 还是直接用像素坐标）
struct cudaTextureDesc {cudaTextureAddressMode addressMode[3]; // 边界模式（x,y,z）cudaTextureFilterMode filterMode; // 采样模式（Point / Linear）cudaTextureReadMode readMode; // 读取模式（按元素类型 / 转 float）unsigned int normalizedCoords; // 坐标是否归一化（0=整数坐标，1=归一化坐标）float maxAnisotropy; // 各向异性过滤（可选）cudaTextureFilterMode mipmapFilterMode; // Mipmap 过滤模式float mipmapLevelBias; // Mipmap 偏移unsigned int mipmapLevelClamp[2]; // Mipmap 级别范围
};
enum cudaTextureAddressMode {cudaAddressModeWrap = 0,cudaAddressModeClamp = 1,cudaAddressModeMirror = 2,cudaAddressModeBorder = 3
};enum cudaTextureFilterMode {cudaFilterModePoint = 0, // 最近邻cudaFilterModeLinear = 1 // 线性插值
};enum cudaTextureReadMode {cudaReadModeElementType = 0, // 原始数据类型cudaReadModeNormalizedFloat = 1 // 转为归一化 float
};
老式API(cuda5.0-)
首先需要将数据声明为texture类型的引用。在内核上能读，但是不能写。
texture<flot> texIn;
之后使用cudaBindTexture()将这些变量绑定到内存缓冲区。这主要是告诉cuda我们希望将指定的缓冲区作为纹理来使用，并且希望将纹理引用作为纹理的名字。
cudaBindTexture(NULL,texIn,data.dev_insrc,imageSize);
之后就可以用texlDfetch()来读这些内存了。
float t = tex1Dfetch(texIn,offset) //有点类似于textIn[offset]，但是只能读不能写
最后，使用完毕后，可以通过cudaUnbindTexture()来解除绑定
cudaUnbindTexture(textIn);
当然，也可以用二维的纹理内存。
texture<float,2> texIn;
cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
/*
struct cudaChannelFormatDesc{int x;//通道x的位宽int y;//通道y的位宽int z;//通道z的位宽int w;//通道w的位宽enum cudaChannelFormatKind f;//数据类型(signed/unsigned/float)
}
*/
cudaBindTexture2D(NULL,texIn,data.dev_insrc,desc,DIM,dIM,sizeof(float)*DIM);
float c = tex2D(texIn,x,y);当使用tex2D()时，我们不再需要担心发生溢出问题。如果x或y小于0，那么tex2D()而且，将返回0处的值。同理，如果某个值大于宽度，那么tex2D()将返回位于宽度处的值。
新API(cuda5.0+)
新式 API 不再使用 texture<> 全局变量和 cudaBindTexture，而是通过 纹理对象（cudaTextureObject_t） 来管理，这样可以更灵活地传参和管理多个纹理。
首先需要创建通道描述符
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
然后定义资源描述符（即数据来源）
cudaResourceDesc resDesc = {};
resDesc.resType = cudaResourceTypeLinear; // 一维线性内存
resDesc.res.linear.devPtr = d_data; // 指向设备内存
resDesc.res.linear.desc = channelDesc; // 数据格式
resDesc.res.linear.sizeInBytes = N * sizeof(float);
resType 指定资源类型，可以是：
- cudaResourceTypeLinear（线性内存）
- cudaResourceTypePitch2D（二维内存）
- cudaResourceTypeArray（CUDA array）
然后定义纹理描述符（采样方式）
cudaTextureDesc texDesc = {};
texDesc.addressMode[0] = cudaAddressModeClamp; // 越界时取边界值
texDesc.filterMode = cudaFilterModePoint; // 最近点采样（不插值）
texDesc.readMode = cudaReadModeElementType;// 按元素类型读取
texDesc.normalizedCoords = 0; // 使用非归一化坐标
之后创建纹理对象
cudaTextureObject_t texObj = 0;
cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL);
后续便可以在内核中读取纹理
__global__ void texKernel(cudaTextureObject_t texObj, float *out, int N) {int i = threadIdx.x + blockIdx.x * blockDim.x;if (i < N) {out[i] = tex1Dfetch<float>(texObj, i); // 按索引读取}
}
用后销毁
cudaDestroyTextureObject(texObj);
当然，也支持2D的，下面是一段2D纹理内存的示例
#include <cuda_runtime.h>
#include <iostream>__global__ void kernel(cudaTextureObject_t texObj, int width, int height) {int x = threadIdx.x + blockIdx.x * blockDim.x;int y = threadIdx.y + blockIdx.y * blockDim.y;if (x < width && y < height) {// 从纹理对象里取数据float val = tex2D<float>(texObj, x, y);printf("Thread (%d,%d): val = %f\n", x, y, val);}
}int main() {const int width = 4, height = 4;float h_data[width * height];// 初始化数据for (int i = 0; i < width * height; i++) {h_data[i] = (float)i;}// 1. 分配 CUDA 数组cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();cudaArray* cuArray;cudaMallocArray(&cuArray, &channelDesc, width, height);// 2. 拷贝数据到 CUDA 数组cudaMemcpy2DToArray(cuArray, 0, 0, h_data,width * sizeof(float), width * sizeof(float), height,cudaMemcpyHostToDevice);// 3. 填写资源描述符cudaResourceDesc resDesc = {};resDesc.resType = cudaResourceTypeArray;resDesc.res.array.array = cuArray;// 4. 填写纹理描述符cudaTextureDesc texDesc = {};texDesc.addressMode[0] = cudaAddressModeClamp; // x方向边界处理texDesc.addressMode[1] = cudaAddressModeClamp; // y方向边界处理texDesc.filterMode = cudaFilterModePoint; // 最近邻采样texDesc.readMode = cudaReadModeElementType;// 直接读元素texDesc.normalizedCoords = 0; // 使用整数坐标// 5. 创建纹理对象cudaTextureObject_t texObj = 0;cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL);// 6. 启动 kerneldim3 block(2, 2);dim3 grid((width + block.x - 1) / block.x,(height + block.y - 1) / block.y);kernel<<<grid, block>>>(texObj, width, height);cudaDeviceSynchronize();// 7. 销毁纹理对象，释放资源cudaDestroyTextureObject(texObj);cudaFreeArray(cuArray);return 0;
}
零拷贝内存
零拷贝内存（Zero-Copy Memory）是 CUDA 里一种特殊的内存分配方式。它让 GPU 内核直接访问主机内存，从而避免显式的数据拷贝。
int *h_ptr; // 主机指针
int *d_ptr; // 设备指针
// 分配零拷贝内存（页锁定+映射）
cudaHostAlloc((void**)&h_ptr, N * sizeof(int), cudaHostAllocMapped);
// 获取对应的设备指针
cudaHostGetDevicePointer((void**)&d_ptr, h_ptr, 0);
//这两个指针指向同一块内存，只是在 CPU 和 GPU 上的“视角”不同。
优点
- 不需要 cudaMemcpy()，简化代码
- GPU 和 CPU 共享一份数据，避免额外显存占用
- 对小数据量、通信频繁的场景很合适
- 对于集成GPU，使用零拷贝内存通常都会带来性能提升，因为内存在物理上与主机是共享的。将缓冲区声明为零拷贝内存的唯一作用就是避免不必要的数据复制
- 当输入内存和输出内存都只能使用一次时，且数据很小（最大就几MB）那么在独立GPU上使用零拷贝内存将带来性能提升。
缺点
- 访问速度慢（因为 GPU 访问的是 PCIe 上的主机内存，带宽和延迟比显存差很多）
- 不适合大规模计算（显存访问远比零拷贝快）
多GPU编程
每个GPU都需要一个不同的CPU线程来控制，可以通过cudaSetDevice()来指定当前线程希望使用的GPU。一旦在某个线程上设置了这个设备，那么将不能再次调用cudaSetDevice()，即便传递的是相同的设备标识符。
cuda_by_example教材中的CUTThread、start_thread()等均已被弃用（cuda4.0）。
对于CPU线程来说，前面所提到的“固定”内存只是对于单个CPU线程来说的。当使用其他线程的时，将仍然被视为标准、可分页的内存。但是我们可以将固定内存分配为可移动的，这意味着可以在主机线程之间移动这块内存，并且每个线程都将其视为固定内存。要达到这个目的，需要使用cudaHostAlloc()来分配内存，并且在调用时使用一个新的标志:cudaHostAllocPortable。
常用基础函数与数据结构
内存显存相关
cudaMalloc()
cudaError_t cudaMalloc(void** devptr,size_t size)
通过这种方法，可以在cuda设备上申请一块普通内存，并返回指针devptr。
例如，我们可以申请一个int类型的
int* dev_c;
cudaMalloc( (void**) &dev_c,sizeof(int));
这里void**是指针的指针，当我们传入&dev_c的时候，实际上是传了一个int**,然后进行了一次强制转换。
注意，请不要在主机代码（即在CPU上的代码）解引用。可以在主机代码上作为参数传递、进行算术运算、转换类型，但是不能用它读写内存。
cudaFree()
cudaError_t cudaFree(void* devPtr);
对于cudaMalloc申请的内存，需要通过cudaFree来释放。
例如，我们将释放掉之前申请的dev_c
int* dev_c;
cudaMalloc( (void**) &dev_c,sizeof(int));
cudaFree(dev_c)
cudaMemcpy()
用来在CPU主机内存和GPU显存之间拷贝数据
cudaError_t cudaMemcpy(void* dst, // 目标地址const void* src, // 源地址size_t count, // 拷贝的字节数cudaMemcpyKind kind // 拷贝的方向
);
/*
kind常见的取值:
cudaMemcpyHostToDevice
cudaMemcpyDeviceToHost
cudaMemcpyHostToHost
cudaMemcpyDeviceToDevice
*/
注意：
- 这是一项同步操作，调度会一直阻塞，直到拷贝完成
- 如果需要异步拷贝，参考cudaMemcpyAsync
- 如果Host内存是页锁内存，那拷贝更快
- 这会把内存复制到全局内存。如果是常量的话，参考cudaMemcpyToSymbol
cudaMemcpyToSymbol()
cudaMemcpyToSymbol() 用于 将主机（Host）内存的数据复制到 GPU 常量内存中的符号变量，常用于传输只读数据给核函数。
cudaError_t cudaMemcpyToSymbol(const void* symbol, // 目标符号（常量变量）const void* src, // 源数据地址（Host 内存或 Device 内存）size_t count, // 要拷贝的数据字节数size_t offset = 0, // 偏移量（可选）cudaMemcpyKind kind = cudaMemcpyHostToDevice // 拷贝类型
);
如果只想更新常量内存的一部分，可以使用 offset：
cudaMemcpyToSymbol(constArray, h_array, 128 * sizeof(float), 128 * sizeof(float));
这表示从 h_array 拷贝 128 个 float 数据到常量内存 constArray 的偏移位置（第 128 个 float 开始）。也就是拷贝的数据是 h_array[0]~h_array[127]，拷贝的位置是 constArray[128]~constArray[255]
注意：
-
必须使用 constant 修饰的变量
普通变量不能用 cudaMemcpyToSymbol()
-
只能从 Host 写入
常量内存不能由核函数直接写入
-
大小限制
常量内存通常 ≤ 64KB，不要超出
-
效率考虑
常量内存有专用缓存，访问速度非常快，但如果所有线程访问不同地址，效率会下降
cudaMemset()
这个函数的原型与标准C函数memset()的原型是相似的，并且这两个函数的行为也基本相同。二者的差异在于，cudaMemset()将返回一个错误码，而C库函数memset()则不是。这个错误码将告诉调用者在设置GPU内存时发生的错误。除了返回错误码外，还有一个不同之处就是，cudaMemset()是在GPU内存上执行，而memset()是在主机内存上运行。
cudaError_t cudaMemset(void *devPtr, int value, size_t count);
//设备内存指针（必须是用 cudaMalloc 分配的地址）。
//要填充的值（按字节填充）。
//要填充的字节数。
cudaHostAlloc()
用于在 主机（CPU）端分配页锁定内存（pinned memory / page-locked memory）。
这种内存的特点是：
- 驻留在物理内存里，不会被操作系统换到磁盘。
- GPU 可以更高效地通过 DMA（直接内存访问）和它进行数据传输。
- 可以在某些模式下让 GPU 直接访问这块内存（零拷贝）。
因此，cudaHostAlloc 常用于 提高 CPU-GPU 数据传输带宽，或者实现 零拷贝内存访问。
cudaError_t cudaHostAlloc(void** ptr, size_t size, unsigned int flags);
//ptr：指向返回的主机内存指针。
//size：要分配的字节数。
//flags：分配模式标志,组合可以用|或者按位结合
//返回值：CUDA 错误码（成功时返回 cudaSuccess）
/*
flags:
cudaHostAllocDefault
默认行为，分配页锁定内存。
cudaHostAllocPortable
分配的内存在所有 CUDA 上下文中可见。
cudaHostAllocMapped
分配的内存可映射到设备地址空间（可实现零拷贝）。
cudaHostAllocWriteCombined
分配写合并内存（写入速度快，读取速度慢），适合 CPU 只写、GPU 只读的场景。
*/
注意事项：
- 内存有限：pinned memory 会占用物理内存，过多分配会影响系统性能。
- 释放方式：必须用 cudaFreeHost 释放，而不是 free()。
- 零拷贝用法：配合 cudaHostAllocMapped 和 cudaHostGetDevicePointer 可以让 GPU 直接访问主机内存
cudaHostGetDevicePointer()
用于获取页锁定主机内存的 GPU 端指针，主要配合 零拷贝内存（Zero-Copy Memory） 使用。
cudaError_t cudaHostGetDevicePointer(void **pDevice,//pDevice：返回的 GPU 端指针（设备指针）。void *pHost,//已分配的主机端页锁定内存指针（由 cudaHostAlloc 分配，并且带有 cudaHostAllocMapped 标志）。unsigned int flags//一般设为 0，暂时没有常用选项。
);
cudaFreeHost()
用于 释放由 cudaHostAlloc 分配的页锁定内存。
cudaError_t cudaFreeHost(void* ptr);
cudaFreeHost 就是 回收 cudaHostAlloc 分配的 pinned memory 的函数，必须配套使用，避免内存泄漏。
设备查询相关
cudaDeviceProp
struct cudaDeviceProp {char name[256]; // 设备名称，比如 "NVIDIA A100"size_t totalGlobalMem; // 全局内存总量（字节）size_t sharedMemPerBlock; // 每个线程块可用的共享内存（字节）int regsPerBlock; // 每个线程块可用的寄存器数量int warpSize; // 一个 warp 的线程数（一般是 32）size_t memPitch; // 内存分配对齐要求int maxThreadsPerBlock; // 一个线程块的最大线程数int maxThreadsDim[3]; // 每个维度线程数上限int maxGridSize[3]; // 每个维度网格数上限size_t totalConstMem; // 常量内存大小（字节）int major; // 计算能力 (Compute Capability) 主版本int minor; // 计算能力 (Compute Capability) 次版本int clockRate; // 核心时钟频率 (kHz)size_t textureAlignment; // 纹理对齐要求int deviceOverlap; // 是否支持异步拷贝和核函数并发int multiProcessorCount; // SM（Streaming Multiprocessor）数量int kernelExecTimeoutEnabled; // 是否支持内核执行超时// ... 还有很多字段，这里只列一些常用的
};
cudaGetDeviceCount()
是 CUDA 的一个运行时 API，用来获取当前系统里 可用 CUDA 设备（GPU） 的数量。
cudaError_t cudaGetDeviceCount(int* count);
cudaGetDeviceProperties()
是 CUDA 提供的一个 API，用来获取 指定 GPU 设备的详细属性，结果存储在 cudaDeviceProp 结构体里。
cudaError_t cudaGetDeviceProperties(cudaDeviceProp* prop, int device);
调用示例
int deviceCount;cudaGetDeviceCount(&deviceCount);for (int i = 0; i < deviceCount; i++) {cudaDeviceProp prop;cudaGetDeviceProperties(&prop, i);cout << "GPU " << i << ": " << prop.name << endl;}
cudaChooseDevice()
是 CUDA 的一个 API，用来根据指定的条件选择最合适的 GPU 设备，适合在多 GPU 系统上自动挑选。
cudaError_t cudaChooseDevice(int* device, const cudaDeviceProp* prop);
我们将给出一个在多GPU系统上筛选合适CPU的例子：
#include <cuda_runtime.h>
#include <iostream>
using namespace std;int main() {int dev;cudaDeviceProp prop;// 清空结构体memset(&prop, 0, sizeof(cudaDeviceProp));// 设置需求，比如至少需要 8.0 以上的计算能力prop.major = 8;prop.minor = 0;// 根据需求选择最合适的 GPUcudaError_t error = cudaChooseDevice(&dev, &prop);if (error != cudaSuccess) {cout << "没有找到符合条件的 GPU: " << cudaGetErrorString(error) << endl;return -1;}cout << "选择的 GPU 编号是 " << dev << endl;// 设置当前设备cudaSetDevice(dev);// 获取它的属性cudaGetDeviceProperties(&prop, dev);cout << "GPU " << dev << " 名称: " << prop.name << endl;
}
cudaMemcpyAsync()
是 异步的数据传输函数，用于在 主机 (Host) 和 设备 (Device) 之间拷贝数据，或者设备内存之间拷贝。
cudaMemcpy()的行为类似于C库函数memcpy()。尤其是，这个函数将以同步方式执行，这意味着，当函数返回时，复制操作就已经完成，并且在输出缓冲区中包含了复制进去的内容。
在调用cudaMemcpyAsync()时，只是放置一个请求，表示在流中执行一次内存复制操作，这个流是通过参数stream来指定的。当函数返回时，我们无法确保复制操作是否已经启动，更无法保证它是否已经结束。我们能够得到的保证是，复制操作肯定会当下一个被放入流中的操作之前执行。任何传递给cudaMemcpyAsync()的主机内存指针都必须已经通过cudaHostAlloc()分配好内存。也就是，你只能以异步方式对页锁定内存进行复制操作。
cudaError_t cudaMemcpyAsync(void* dst, // 目标地址const void* src, // 源地址size_t count, // 拷贝字节数cudaMemcpyKind kind, // 拷贝方向cudaStream_t stream = 0 // 所属流（默认是 0 流）
);
异步：调用后 CPU 不会阻塞，立即返回。
需要页锁定内存 (pinned memory) 才能实现真正的异步拷贝。
- 如果主机端用 cudaMallocHost() 或 cudaHostAlloc() 分配内存，就能保证异步。
- 否则 CUDA 可能会偷偷同步，导致看起来变成“假异步”。
支持流 (Stream)：
- 不同流里的拷贝和计算可能并发执行（取决于 GPU 是否支持设备重叠）。
同步与异步
__syncthreads()
__syncthreads() 是 一个线程块（Block）内部的同步屏障，用于：
- 同步线程：确保线程块内所有线程在执行到该点时全部停下来，等待所有线程到达后再继续执行。
- 协调共享内存访问：保证线程读写共享内存时数据一致。
它只能作用于同一个Block内的线程，不能跨Block同步。
常见用途：
- 共享内存读写同步
- 避免线程间竞争条件
- 协同计算，比如归约、扫描等
使用示例（共享内存同步）
__global__ void exampleKernel(int *data, int N) {__shared__ int temp[256]; // 共享内存int idx = threadIdx.x + blockIdx.x * blockDim.x;// 每个线程拷贝到共享内存if (idx < N) temp[threadIdx.x] = data[idx];__syncthreads(); // 等待所有线程完成写入共享内存// 读取共享内存if (idx < N) data[idx] = temp[threadIdx.x] * 2;
}
__syncthreads() 保证所有线程都完成了 temp[threadIdx.x] = data[idx] 操作。否则有的线程会提前读取还未写入的共享内存数据，造成错误。
注意：
- 必须在 Block 内所有线程都会执行的路径上调用，如果条件分支里只有部分线程调用 __syncthreads()，会导致 死锁。
cudaStreamSynchronize()
是 CUDA 里用于等待指定流完成所有任务的 API。
cudaError_t cudaStreamSynchronize(cudaStream_t stream);
- 让 CPU 等待某个 CUDA 流（stream） 中的所有操作（包括 kernel、异步内存拷贝等）都执行完毕。
- 调用这个函数会阻塞 CPU 线程，直到该流里的所有任务完成。
- 不会影响其他流，也就是说只同步单独一个流，而不是全局设备。
与其他同步函数对比
-
cudaDeviceSynchronize()
等待整个设备上的所有任务完成，影响全局。
-
cudaStreamSynchronize(stream)
只等待一个流中的任务完成，更细粒度。
-
cudaEventSynchronize(event)
等待一个事件（event）完成，通常用于精确控制。
性能度量
cudaEventCreate()
用来 创建一个 CUDA 事件对象（cudaEvent_t）
cudaError_t cudaEventCreate(cudaEvent_t* event);
常用于
- 计时（配合 cudaEventRecord() 和 cudaEventElapsedTime()）
- 同步（配合 cudaEventRecord()、cudaEventSynchronize()、cudaStreamWaitEvent()）
cudaEventRecord()
用来在 指定的 CUDA 流（stream）上记录一个事件（event）。
cudaError_t cudaEventRecord(cudaEvent_t event, // 要记录的事件cudaStream_t stream = 0 // 事件绑定的流，默认是 0（即默认流）
);
当流执行到 cudaEventRecord() 时，事件会被“标记”为已发生。之后可以用 cudaEventQuery()、cudaEventSynchronize() 来检测/等待事件完成。
cudaEventSynchronize()
是 CUDA 里的一个同步函数，用来等待某个事件完成。
cudaError_t cudaEventSynchronize(cudaEvent_t event);
- 如果 event 还没有完成（比如记录在某个 kernel 执行之后），那么 CPU 会阻塞在这里，直到这个事件对应的所有操作都完成。
- 如果事件已经完成，函数会立即返回。
cudaEventElapsedTime()
是 CUDA 提供的一个函数，用来计算两个事件（cudaEvent_t）之间的时间差（毫秒级）。
cudaError_t cudaEventElapsedTime(float* ms,//指向一个 float 类型的指针，函数会把计算出来的时间（单位：毫秒）存放到这里。cudaEvent_t start,//起始事件，一般用 cudaEventRecord(start, 0) 记录。cudaEvent_t end//结束事件，一般用 cudaEventRecord(end, 0) 记录。
);
一个计算执行时间的例子：
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);cudaEventRecord(start, 0);// 执行一些需要计时的 CUDA kernel
kernel<<<grid, block>>>(...);cudaEventRecord(stop, 0);
cudaEventSynchronize(stop); // 等待 stop 事件完成float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop);printf("Kernel execution time: %f ms\n", elapsedTime);cudaEventDestroy(start);
cudaEventDestroy(stop);
- 结果是毫秒 (ms)，可以精确到 0.5 微秒（根据硬件和驱动）。
- 由于 CUDA 的 异步执行机制，必须在 stop 上调用 cudaEventSynchronize() 才能确保内核执行完成，否则测到的是错误的时间。
由于CUDA事件是直接在GPU上实现的，因此它们不适用于对同时包含设备代码和主机代码的混合代码计时。也就是说，如果你试图通过CUDA事件对核函数和设备内存复制之外的代码进行计时，将得到不可靠的结果。
cudaEventDestroy()
是 CUDA 提供的 API，用来销毁事件对象（cudaEvent_t），释放它所占用的资源。
cudaError_t cudaEventDestroy(cudaEvent_t event);
- 调用后，该事件就不可再使用。
- CUDA 会释放与该事件相关的资源（如驱动端的记录数据）。
- 如果事件仍然在使用中（例如在某个流还没执行到这条记录），销毁行为是 延迟的 —— 会等事件完成后再释放。
纹理内存相关
大部分函数直接在概念中被介绍了，请直接跳转。
原子计算
atomicAdd()
原子加法，有整数、浮点数（cuda8.0)、甚至double类型（cuda11+)
__device__ int atomicAdd(int* address, int val);
__device__ float atomicAdd(float* address, float val);
__device__ double atomicAdd(double* address, double val);
- 普通加法：多线程同时对一个地址加值，可能会出现“读-改-写”冲突，导致结果不正确。
- 原子加法：硬件会保证一个线程完成加法后才允许另一个线程进入这个加法过程，确保结果正确。
常用编程技巧
计算全局索引
- blockIdx.x → 当前 Block 的编号
- blockDim.x → 每个 Block 有多少线程
- threadIdx.x → 当前线程在 Block 内的编号
注意：在使用索引时，很有可能线程会大于数据量，所以一定要判断是否越界
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {C[idx] = A[idx] + B[idx];
}
一维索引
适合处理一维的数组 如A[idx]
int idx = blockIdx.x * blockDim.x + threadIdx.x;
二维索引
适合处理矩阵等二维结构，可以想象成大格套小格。
ID是从 0 开始的，blockIdx.x*blockDim.x是在计算大格中的“左上角”起始点，thread.x算的是小格偏移。
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
int idx = row * (gridDim.x * blockDim.x) + col;//展开为一维
三维索引
适合处理三维数组（如图像数据等）
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
int idx = (z * gridDim.y * blockDim.y + y) * (gridDim.x * blockDim.x) + x;
数据大于线程数时（stride)
当数据量N远远大于总线程数时，应当转换思路，让一个线程处理多个元素。
示例：
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int stride = blockDim.x * gridDim.x;
for (int i = idx; i < N; i += stride) {C[i] = A[i] + B[i];
}
volatile防止缓存错误
在声明变量时，对于常更改且有关键作用（分支作用）的变量，可以使用该关键字，告诉编译器每次都需要真的去内存中查找这个值，以防止编译器“优化”，读到错误的值。
volatile bool dstOut = true;
缩小原子竞争资源范围加速
具体可以看cuda_by_example中第九章原子性的直方图计算核函数例子，它的优化核心思想在于：
通过共享内存，让块间也存在的竞争消失，仅缩小为块内竞争，缩小了资源竞争的范围
宽度优先放置流操作
具体可以看cuda_by_example中第十章高效使用多个CUDA流的例子，它的优化核心思想在于：
根据硬件实际调度流操作的原理，将流操作宽度优先插入，而不是深度优先。
)
