Cuda 中的嵌套并行性

Question

在下面的代码中，我想使用嵌套并行计算数组元素的 10 次。我使用这个简单的示例来了解有关 Cuda 中的动态并行性的更多信息。代码的工作方式是，对于 parentArray 的每个元素，都有另一个内核将此元素保存在 childArray 的位置（0 到 9）。因此，对于 parentArray 的每个元素，我有另一个包含 10 个元素的数组，每个元素都等于 parentArray 的元素。最后，我计算所有 childArray 的总和并将结果保存在 parentArray 中。

因此结果应该是：

parentArray 的元素 0，结果 = 0
parentArray 的元素 1，结果 = 10
parentArray 的元素 2，Result = 20 等等

目前，代码可以编译但没有给出预期的结果。当前代码有什么问题？

计算元素总和的函数

__device__ double summe(double *arr, int size)
{
  double result = 0.0;
  for(int i = 0; i < size; i++)
  {
    result += arr[i];
  }
  return result;
}

子内核调用的函数

__device__ double getElement(double arrElement)
{
  return arrElement;
}

存储结果的数组

__device__ double childArr[10];

子内核

__global__ void childKernel(double *arr, double arrElement,int N)
{
  int cidx = blockIdx.x * blockDim.x + threadIdx.x;
  if (cidx < N)
  {
    arr[cidx] = getElement(arrElement);
  }
}

父内核

__global__ void parentKernel(double *parentArray, int N)
{
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  if (idx < N)
  {
    childKernel<<<1,10>>>(childArr,parentArray[idx],N);
    __syncthreads();
    parentArray[idx] = summe(childArr,10);

  }

}

主要部分

 int main(void)
    {

      double *host_array;
      double *device_array;

      // Number of elements in arrays
      const int N_array = 10;

      // size of array
      const size_t size_array = N_array * sizeof(double);

      // Allocate array on host
      host_array = (double *)malloc(size_array);

      // Allocate array on device
      CUDA_CALL(cudaMalloc((void **) &device_array, size_array));

      // Initialize host array
      for (int i=0; i<N_array; i++)
      {
        host_array[i] = (double)i;
      }

      // and copy it to CUDA device
      CUDA_CALL(cudaMemcpy(device_array, host_array, size_array, cudaMemcpyHostToDevice));

      // Do calculation on device:
      int block_size = 4;
      // if N = 10, then n_blocks = 3
      int n_blocks = N_array/block_size + (N_array % block_size == 0 ? 0:1);

      parentKernel<<<n_blocks, block_size>>>(device_array,N_array);

      // Retrieve result from device and store it in host array
      CUDA_CALL(cudaMemcpy(host_array, device_array, sizeof(double)*N_array, cudaMemcpyDeviceToHost));

  // Print results
  for (int i=0; i<N_array; i++)
  {
    printf("Element %d of parentArray, Result = %f\n", i, host_array[i]);
  }

  // Cleanup
  free(host_array);
  CUDA_CALL(cudaFree(device_array));

}

我得到的结果是：

0 52.000000
1 52.000000
2 52.000000
3 52.000000
4 48.000000
5 48.000000
6 48.000000
7 48.000000
8 48.000000
9 48.000000

我用的是 Cuda 6.5
NVCCFLAGS= -arch=sm_35 -rdc=true -G -O3 --compiler-options -Wall

/opt/cuda-6.5/bin/nvcc -V 

nvcc: NVIDIA (R) Cuda compiler driver    
Copyright (c) 2005-2014 NVIDIA Corporation    
Built on Thu_Jul_17_21:41:27_CDT_2014   
Cuda compilation tools, release 6.5, V6.5.12

Answer 1

此时您将启动 10 个内核（每个子内核也有 10 个线程），10 个活动父内核线程中的每一个：

childKernel<<<1,10>>>(childArr,parentArray[idx],N);

这 10 个内核将运行以任何顺序排列，彼此完全异步。此外，这 10 个内核中的每一个 都试图将值写入 childArr 中相同的 10 个位置。所以这是一个竞争条件。此时最终结果在childArr：

__syncthreads();

将无法预测。

避免竞争条件的一种可能方法是让每个子内核写入 childArr 的单独部分。

另一个问题是在内核中使用 __syncthreads() 而不是 cudaDeviceSynchronize() 作为屏障。内核启动，无论是来自主机还是设备代码，都是异步的，并且 __syncthreads() 不保证异步启动的先前工作已完成。 cudaDeviceSynchronize() 导致调用线程暂停，直到该线程启动的所有先前内核都完成。（见下面的注释）

通过这两项更改，您的代码可以产生您期望的输出：

$ cat t11.cu
#include <stdio.h>
#define CUDA_CALL(x) x
#define MY_M 10
#define MY_N 10

__device__ double childArr[MY_M*MY_N];

__device__ double summe(double *arr, int size)
{
  double result = 0.0;
  for(int i = 0; i < size; i++)
  {
    result += arr[i];
  }
  return result;
}

__device__ double getElement(double arrElement)
{
  return arrElement;
}

__global__ void childKernel(double *arr, double arrElement,int N)
{
  int cidx = blockIdx.x * blockDim.x + threadIdx.x;
  if (cidx < N)
  {
    arr[cidx] = getElement(arrElement);
  }
}

__global__ void parentKernel(double *parentArray, int N)
{
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  if (idx < N)
  {
    childKernel<<<1,MY_M>>>(childArr+MY_M*idx,parentArray[idx],N);
    cudaDeviceSynchronize();
    parentArray[idx] = summe(childArr+MY_M*idx,MY_M);

  }

}

int main(void)
    {

      double *host_array;
      double *device_array;

      // Number of elements in arrays
      const int N_array = MY_N;

      // size of array
      const size_t size_array = N_array * sizeof(double);

      // Allocate array on host
      host_array = (double *)malloc(size_array);

      // Allocate array on device
      CUDA_CALL(cudaMalloc((void **) &device_array, size_array));

      // Initialize host array
      for (int i=0; i<N_array; i++)
      {
        host_array[i] = (double)i;
      }

      // and copy it to CUDA device
      CUDA_CALL(cudaMemcpy(device_array, host_array, size_array, cudaMemcpyHostToDevice));

      // Do calculation on device:
      int block_size = 4;
      // if N = 10, then n_blocks = 3
      int n_blocks = N_array/block_size + (N_array % block_size == 0 ? 0:1);

      parentKernel<<<n_blocks, block_size>>>(device_array,N_array);

      // Retrieve result from device and store it in host array
      CUDA_CALL(cudaMemcpy(host_array, device_array, sizeof(double)*N_array, cudaMemcpyDeviceToHost));

  // Print results
  for (int i=0; i<N_array; i++)
  {
    printf("Element %d of parentArray, Result = %f\n", i, host_array[i]);
  }

  // Cleanup
  free(host_array);
  CUDA_CALL(cudaFree(device_array));

}


$ nvcc -arch=sm_52 -rdc=true -o t11 t11.cu -lcudadevrt
$ cuda-memcheck ./t11
========= CUDA-MEMCHECK
Element 0 of parentArray, Result = 0.000000
Element 1 of parentArray, Result = 10.000000
Element 2 of parentArray, Result = 20.000000
Element 3 of parentArray, Result = 30.000000
Element 4 of parentArray, Result = 40.000000
Element 5 of parentArray, Result = 50.000000
Element 6 of parentArray, Result = 60.000000
Element 7 of parentArray, Result = 70.000000
Element 8 of parentArray, Result = 80.000000
Element 9 of parentArray, Result = 90.000000
========= ERROR SUMMARY: 0 errors
$

注意CDP代码通常要编译：

对于 3.5 或更高的计算能力
使用 -rdc=true 开关（或等效序列，例如 -dc 后跟设备 link）
配合-lcudadevrt开关，拾取设备运行时间库。

注意：事实上，在先前的子内核调用之后，从父线程调用的 cudaDeviceSynchronize() 将暂停该线程，直到所有先前从块中的 任何线程启动的内核 完成。 (documentation) 但是，由于不能保证块中的线程彼此锁步执行，因此其他线程中的哪些内核在特定点启动可能并不明显。因此，正确的用法可能涉及 __syncthreads()（以保证其他线程中的子内核已启动），然后紧接着 cudaDeviceSynchronize() 以保证那些子内核已完成，如果这是所需的行为。但是，在这种特殊情况下，给定父线程的结果 不依赖于 其他父线程子内核的完成，因此我们可以在这种情况下省略 __syncthreads()，只需替换它与 cudaDeviceSynchronize().

Cuda 中的嵌套并行性

Nested parallelism in Cuda

parallel-processing

cuda

gpu

nested-loops