CUDA编程之CUDA Sample-3_CUDA_Features-cdpSimpleQuicksort

翊桐公羽

于 2024-07-12 11:35:07 发布

阅读量867

点赞数 16

文章标签：算法

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

版权

CUDA sample中3_CUDA_Features里包含一些展示 CUDA 各种特性的sample，cdpSimpleQuicksort这个sample使用 CUDA 动态并行性（CDP）实现了快速排序。

CDP (CUDA Dynamic Paralellism) 允许在运行在 GPU 上的线程中启动内核。动态并行性在 CUDA 5.0 及更高版本的设备上可用,这些设备的计算能力为 3.5 或更高(sm_35)。

早期的 CUDA 程序必须遵循一种扁平、批量并行的编程模型。程序必须执行一系列的内核启动,而且为了获得最佳性能,每个内核都必须暴露足够的并行性来有效地利用 GPU。对于由"并行 for"循环组成的应用程序,批量并行模型并不太限制性,但是某些并行模式(如嵌套并行性)就无法轻易表达。

嵌套并行性自然会出现在许多应用程序中,例如使用自适应网格的应用程序,这种网格常用于现实世界的应用程序中,以减少计算复杂度,同时捕捉相关的细节级别。扁平的批量并行应用程序必须使用细粒度的网格,并执行不需要的计算,或使用粗粒度的网格并丢失更细的细节。

CUDA 5.0 引入了动态并行性,这使得从device上运行的线程启动内核成为了可能;线程可以启动更多的线程。应用程序可以启动一个粗粒度的内核,该内核反过来启动更细粒度的内核来执行所需的工作。这避免了不需要的计算,同时捕捉所有有趣的细节,正下图所示：

动态并行性通常对于无法避免嵌套并行性的问题很有用。这包括但不限于以下算法类别:

使用层次数据结构(如自适应网格)的算法;
使用递归的算法,其中每个递归级别都有并行性,如快速排序;
工作自然分割成独立批次的算法,其中每个批次涉及复杂的并行处理,但无法完全利用单个 GPU。

这些算法通常难以用扁平的批量并行模型来表达和实现。动态并行性允许在 GPU 上灵活地创建和管理嵌套的并行计算。这不仅提高了性能,还让程序员能够更自然地表达并行算法的结构。

#include <iostream>
#include <cstdio>
#include <helper_cuda.h>
#include <helper_string.h>

#define MAX_DEPTH 16
#define INSERTION_SORT 32


// Selection sort used when depth gets too big or the number of elements drops
// below a threshold.

__device__ void selection_sort(unsigned int *data, int left, int right) {
  for (int i = left; i <= right; ++i) {
    unsigned min_val = data[i];
    int min_idx = i;

    // Find the smallest value in the range [left, right].
    for (int j = i + 1; j <= right; ++j) {
      unsigned val_j = data[j];

      if (val_j < min_val) {
        min_idx = j;
        min_val = val_j;
      }
    }

    // Swap the values.
    if (i != min_idx) {
      data[min_idx] = data[i];
      data[i] = min_val;
    }
  }
}


// Very basic quicksort algorithm, recursively launching the next level.

__global__ void cdp_simple_quicksort(unsigned int *data, int left, int right,
                                     int depth) {
  // If we're too deep or there are few elements left, we use an insertion
  // sort...
  if (depth >= MAX_DEPTH || right - left <= INSERTION_SORT) {
    selection_sort(data, left, right);
    return;
  }

  unsigned int *lptr = data + left;
  unsigned int *rptr = data + right;
  unsigned int pivot = data[(left + right) / 2];

  // Do the partitioning.
  while (lptr <= rptr) {
    // Find the next left- and right-hand values to swap
    unsigned int lval = *lptr;
    unsigned int rval = *rptr;

    // Move the left pointer as long as the pointed element is smaller than the
    // pivot.
    while (lval < pivot) {
      lptr++;
      lval = *lptr;
    }

    // Move the right pointer as long as the pointed element is larger than the
    // pivot.
    while (rval > pivot) {
      rptr--;
      rval = *rptr;
    }

    // If the swap points are valid, do the swap!
    if (lptr <= rptr) {
      *lptr++ = rval;
      *rptr-- = lval;
    }
  }

  // Now the recursive part
  int nright = rptr - data;
  int nleft = lptr - data;

  // Launch a new block to sort the left part.
  if (left < (rptr - data)) {
    cudaStream_t s;
    cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking);
    cdp_simple_quicksort<<<1, 1, 0, s>>>(data, left, nright, depth + 1);
    cudaStreamDestroy(s);
  }

  // Launch a new block to sort the right part.
  if ((lptr - data) < right) {
    cudaStream_t s1;
    cudaStreamCreateWithFlags(&s1, cudaStreamNonBlocking);
    cdp_simple_quicksort<<<1, 1, 0, s1>>>(data, nleft, right, depth + 1);
    cudaStreamDestroy(s1);
  }
}


// Call the quicksort kernel from the host.

void run_qsort(unsigned int *data, unsigned int nitems) {
  // Prepare CDP for the max depth 'MAX_DEPTH'.

  // Launch on device
  int left = 0;
  int right = nitems - 1;
  std::cout << "Launching kernel on the GPU" << std::endl;
  cdp_simple_quicksort<<<1, 1>>>(data, left, right, 0);
  checkCudaErrors(cudaDeviceSynchronize());
}


// Initialize data on the host.

void initialize_data(unsigned int *dst, unsigned int nitems) {
  // Fixed seed for illustration
  srand(2047);

  // Fill dst with random values
  for (unsigned i = 0; i < nitems; i++) dst[i] = rand() % nitems;
}


// Verify the results.

void check_results(int n, unsigned int *results_d) {
  unsigned int *results_h = new unsigned[n];
  checkCudaErrors(cudaMemcpy(results_h, results_d, n * sizeof(unsigned),
                             cudaMemcpyDeviceToHost));

  for (int i = 1; i < n; ++i)
    if (results_h[i - 1] > results_h[i]) {
      std::cout << "Invalid item[" << i - 1 << "]: " << results_h[i - 1]
                << " greater than " << results_h[i] << std::endl;
      exit(EXIT_FAILURE);
    }

  std::cout << "OK" << std::endl;
  delete[] results_h;
}


// Main entry point.

int main(int argc, char **argv) {
  int num_items = 128;
  bool verbose = false;

  if (checkCmdLineFlag(argc, (const char **)argv, "help") ||
      checkCmdLineFlag(argc, (const char **)argv, "h")) {
    std::cerr << "Usage: " << argv[0]
              << " num_items=<num_items>\twhere num_items is the number of "
                 "items to sort"
              << std::endl;
    exit(EXIT_SUCCESS);
  }

  if (checkCmdLineFlag(argc, (const char **)argv, "v")) {
    verbose = true;
  }

  if (checkCmdLineFlag(argc, (const char **)argv, "num_items")) {
    num_items = getCmdLineArgumentInt(argc, (const char **)argv, "num_items");

    if (num_items < 1) {
      std::cerr << "ERROR: num_items has to be greater than 1" << std::endl;
      exit(EXIT_FAILURE);
    }
  }

  // Find/set device and get device properties
  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);
  }

  // Create input data
  unsigned int *h_data = 0;
  unsigned int *d_data = 0;

  // Allocate CPU memory and initialize data.
  std::cout << "Initializing data:" << std::endl;
  h_data = (unsigned int *)malloc(num_items * sizeof(unsigned int));
  initialize_data(h_data, num_items);

  if (verbose) {
    for (int i = 0; i < num_items; i++)
      std::cout << "Data [" << i << "]: " << h_data[i] << std::endl;
  }

  // Allocate GPU memory.
  checkCudaErrors(
      cudaMalloc((void **)&d_data, num_items * sizeof(unsigned int)));
  checkCudaErrors(cudaMemcpy(d_data, h_data, num_items * sizeof(unsigned int),
                             cudaMemcpyHostToDevice));

  // Execute
  std::cout << "Running quicksort on " << num_items << " elements" << std::endl;
  run_qsort(d_data, num_items);

  // Check result
  std::cout << "Validating results: ";
  check_results(num_items, d_data);

  free(h_data);
  checkCudaErrors(cudaFree(d_data));

  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之间的隐式同步。

以下是非CDP和使用CDP之间的区别：

从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计算工具。

快速排序

快速排序是典型的分治算法
递归地将数据分区和排序
完全依赖于数据的执行
在 Fermi 架构上实现高效的分治算法比较困难的

分治算法通常将一个大问题分解成多个较小的子问题,然后递归地解决这些子问题,最后将结果合并。这种方法在许多算法中都有广泛应用,如快速排序、归并排序等。

快速排序它的工作原理如下:

选择一个基准元素(通常选择数组的第一个或最后一个元素)。
将数组中小于基准元素的元素移到左边,大于基准元素的元素移到右边。这个过程称为分区。
对左右两个子数组递归地应用步骤 1 和 2。
当子数组只有 0 或 1 个元素时,排序结束。

快速排序的时间复杂度通常为 O(n log n),是一种非常高效的排序算法。但是,它的最坏情况时间复杂度会退化到 O(n^2),比如当输入数组已经是有序的或逆序的情况下。

Kernel解读

这个sample的kernel实现了一个基本的快速排序算法，并在排序过程中结合了选择排序。当递归深度过大或元素数量低于某个阈值时，选择使用选择排序来处理。这段代码首先定义了选择排序的设备函数，然后定义了在GPU上运行的快速排序内核，并在主机上调用该内核。

以下是代码的解释：


// 当递归深度太大或元素数量低于阈值时使用选择排序。

__device__ void selection_sort(unsigned int *data, int left, int right) {
  // 对区间 [left, right] 内的元素进行选择排序。
  for (int i = left; i <= right; ++i) {
    unsigned min_val = data[i];
    int min_idx = i;

    // 找到区间 [left, right] 内的最小值。
    for (int j = i + 1; j <= right; ++j) {
      unsigned val_j = data[j];

      if (val_j < min_val) {
        min_idx = j;
        min_val = val_j;
      }
    }

    // 交换值。
    if (i != min_idx) {
      data[min_idx] = data[i];
      data[i] = min_val;
    }
  }
}


// 非常基本的快速排序算法，递归地启动下一层。

__global__ void cdp_simple_quicksort(unsigned int *data, int left, int right,
                                     int depth) {
  // 如果递归深度太大或剩余的元素太少，使用选择排序。
  if (depth >= MAX_DEPTH || right - left <= INSERTION_SORT) {
    selection_sort(data, left, right);
    return;
  }

  unsigned int *lptr = data + left;
  unsigned int *rptr = data + right;
  unsigned int pivot = data[(left + right) / 2];

  // 进行分区操作。
  while (lptr <= rptr) {
    // 找到需要交换的左右值。
    unsigned int lval = *lptr;
    unsigned int rval = *rptr;

    // 当左侧指针指向的值小于枢纽值时，移动左侧指针。
    while (lval < pivot) {
      lptr++;
      lval = *lptr;
    }

    // 当右侧指针指向的值大于枢纽值时，移动右侧指针。
    while (rval > pivot) {
      rptr--;
      rval = *rptr;
    }

    // 如果交换点有效，交换值。
    if (lptr <= rptr) {
      *lptr++ = rval;
      *rptr-- = lval;
    }
  }

  // 现在进行递归部分。
  int nright = rptr - data;
  int nleft = lptr - data;

  // 启动一个新的线程块来排序左侧部分。
  if (left < (rptr - data)) {
    cudaStream_t s;
    cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking);
    cdp_simple_quicksort<<<1, 1, 0, s>>>(data, left, nright, depth + 1);
    cudaStreamDestroy(s);
  }

  // 启动一个新的线程块来排序右侧部分。
  if ((lptr - data) < right) {
    cudaStream_t s1;
    cudaStreamCreateWithFlags(&s1, cudaStreamNonBlocking);
    cdp_simple_quicksort<<<1, 1, 0, s1>>>(data, nleft, right, depth + 1);
    cudaStreamDestroy(s1);
  }
}


// 从主机调用快速排序内核。

void run_qsort(unsigned int *data, unsigned int nitems) {
  // 为最大深度 'MAX_DEPTH' 准备 CDP。

  // 在设备上启动快速排序内核。
  int left = 0;
  int right = nitems - 1;
  std::cout << "Launching kernel on the GPU" << std::endl;
  cdp_simple_quicksort<<<1, 1>>>(data, left, right, 0);
  checkCudaErrors(cudaDeviceSynchronize());
}

关键点解释：

选择排序 (selection_sort)：
- 这是一个device函数，用于对较小的数组进行排序。当递归深度太大或数组元素数量太少时使用。
快速排序内核 (cdp_simple_quicksort)：
- 这是一个在GPU上运行的内核函数，使用快速排序算法对数组进行递归排序。
- 使用枢纽值将数组分成两部分，并对每部分递归调用快速排序。
- 当递归深度超过 MAX_DEPTH 或元素数量少于 INSERTION_SORT 时，改用选择排序。
主机函数 (run_qsort)：
- 这是在host上调用快速排序内核的函数。
- 初始化排序范围，并启动快速排序内核。
- 使用 cudaDeviceSynchronize 确保设备上的排序操作完成。

注意事项：

递归深度和元素阈值：在快速排序中，递归深度和元素数量阈值对于性能和资源使用至关重要。合理设置这些参数可以避免过度递归和提升性能。
CUDA流 (cudaStream_t)：使用非阻塞的CUDA流来并行执行多个内核，提升排序效率。

注意事项：

该sample的深度最多可以达到24层，可以通过MAX_DEPTH设置深度。

尽量避免在device端使用cudaDeviceSynchronize，会造成资源浪费。

Runtime的exception在Child kernel里捕获不到，需要在host端调用。

运行结果：

> .\cdpSimpleQuicksort.exe num_items=10000
GPU Device 0: "Ada" with compute capability 8.9

Initializing data:
Running quicksort on 10000 elements
Launching kernel on the GPU
Validating results: OK

翊桐公羽

关注

16
点赞
踩
10

收藏

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

CUDA sample中3_CUDA_Features里包含一些展示 CUDA 各种特性的sample，cdpSimpleQuicksort这个sample使用 CUDA 动态并行性（CDP）实现了快速排序。CDP (CUDA Dynamic Paralellism) 允许在运行在 GPU 上的线程中启动内核。动态并行性在 CUDA 5.0 及更高版本的设备上可用,这些设备的计算能力为 3.5 或更高(sm_35)。早期的 CUDA 程序必须遵循一种扁平、批量并行的编程模型。
复制链接

扫一扫