《CUDA Programming Guide》3. Advanced CUDA

Advanced CUDA APIs and Features

`cudaLaunchKernelEx`

cudaLaunchKernelEx：一般的 Kernel 启动只能直接指定四类配置，这是对普通 <<<>>> 语法的拓展，可以添加一些高级属性：

Thread Block Cluster 的大小
L1 Cache 与 Shared Memory 的偏好
是否使用 Programmatic Dependent Launch（PDL）

使用时大概这样：

cudaLaunchConfig_t config = {};
config.gridDim = grid;
config.blockDim = block;
config.dynamicSmemBytes = 0;
config.stream = stream;

cudaLaunchAttribute attributes[1];
attributes[0].id = /* Attritibute ID */;
attributes[0].val = /* Attritibute Value */;

config.attrs = attributes;
config.numAttrs = 1;

// Same as kernel<<<grid, block, 0, stream>>>(arg1, arg2) but with attrs
cudaLaunchKernelEx(&config, kernel, arg1, arg2);

Launching Clusters

如何指定 Cluster Size：

编译期指定：
- 在 Kernel 定义时指定，__cluster_dims__(2, 1, 1) __global__ void kernel(...) { ... }，表示一个 Cluster 包括 $2\times 1\times 1$ 个 Block
- 在 Kernel 定义时也可以指定 __block_size__((1024, 1, 1), (2, 2, 2))，这分别表示 Threads Per Block 和 Blocks Per Cluster，此时 Launch Kernel 时 kernel<<<dim3(8, 8, 8)>>> 就表示的是 Cluster 数量
运行时指定：使用 cudaLaunchAttribute 来指定，有两种：
- cudaLaunchAttributeClusterDimension：表示 Kernel 必须按照该 Cluster Size 执行，与 __cluster_dims__ 含义一致，都表示 Minimum Cluster Size
- cudaLaunchAttributePreferredClusterDimension：表示 Preferred Cluster Size，在满足前面 Minimum Cluster Size 条件的基础上，如果可以按照 Preferred Cluster Size 指定的大小执行，那就按照这个执行，如果不行那就退回 Minimum Cluster Size。必须是 Minimum Cluster Size 的整数倍

注意，为了可以完全分组，Grid 每个维度的 Block 数必须能够被对应的 Cluster 维度整除

Programmatic Dependent Kernel Launch

假设同一 Stream 中有两个 Kernel：

primary<<<..., stream>>>();
secondary<<<..., stream>>>();

这两个 Kernel 会严格地顺序执行，但是某些场景中，secondary 只依赖 primary 的一部分结果，如：

Primary：   Data Generation ----------- Other Work
                                 |
Secondary： Independent Work -- Wait -- Dependent Work

PDL 允许开发者显式标记这两个时间点，从而让两个存在依赖的 Kernel 部分重叠：

__global__ void primary(...) {
    // Data Generation
    cudaTriggerProgrammaticLaunchCompletion();
    // Other Work
}

__global__ void secondary(...) {
    // Independent Work
    cudaGridDependencySynchronize();
    // Dependent Work
}

secondary 需要通过 cudaLaunchKernelEx() 启动，并添加一个特殊属性：

attribute[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
attribute[0].val.programmaticStreamSerializationAllowed = 1;

Batched Memory Transfers

避免对若干个小 Buffer 单独数据传输，将这些数据传输任务批量提交

std::vector<void*> srcs(batch_size);
std::vector<void*> dsts(batch_size);
std::vector<size_t> sizes(batch_size);

// Allocate the src and dst buffers
for (size_t i = 0; i < batch_size - 10; i++) {
    cudaMallocHost(&srcs[i], sizes[i]);
    cudaMalloc(&dsts[i], sizes[i]);
}

int buffer[10];

for (size_t i = batch_size - 10; i < batch_size; i++) {
    srcs[i] = &buffer[10 - (batch_size - i)];
    cudaMalloc(&dsts[i], sizes[i]);
}

// Setup attributes for this batch of copies
cudaMemcpyAttributes attrs[2] = {};
attrs[0].srcAccessOrder = cudaMemcpySrcAccessOrderStream;
attrs[1].srcAccessOrder = cudaMemcpySrcAccessOrderDuringApiCall;

size_t attrsIdxs[2];
attrsIdxs[0] = 0;
attrsIdxs[1] = batch_size - 10;

// Launch the batched memory transfer
cudaMemcpyBatchAsync(&dsts[0], &srcs[0], &sizes[0], batch_size,
    &attrs, &attrsIdxs, 2 /* numAttrs */, nullptr /* failIdx */, stream);

其中可以传递属性值，一个重要的属性值是 srcAccessOrder，用来说明 Runtime 何时被允许读取源数据，有下面的值：

cudaMemcpySrcAccessOrderStream：CUDA 必须等该 Batch Copy 在 Stream 中轮到执行时，才可以访问源数据
cudaMemcpySrcAccessOrderDuringApiCall：对于栈上的临时数据，需要立刻进行数据传输，API 返回前必须把源数据处理好
cudaMemcpyAccessOrderAny：不是短生命周期临时对象，也可脱离 Stream 顺序，但具体最优处理方式依赖平台时，可以允许 Runtime 自行选择合适的访问顺序

不同复制可以使用不同属性，上面代码中 attrsIdxs 的值表示从哪个 Transfer 开始使用下一套属性

批量复制还可以提供源和目标位置提示，如：

// Allocate the source and destination buffers
std::vector<void*> srcs(batch_size);
std::vector<void*> dsts(batch_size);
std::vector<size_t> sizes(batch_size);

// cudaMemLocation structures we will use tp provide location hints
// Device device_id
cudaMemLocation srcLoc = {cudaMemLocationTypeDevice, dev_id};

// Host with numa Node numa_id
cudaMemLocation dstLoc = {cudaMemLocationTypeHostNuma, numa_id};

// Allocate the src and dst buffers
for (size_t i = 0; i < batch_size; i++) {
    cudaMallocManaged(&srcs[i], sizes[i]);
    cudaMallocManaged(&dsts[i], sizes[i]);

    cudaMemPrefetchAsync(srcs[i], sizes[i], srcLoc, 0, stream);
    cudaMemPrefetchAsync(dsts[i], sizes[i], dstLoc, 0, stream);
    cudaMemsetAsync(srcs[i], sizes[i], stream);
}

// Setup attributes for this batch of copies
cudaMemcpyAttributes attrs = {};

// These are managed memory pointers so Stream Order is appropriate
attrs.srcAccessOrder = cudaMemcpySrcAccessOrderStream;

// Now we can specify the location hints here.
attrs.srcLocHint = srcLoc;
attrs.dstlocHint = dstLoc;

// All copies in the batch have same copy attributes.
size_t attrsIdxs = 0;

// Launch the batched memory transfer
cudaMemcpyBatchAsync(&dsts[0], &srcs[0], &sizes[0], batch_size,
    &attrs, &attrsIdxs, 1 /* numAttrs */, nullptr /* failIdx */, stream);

依然是 Hint，不能简单理解为强制数据一定驻留在该位置，可以帮助 Runtime 更好地去优化数据传输

Advanced CUDA APIs and Features

cudaLaunchKernelEx

Launching Clusters

Programmatic Dependent Kernel Launch

Batched Memory Transfers

Advanced Kernel Programming

`cudaLaunchKernelEx`