【CUDA编程】系列博客参考NVIDIA官方文档“CUDA C++ Programming Guide（v12.6）”。
本文为原创文章，未经本人允许，禁止转载。转载请注明出处。
1.Device Enumeration
host系统可以包含多个device。以下代码展示了如何枚举这些device、查询它们的属性并确定支持CUDA的device数量。
1
2
3
4
5
6
7
8
9
int deviceCount;
cudaGetDeviceCount(&deviceCount); //返回计算能力大于等于2.0的device数量
int device;
for (device = 0; device < deviceCount; ++device) {
cudaDeviceProp deviceProp; //CUDA device属性结构体
cudaGetDeviceProperties(&deviceProp, device); //返回device属性
printf("Device %d has compute capability %d.%d.\n",
device, deviceProp.major, deviceProp.minor);
}
2.Device Selection
host线程可以通过调用cudaSetDevice()随时设置其操作的device。device的内存分配和kernel启动均在当前设置的device上完成，stream和event也与当前设置的device相关联。如果没有调用cudaSetDevice()，则默认当前device为device 0。
1
2
3
4
5
6
7
8
9
size_t size = 1024 * sizeof(float);
cudaSetDevice(0); // Set device 0 as current
float* p0;
cudaMalloc(&p0, size); // Allocate memory on device 0
MyKernel<<<1000, 128>>>(p0); // Launch kernel on device 0
cudaSetDevice(1); // Set device 1 as current
float* p1;
cudaMalloc(&p1, size); // Allocate memory on device 1
MyKernel<<<1000, 128>>>(p1); // Launch kernel on device 1
3.Stream and Event Behavior
如果在一个与当前device不关联的stream中启动kernel，启动将失败，如以下代码所示：
1
2
3
4
5
6
7
8
9
10
11
cudaSetDevice(0); // Set device 0 as current
cudaStream_t s0;
cudaStreamCreate(&s0); // Create stream s0 on device 0
MyKernel<<<100, 64, 0, s0>>>(); // Launch kernel on device 0 in s0
cudaSetDevice(1); // Set device 1 as current
cudaStream_t s1;
cudaStreamCreate(&s1); // Create stream s1 on device 1
MyKernel<<<100, 64, 0, s1>>>(); // Launch kernel on device 1 in s1
// This kernel launch will fail:
MyKernel<<<100, 64, 0, s0>>>(); // Launch kernel on device 1 in s0
但内存拷贝即使是针对与当前device不关联的stream，也可以成功完成。
如果输入event和输入stream属于不同device，则cudaEventRecord()的调用将会失败。
如果两个输入event属于不同device，则cudaEventElapsedTime()的调用将会失败。
即使输入event关联的device和当前device不同，cudaEventSynchronize()和cudaEventQuery()的调用也会成功。
即使输入stream和输入event属于不同device，cudaStreamWaitEvent()的调用也会成功。因此，cudaStreamWaitEvent()可用于同步多个device之间的操作。
每个device都有自己的默认stream，因此，一个device的默认stream可以无序或并发地执行，与其他device的默认stream发出的命令无冲突。
4.Peer-to-Peer Memory Access
根据系统属性，特别是PCIe和/或NVLINK拓扑结构，device之间可以访问彼此的内存（例如，一个device上运行的kernel可以解引用指向另一个device内存的指针）。如果两个device的cudaDeviceCanAccessPeer()调用都返回true，则说明这两个device之间支持这种点对点的内存访问（peer-to-peer memory access）功能。
点对点的内存访问仅在64位应用程序中支持，且必须通过调用cudaDeviceEnablePeerAccess()在两个device之间启用。在不支持NVSwitch的系统中，每个device最多支持8个点对点连接。
对于启用了统一虚拟地址空间的device，可以使用相同的指针从两个device访问内存。
1
2
3
4
5
6
7
8
9
10
11
12
13
14
cudaSetDevice(0); // Set device 0 as current
float* p0;
size_t size = 1024 * sizeof(float);
cudaMalloc(&p0, size); // Allocate memory on device 0
MyKernel<<<1000, 128>>>(p0); // Launch kernel on device 0
cudaSetDevice(1); // Set device 1 as current
//第一个参数0表示要和device 0启用点对点内存访问
//第二个参数0是标志位，目前只能设置为0
cudaDeviceEnablePeerAccess(0, 0); // Enable peer-to-peer access
// with device 0
// Launch kernel on device 1
// This kernel launch can access memory on device 0 at address p0
MyKernel<<<1000, 128>>>(p0);
4.1.IOMMU on Linux
在Linux中使用IOMMU（输入输出内存管理单元）存在一些限制，而Windows系统中则不存在这些限制。
5.Peer-to-Peer Memory Copy
可以在两个不同device之间执行内存拷贝。
当两个device使用统一虚拟地址空间时，这可以通过常规的内存拷贝函数（见：Device Memory）完成。
否则，可以使用cudaMemcpyPeer()、cudaMemcpyPeerAsync()、cudaMemcpy3DPeer()、cudaMemcpy3DPeerAsync()等点对点内存拷贝函数完成。
1
2
3
4
5
6
7
8
9
10
11
12
cudaSetDevice(0); // Set device 0 as current
float* p0;
size_t size = 1024 * sizeof(float);
cudaMalloc(&p0, size); // Allocate memory on device 0
cudaSetDevice(1); // Set device 1 as current
float* p1;
cudaMalloc(&p1, size); // Allocate memory on device 1
cudaSetDevice(0); // Set device 0 as current
MyKernel<<<1000, 128>>>(p0); // Launch kernel on device 0
cudaSetDevice(1); // Set device 1 as current
cudaMemcpyPeer(p1, 1, p0, 0, size); // Copy p0 to p1
MyKernel<<<1000, 128>>>(p1); // Launch kernel on device 1
cudaMemcpyPeer函数用于在两个device之间拷贝内存：
1
2
3
4
5
6
7
8
//参数详解：
//dst：目标device的指针
//dstDevice：目标device
//src：源device的指针
//srcDevice：源device
//count：拷贝的内存大小（单位：字节）
__host__cudaError_t cudaMemcpyPeer (void *dst, int
dstDevice, const void *src, int srcDevice, size_t count)
两个不同device之间的内存拷贝（在隐式的NULL stream中）：
- 必须在两个device完成之前发出的所有命令后，才能开始内存拷贝。
- 必须在内存拷贝完成之后，才会开始其他的任务。
在点对点内存拷贝中，内存拷贝可以与其他stream中的内存拷贝或kernel并行执行，从而提高效率。
如果两个device通过cudaDeviceEnablePeerAccess()启用了点对点内存访问，则内存拷贝可以直接在device间进行，而无需通过host内存中转。这将显著提升内存拷贝速度。
