CUDA sample中3_CUDA_Features里包含一些展示 CUDA 各种特性的sample,cdpSimplePrint这个sample使用 CUDA 动态并行性(CDP)实现的简单 printf。
CDP (CUDA Dynamic Paralellism) 允许在运行在 GPU 上的线程中启动内核。CDP 仅在具有 3.5 或更高 SM 架构的 GPU 上可用。
#include <helper_cuda.h>
#include <helper_string.h>
#include <cstdio>
#include <cstdlib>
#include <iostream>
// Variable on the GPU used to generate unique identifiers of blocks.
__device__ int g_uids = 0;
// Print a simple message to signal the block which is currently executing.
__device__ void print_info(int depth, int thread, int uid, int parent_uid) {
if (threadIdx.x == 0) {
if (depth == 0)
printf("BLOCK %d launched by the host\n", uid);
else {
char buffer[32];
for (int i = 0; i < depth; ++i) {
buffer[3 * i + 0] = '|';
buffer[3 * i + 1] = ' ';
buffer[3 * i + 2] = ' ';
}
buffer[3 * depth] = '\0';
printf("%sBLOCK %d launched by thread %d of block %d\n", buffer, uid,
thread, parent_uid);
}
}
__syncthreads();
}
// The kernel using CUDA dynamic parallelism.
//
// It generates a unique identifier for each block. Prints the information
// about that block. Finally, if the 'max_depth' has not been reached, the
// block launches new blocks directly from the GPU.
__global__ void cdp_kernel(int max_depth, int depth, int thread,
int parent_uid) {
// We create a unique ID per block. Thread 0 does that and shares the value
// with the other threads.
__shared__ int s_uid;
if (threadIdx.x == 0) {
s_uid = atomicAdd(&g_uids, 1);
}
__syncthreads();
// We print the ID of the block and information about its parent.
print_info(depth, thread, s_uid, parent_uid);
// We launch new blocks if we haven't reached the max_depth yet.
if (++depth >= max_depth) {
return;
}
cdp_kernel<<<gridDim.x, blockDim.x>>>(max_depth, depth, threadIdx.x, s_uid);
}
// Main entry point.
int main(int argc, char **argv) {
printf("starting Simple Print (CUDA Dynamic Parallelism)\n");
// Parse a few command-line arguments.
int max_depth = 2;
if (checkCmdLineFlag(argc, (const char **)argv, "help") ||
checkCmdLineFlag(argc, (const char **)argv, "h")) {
printf(
"Usage: %s depth=<max_depth>\t(where max_depth is a value between 1 "
"and 8).\n",
argv[0]);
exit(EXIT_SUCCESS);
}
if (checkCmdLineFlag(argc, (const char **)argv, "depth")) {
max_depth = getCmdLineArgumentInt(argc, (const char **)argv, "depth");
if (max_depth < 1 || max_depth > 8) {
printf("depth parameter has to be between 1 and 8\n");
exit(EXIT_FAILURE);
}
}
// Find/set the device.
int device = -1;
cudaDeviceProp deviceProp;
device = findCudaDevice(argc, (const char **)argv);
checkCudaErrors(cudaGetDeviceProperties(&deviceProp, device));
if (!(deviceProp.major > 3 ||
(deviceProp.major == 3 && deviceProp.minor >= 5))) {
printf("GPU %d - %s does not support CUDA Dynamic Parallelism\n Exiting.",
device, deviceProp.name);
exit(EXIT_WAIVED);
}
// Print a message describing what the sample does.
printf(
"*********************************************************************"
"******\n");
printf(
"The CPU launches 2 blocks of 2 threads each. On the device each thread "
"will\n");
printf(
"launch 2 blocks of 2 threads each. The GPU we will do that "
"recursively\n");
printf("until it reaches max_depth=%d\n\n", max_depth);
printf("In total 2");
int num_blocks = 2, sum = 2;
for (int i = 1; i < max_depth; ++i) {
num_blocks *= 4;
printf("+%d", num_blocks);
sum += num_blocks;
}
printf("=%d blocks are launched!!! (%d from the GPU)\n", sum, sum - 2);
printf(
"************************************************************************"
"***\n\n");
// Launch the kernel from the CPU.
printf("Launching cdp_kernel() with CUDA Dynamic Parallelism:\n\n");
cdp_kernel<<<2, 2>>>(max_depth, 0, 0, -1);
checkCudaErrors(cudaGetLastError());
// Finalize.
checkCudaErrors(cudaDeviceSynchronize());
exit(EXIT_SUCCESS);
}
CUDA Dynamic Parallelism(动态并行)
动态并行性是CUDA编程模型的一种扩展,它使CUDA内核能够直接在GPU上创建并同步新的工作。在程序需要的任何点动态创建并行性。能够直接从GPU创建工作,可以减少在host和device之间转移执行控制和数据的需求,因为启动配置决策现在可以由在device上执行的线程在运行时做出。此外,数据依赖的并行工作可以在内核中内嵌生成,利用GPU的硬件调度器和负载均衡器动态地进行,并根据数据驱动的决策或工作负载进行适应。之前需要进行修改以消除递归、不规则循环结构或其他不适合平面、单级并行性的结构的算法和编程模式,现在可以更加透明地表达。
动态并行性仅由计算能力3.5及更高的设备支持。
CUDA执行模型基于thread、block和grid这些原语,其中内核函数定义了block和grid中的单个线程执行的程序。当调用内核函数时,grid的属性由执行配置描述,在CUDA中具有特殊的语法。CUDA中对动态并行性的支持扩展了在设备上运行的线程配置、启动和隐式同步新网格的能力。配置和启动新grid的设备线程属于Parent Grid,而由其调用创建的grid是Child Grid。
Child Grid的调用和完成被很好的嵌套,这意味着在其线程创建的所有Child Grid完成之前,Parent Grid不会被视为完成,运行时还保证了Parent Grid和Child Grid之间的隐式同步。
从device端启动Kernels的语法:
kernel_name<<< Dg, Db, Ns, S >>>([kernel arguments]);
-
Dg
是dim3
类型的,指定grid的维度和大小。 -
Db
是dim3
类型的,指定每个block的维度和大小。 -
Ns
是size_t
类型的,指定除了静态分配的内存之外,为此调用在每个线程块上动态分配的共享内存字节数。Ns
是一个可选参数,默认为 0。 -
S
是cudaStream_t
类型的,指定与此调用关联的Stream。该Stream必须是在与调用相同的grid中分配的。S
是一个可选参数,默认为 NULL 流。
CUDA 中的 CDP (CUDA Dynamic Parallelism) 有以下几个主要的好处:
-
提高并行性和资源利用率:
- CDP 使得子内核可以根据运算需求动态地启动新的子网格,从而更好地利用可用的硬件资源。
- 这样可以增加整个应用的并行度,提高资源利用率。
-
简化编程模型:
- 使用 CDP,开发人员可以更自然地表达算法的并行性,无需显式地管理启动和同步多个内核。
- 这样可以简化编程复杂度,提高开发效率。
-
支持递归和动态任务创建:
- CDP 支持内核内部动态启动新的内核,使得可以实现递归和动态任务创建的编程模式。
- 这样可以更好地适配一些需要动态任务创建的算法,如oct-tree、adaptive mesh refinement 等。
-
增强异构计算能力:
- CDP 使得CPU和GPU之间的协作更加灵活,可以让CPU发起GPU内部的计算任务。
- 这样有助于构建更复杂的异构计算应用。
总的来说,CUDA Dynamic Parallelism 通过提升并行性、简化编程、支持动态任务创建等方式,增强了 CUDA 编程的表达能力和性能,为开发人员提供了更强大的GPU计算工具。
Kernel解读
该Sample的kernel使用了 CUDA Dynamic Parallelism (CDP) ,它演示了如何在 GPU 上动态地创建新的线程块。以下是对代码的解读:
-
全局变量
__device__ int g_uids = 0;
这是一个存储在 GPU device内存上的全局变量,用于生成唯一的block标识符。 -
打印当前block的 ID 及其Parent block和thread
__device__ void print_info(int depth, int thread, int uid, int parent_uid)
- 这是一个device函数,用于打印有关当前执行block的信息,包括它的depth、线程索引、唯一标识符以及Parent block的标识符。
- 该函数使用
printf
打印信息,并确保所有线程在函数返回前都同步。
-
从 GPU 启动嵌套内核
__global__ void cdp_kernel(int max_depth, int depth, int thread, int parent_uid)
- 这是使用 CDP 的主要 CUDA kernel函数。
- 它首先生成一个唯一标识符
s_uid
,并将其存储在共享内存中。 - 接下来,它使用
print_info
函数打印有关当前块的信息。 - 如果当前深度小于最大深度
max_depth
,内核会递归地启动新的Child block。 - 新block的深度增加 1,线程索引设置为当前线程的索引,Parent block的标识符设置为当前块的标识符。
如上图所示,该Sample启动了2个block,每个block中有2个线程,每个线程又启动2个block,由于深度是2,递归到此结束。
这个 CUDA kernel演示了如何使用 CDP 在 GPU 上动态地创建新的block。这可以用于实现复杂的、具有动态任务创建特性的算法,如递归算法、自适应网格细化等。
运行结果:
starting Simple Print (CUDA Dynamic Parallelism)
GPU Device 0: "Ada" with compute capability 8.9
***************************************************************************
The CPU launches 2 blocks of 2 threads each. On the device each thread will
launch 2 blocks of 2 threads each. The GPU we will do that recursively
until it reaches max_depth=2
In total 2+8=10 blocks are launched!!! (8 from the GPU)
***************************************************************************
Launching cdp_kernel() with CUDA Dynamic Parallelism:
BLOCK 0 launched by the host
BLOCK 1 launched by the host
| BLOCK 2 launched by thread 0 of block 1
| BLOCK 5 launched by thread 0 of block 0
| BLOCK 4 launched by thread 0 of block 0
| BLOCK 3 launched by thread 0 of block 1
| BLOCK 7 launched by thread 1 of block 0
| BLOCK 6 launched by thread 1 of block 0
| BLOCK 8 launched by thread 1 of block 1
| BLOCK 9 launched by thread 1 of block 1