CUDA编程之CUDA Sample-3_CUDA_Features-cdpSimplePrint

最新推荐文章于 2024-09-06 18:06:22 发布

翊桐公羽

最新推荐文章于 2024-09-06 18:06:22 发布

阅读量620

点赞数 18

文章标签： linux 运维服务器

本文链接：https://blog.csdn.net/zhujingwen_tongji/article/details/140301104

版权

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

翊桐公羽

关注

18
点赞
踩
18

收藏

觉得还不错? 一键收藏
0
评论
CUDA编程之CUDA Sample-3_CUDA_Features-cdpSimplePrint

CUDA sample中3_CUDA_Features里包含一些展示 CUDA 各种特性的sample，cdpSimplePrint这个sample使用 CUDA 动态并行性（CDP）实现的简单 printf。CDP (CUDA Dynamic Paralellism) 允许在运行在 GPU 上的线程中启动内核。CDP 仅在具有 3.5 或更高 SM 架构的 GPU 上可用。
复制链接

扫一扫