动态并行的循环展开会降低时间性能

Question

我有一个简单的程序来计算平方根，循环展开是这样的

循环展开

 #include <stdio.h>
#include <cuda.h>
__global__ void square(float *a, int N,int idx);


// Kernel that executes on the CUDA device
__global__ void first(float *arr, int N)
{
  int idx = 2*(blockIdx.x * blockDim.x + threadIdx.x);
  int n=N;
  //printf("%d\n",n);
  for(int q=0;q<2;q++)
  {
  if(N<2000)
  {
  arr[idx+q] = arr[idx+q] * arr[idx+q];
  }
  }

}



// main routine that executes on the host
int main(void)
{
  clock_t start = clock(),diff;
  float *a_h, *a_d;  // Pointer to host & device arrays
  const int N = 1000;  // Number of elements in arrays
  size_t size = N * sizeof(float);
  a_h = (float *)malloc(size);        // Allocate array on host
  cudaMalloc((void **) &a_d, size);   // Allocate array on device
  // Initialize host array and copy it to CUDA device
  for (int i=0; i<N; i++) a_h[i] = (float)i;
  cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice);
  // Do calculation on device:
  int block_size = 4;
  //int n_blocks = N/block_size + (N%block_size == 0 ? 0:1);
  first <<< 4, 128 >>> (a_d, N);
  //cudaThreadSynchronize();
  // Retrieve result from device and store it in host array
  cudaMemcpy(a_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost);
  // Print results
  for (int i=0; i<N; i++) printf("%d %f\n", i, a_h[i]);
  // Cleanup
  free(a_h); cudaFree(a_d);
  diff = clock() - start;
 int msec = diff * 1000 / CLOCKS_PER_SEC;

 printf("Time taken %d seconds %d milliseconds\n", msec/1000, msec%1000);

}

然后意识到循环计算可以通过动态并行最小化。

动态并行展开被实现为

动态并行展开

#include <stdio.h>
#include <cuda.h>
__global__ void square(float *a, int N,int idx);


// Kernel that executes on the CUDA device
__global__ void first(float *arr, int N)
{
  int idx = 2*(blockIdx.x * blockDim.x + threadIdx.x);
  int n=N;
  square <<< 1,2 >>> (arr, n,idx);


}

__global__ void square(float *a, int N,int idx)
{
  int tdx = blockIdx.x * blockDim.x + threadIdx.x;
  printf("%d\n",N);
  if(N<2000)
  {
  a[tdx+idx] = a[tdx+idx] * a[tdx+idx];
  }
}

// main routine that executes on the host
int main(void)
{
  clock_t start = clock(),diff;
  float *a_h, *a_d;  // Pointer to host & device arrays
  const int N = 1000;  // Number of elements in arrays
  size_t size = N * sizeof(float);
  a_h = (float *)malloc(size);        // Allocate array on host
  cudaMalloc((void **) &a_d, size);   // Allocate array on device
  // Initialize host array and copy it to CUDA device
  for (int i=0; i<N; i++) a_h[i] = (float)i;
  cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice);
  // Do calculation on device:
  int block_size = 4;
  //int n_blocks = N/block_size + (N%block_size == 0 ? 0:1);
  first <<< 4, 128 >>> (a_d, N);
  //cudaThreadSynchronize();
  // Retrieve result from device and store it in host array
  cudaMemcpy(a_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost);
  // Print results
  for (int i=0; i<N; i++) printf("%d %f\n", i, a_h[i]);
  // Cleanup
  free(a_h); cudaFree(a_d);
  diff = clock() - start;
 int msec = diff * 1000 / CLOCKS_PER_SEC;

 printf("Time taken %d seconds %d milliseconds\n", msec/1000, msec%1000);

}

通过展开实现动态并行性比仅展开需要更多的执行时间。在这种情况下，我们不应该通过动态并行来提高执行时间吗？

Answer 1

动态并行主要在您具有动态并行的情况下很有用。那就是：在进行一些计算之前，您不知道需要多少并行度。您无需将数据传输回主机，然后立即将其输入参数化另一个启动，而是从内核中启动。在这种模式下，由于避免了内核启动之间的 memcpy，您会看到加速。

在您上面的示例中，情况并非如此。您可能刚刚从主机启动了两倍的线程。不需要任何动态，因为那里没有可用的并行性，这是您在第一次内核启动时不知道的。

此外，使用动态并行启动的内核的性能要求与从主机启动的内核相似。您必须启动合理数量的工作，否则启动延迟将支配您的计算时间。