线程块：

• 三维向量，
  线程索引： index = x + yDx + zDx * Dy
  假如 Dx = 10, Dy = 20 , Dz = 30,对于线程(2, 5, 9) 的index = 2 + 5 * 10 + 9 * 10 * 20 = 1852, 总线程数 Dx * Dy * Dz = 10 * 20 * 30 = 6000 = 6k

假如Dx = 6, Dy = 4, Dz = 2 ，如下是对应的线程id编号

第一层：

第二层

内存模型：

静态与动态内存的区别

• 统一内存
• 异步SIMT编程
  异步屏障（Asynchronous Barrier）
  用于实现CUDA线程间的同步机制。

异步内存拷贝（cuda::memcpy_async）
允许在GPU执行计算的同时，异步地从全局内存移动数据。
同步对象：cuda::barrier or a cuda::pipeline，
同步范围： cuda::thread_scope::thread_scope_thread， cuda::thread_scope::thread_scope_block，cuda::thread_scope::thread_scope_device，cuda::thread_scope::thread_scope_system

• 计算能力 SM version

编译

• 即时编译（JIT）: ptx 动态加载到设备驱动中， 使用 NVRTC 可以支持运行时编译 cuda c++
• 二进制兼容 ， 通过 -code=sm_80 编译成针对计算能力8.0的二进制，
• ptx兼容， 通过 -arch=compute_50 指定C++编译为PTX代码时采用的计算能力，

cuda 运行态

• 每个设备创建一个CUDA上下文
• cudaInitDevice() 和 cudaSetDevice() 会初始化指定设备的上下文
• cudaDeviceReset() 销毁上下文
• CUDA 数组（CUDA arrays for texture fetching） & 线性内存（Linear memory）
• cudaMallocPitch(&devPtr, &pitch, width * sizeof(float), height); // pitch 表示间距
• cudaMallocPitch cudaMalloc3D, cudaMemcpy2D, cudaMemcpy3D // 更好的满足 存储对齐要求
• 支持L2缓存的管理，为持久化访问分配大小
• 页锁主机内存 cudaHostAlloc() and cudaFreeHost() ， 可以通过 cudaHostAllocPortable 为所有设备设置可移植性， 通过 cudaHostAllocWriteCombined 支持写合并模式（应用场景：CPU向写入、GPU读取​​）， 写合并内存会释放cpu L1 L2 缓存，同时PCIe总线上传输数据时，写合并内存不会被缓存系统窥探，缺点： cpu读取很慢， 写入很快。
• 内存映射 cudaHostAlloc() 传入 cudaHostAllocMapped ，将一块锁页主机内存映射到设备的地址空间， 主机内存中的地址cudaHostAlloc返回，cudaHostGetDevicePointer获取设备内存地址，好处： 可以在kernel函数中直接访问主机内存，没有设备内存的占用、拷贝和传输，无需使用流。在使用之前 需要通过cudaSetDeviceFlags() 并传入 cudaDeviceMapHost。
• 环境变量 CUDA_LAUNCH_BLOCKING=1 用于禁止内核启动的异步性，用于调试。
• 并行： cpu与gpu并行， kernel的并行， 数据传输与kernel的并行，数据传输与数据传输之间的并行，流与流之间的并行。

要通过流实现内存拷贝间的重叠行为， 页锁是必须的：因为如果使用普通内存，则cudaMemcpyAsync会退化为同步操作， 因为要等待cpu固定页面，而页锁可以使GPU的直接访问内存（DMA）。

cudaStream_t stream[2];
for (int i = 0; i < 2; ++i)
    cudaStreamCreate(&stream[i]);
float* hostPtr;
cudaMallocHost(&hostPtr, 2 * size);

for (int i = 0; i < 2; ++i) {
    cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,
                    size, cudaMemcpyHostToDevice, stream[i]);
    MyKernel <<<100, 512, 0, stream[i]>>>
          (outputDevPtr + i * size, inputDevPtr + i * size, size);
    cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size,
                    size, cudaMemcpyDeviceToHost, stream[i]);
}

如果GPU设备不支持并行的数据传输 ， 其优化方案如下：

for (int i = 0; i < 2; ++i)
    cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,
                    size, cudaMemcpyHostToDevice, stream[i]);
for (int i = 0; i < 2; ++i)
    MyKernel<<<100, 512, 0, stream[i]>>>
          (outputDevPtr + i * size, inputDevPtr + i * size, size);
for (int i = 0; i < 2; ++i)
    cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size,
                    size, cudaMemcpyDeviceToHost, stream[i]);

这样， 提交到流[1]的主机到设备内存拷贝操作，将与提交到流[0]的内核启动操作产生执行重叠。

• 流之间的显示同步： cudaDeviceSynchronize(), cudaStreamSynchronize(), cudaStreamWaitEvent(),cudaStreamQuery(),
• cudaLaunchHostFunc()在流中任意位置插入CPU函数调用的机制,该函数将在主机上执行， 会等待该函数执行完成后，才继续执行后续的流中的操作， 被加入流中的主机函数不得直接或间接调用任何CUDA API函数，否则可能导致函数等待自身执行完成，从而引发死锁。

void CUDART_CB MyCallback(void *data){
    printf("Inside callback %d\n", (size_t)data);
}
...
for (size_t i = 0; i < 2; ++i) {
    cudaMemcpyAsync(devPtrIn[i], hostPtr[i], size, cudaMemcpyHostToDevice, stream[i]);
    MyKernel<<<100, 512, 0, stream[i]>>>(devPtrOut[i], devPtrIn[i], size);
    cudaMemcpyAsync(hostPtr[i], devPtrOut[i], size, cudaMemcpyDeviceToHost, stream[i]);
    cudaLaunchHostFunc(stream[i], MyCallback, (void*)i);
}

• 创建流时通过cudaStreamCreateWithPriority()指定流的相对优先级。通过cudaDeviceGetStreamPriorityRange()函数可获取设备的优先级范围,

// get the range of stream priorities for this device
int leastPriority, greatestPriority;
cudaDeviceGetStreamPriorityRange(&leastPriority, &greatestPriority);
// create streams with highest and lowest available priorities
cudaStream_t st_high, st_low;
cudaStreamCreateWithPriority(&st_high, cudaStreamNonBlocking, greatestPriority));
cudaStreamCreateWithPriority(&st_low, cudaStreamNonBlocking, leastPriority);

• 编程式依赖启动与同步， （preamble section 部分的并行）， 场景：第二个kernel依赖第一个，但是第二个和第一个之间，并不是完全的串行

实现如下: 注意并发和死锁

__global__ void primary_kernel() {
   // Initial work that should finish before starting secondary kernel

   // Trigger the secondary kernel
   cudaTriggerProgrammaticLaunchCompletion();

   // Work that can coincide with the secondary kernel
}

__global__ void secondary_kernel()
{
   // Independent work

   // Will block until all primary kernels the secondary kernel is dependent on have completed and flushed results to global memory
   cudaGridDependencySynchronize();

   // Dependent work
}

cudaLaunchAttribute attribute[1];
attribute[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
attribute[0].val.programmaticStreamSerializationAllowed = 1;
configSecondary.attrs = attribute;
configSecondary.numAttrs = 1;

primary_kernel<<<grid_dim, block_dim, 0, stream>>>();
cudaLaunchKernelEx(&configSecondary, secondary_kernel);

也可以使用cuda 图的方式，待研究...

*cuda图： 可以通过api创建， 也可以通过流捕获，cudaStreamBeginCapture() 会将流置于捕获模式。当流处于捕获状态时，提交到该流的工作不会立即排队执行，而是逐步追加到一个正在构建的内部图中

• 跨流的之间的依赖，可以通过在主流中使用cudaEventRecord(), 在分流中使用cudaStreamWaitEvent(), 这样相当于告诉捕获程序，这儿有一个分叉，此时，分流会被自动置于捕获模式，并且将后续的操作捕获在主流的cuda图中，并且要求流2要合并如主流中，cudaEventRecord(event2, stream2);cudaStreamWaitEvent(stream1, event2);
• cudaUserObject_t 用于cuda图捕获的生命周期
• cuda 图更新
• 设备图启动机制：工作流需要在运行时根据数据做出决策，并执行不同的操作。相较于将这些决策过程交由主机处理, 可能导致设备与主机间的往返延迟
• 可从设备端启动的图称为"设备图"
• 无法从设备端启动的图称为"主机图"
• 设备图的创建 cudaGraphInstantiate()传递cudaGraphInstantiateFlagDeviceLaunch
• 设备图要求： 图内所有节点必须位于同一设备， 仅允许包含以下节点类型：内核节点、内存拷贝节点、内存设置节点和子图节点。
• 设备图上传： 使用cudaGraphUpload()函数直接上传，通过cudaGraphInstantiateWithParams()在实例化时请求上传
• 设备图启动：

即发即弃启动：

尾部启动：

同级启动：

• 图执行环境满足层级结构​
• 条件节点： cudaGraphConditionalHandleCreate() 创建条件句柄，cudaGraphAddConditionalNode() 添加条件节点， cudaGraphSetConditional()更新条件值，在创建条件节点时，系统会生成一个空白的子图并返回其句柄，以便用户填充子图内容。该条件体子图可通过以下两种方式构建：使用标准图API直接编辑， 通过cudaStreamBeginCaptureToGraph()函数从流捕获。

事件

• 为了监视设备的执行过程及进度， 在程序任意位置插入记录事件。
• 事件机制构成：cudaEventRecord() + cudaEventQuery()/cudaEventSynchronize()
• 支持cudaEventElapsedTime()计算纳秒级时间间隔

多设备系统

• cudaSetDevice， cudaGetDeviceCount， cudaDeviceProp ， cudaGetDeviceProperties
• 在 PCIe/NVLINK 连接的情况下， 可以cudaDeviceEnablePeerAccess(peerDeviceId, 0);激活访问权限， 支持直接指定访问和同一地址空间访问。
• 设备间内存复制：统一地址空间，直接拷贝，非统一地址空间： cudaMemcpyPeer(), cudaMemcpyPeerAsync(), cudaMemcpy3DPeer()，cudaMemcpy3DPeerAsync()等拷贝

对等拷贝：

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

统一虚拟地址空间

当应用程序以64位进程运行时，主机与所有计算能力2.0及以上的设备将共享统一的虚拟地址空间，

cudaPointerAttributes attributes;  
cudaPointerGetAttributes(&attributes, ptr);  // 可确定指针位于主机/设备内存

cudaMemcpy(dest, src, size, cudaMemcpyDefault);  // 自动识别指针位置

• 错误检查： cudaDeviceSynchronize()后使用 cudaGetLastError()进行错误检查

Texture and Surface Memory //TODO