如何在内核中动态分配数组？

Question

我需要在内核函数中动态分配一些数组.我怎么能这样做？

我的代码是这样的:

__global__ func(float *grid_d,int n, int nn){  
    int i,j;  
    float x[n],y[nn];  
    //Do some really cool and heavy computations here that takes hours.  
}

但那不行.如果这是在主机代码中我可以使用malloc.cudaMalloc需要主机上的指针,以及设备上的其他指针.在内核函数内部,我没有主机指针.

所以我该怎么做？

如果花费太长时间(几秒钟)来分配所有数组(我需要大约4的大小为n和5大小为nn),这不会是一个问题.因为内核可能至少运行20分钟.

Answer 1

动态内存分配仅支持计算功能2.x和更新的硬件.您可以在内核中使用C++ new关键字或malloc,因此您的示例可能变为:

__global__ func(float *grid_d,int n, int nn){  
    int i,j;  
    float *x = new float[n], *y = new float[nn];   
}

这会在具有上下文生命周期的本地内存运行时堆上分配内存,因此如果您打算不再使用内存,请确保在内核完成运行后释放内存.您还应注意,无法直接从主机API访问运行时堆内存,因此您无法将内核中分配的指针作为参数传递给cudaMemcpy,例如.

Answer 2

@talonmies回答了关于如何在内核中动态分配内存的问题.这是一个补充答案,可以解决__device__ malloc()您可能想要考虑的性能和替代方案.

在内核中动态分配内存可能很诱人,因为它允许GPU代码看起来更像CPU代码.但它会严重影响性能.我写了一个自包含的测试,并将其包含在下面.该测试启动了约260万个线程.每个线程使用从线程索引派生的一些值填充16个全局内存整数,然后对值求和并返回总和.

该测试实现了两种方法.第一种方法使用__device__ malloc(),第二种方法使用在内核运行之前分配的内存.

在我的2.0设备上,内核在使用时运行1500ms,__device__ malloc()在使用预分配内存时运行27ms .换句话说,当在内核中动态分配内存时,测试运行时间要长56倍.时间包括外部循环cudaMalloc()/ cudaFree(),它不是内核的一部分.如果使用相同数量的线程多次启动相同的内核(通常是这种情况),则cudaMalloc()/ 的成本将cudaFree()在所有内核启动时分摊.这使得差异更大,达到60倍左右.

推测,我认为性能损失部分是由隐式序列化引起的.GPU必须序列化所有同时调用__device__ malloc(),以便为每个调用者提供单独的内存块.

不使用的版本__device__ malloc()在运行内核之前分配所有GPU内存.指向内存的指针传递给内核.每个线程计算先前分配的内存的索引,而不是使用a __device__ malloc().

预先分配内存的潜在问题是,如果只有一些线程需要分配内存,并且不知道它们是哪些线程,则需要为所有线程分配内存.如果没有足够的内存,那么减少每个内核调用的线程数然后使用可能会更有效__device__ malloc().其他解决方法可能最终会重新实现 __device__ malloc()后台正在进行的操作,并会看到类似的性能影响.

测试性能__device__ malloc():

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>

const int N_ITEMS(16);

#define USE_DYNAMIC_MALLOC

__global__ void test_malloc(int* totals)
{
  int tx(blockIdx.x * blockDim.x + threadIdx.x);

  int* s(new int[N_ITEMS]);

  for (int i(0); i < N_ITEMS; ++i) {
    s[i] = tx * i;
  }

  int total(0);
  for (int i(0); i < N_ITEMS; ++i) {
    total += s[i];
  }

  totals[tx] = total;

  delete[] s;
}

__global__ void test_malloc_2(int* items, int* totals)
{
  int tx(blockIdx.x * blockDim.x + threadIdx.x);

  int* s(items + tx * N_ITEMS);

  for (int i(0); i < N_ITEMS; ++i) {
    s[i] = tx * i;
  }

  int total(0);
  for (int i(0); i < N_ITEMS; ++i) {
    total += s[i];
  }

  totals[tx] = total;
}

int main()
{
  cudaError_t cuda_status;

  cudaSetDevice(0);

  int blocks_per_launch(1024 * 10);
  int threads_per_block(256);

  int threads_per_launch(blocks_per_launch * threads_per_block);

  int* totals_d;
  cudaMalloc((void**)&totals_d, threads_per_launch * sizeof(int));

  cudaEvent_t start, stop;
  cudaEventCreate(&start);
  cudaEventCreate(&stop);

  cudaDeviceSynchronize();
  cudaEventRecord(start, 0);

#ifdef USE_DYNAMIC_MALLOC
  cudaDeviceSetLimit(cudaLimitMallocHeapSize, threads_per_launch * N_ITEMS * sizeof(int));

  test_malloc<<<blocks_per_launch, threads_per_block>>>(totals_d);
#else
  int* items_d;
  cudaMalloc((void**)&items_d, threads_per_launch * sizeof(int) * N_ITEMS);

  test_malloc_2<<<blocks_per_launch, threads_per_block>>>(items_d, totals_d);

  cudaFree(items_d);
#endif

  cuda_status = cudaDeviceSynchronize();
  if (cuda_status != cudaSuccess) {
    printf("Error: %d\n", cuda_status);
    exit(1);
  }

  cudaEventRecord(stop, 0);
  cudaEventSynchronize(stop);
  float elapsedTime;
  cudaEventElapsedTime(&elapsedTime, start, stop);

  printf("Elapsed: %f\n", elapsedTime);

  int* totals_h(new int[threads_per_launch]);
  cuda_status = cudaMemcpy(totals_h, totals_d, threads_per_launch * sizeof(int), cudaMemcpyDeviceToHost);
  if (cuda_status != cudaSuccess) {
    printf("Error: %d\n", cuda_status);
    exit(1);
  }

  for (int i(0); i < 10; ++i) {
    printf("%d ", totals_h[i]);
  }
  printf("\n");

  cudaFree(totals_d);
  delete[] totals_h;

  return cuda_status;
}

输出:

C:\rd\projects\test_cuda_malloc\Release>test_cuda_malloc.exe
Elapsed: 27.311169
0 120 240 360 480 600 720 840 960 1080

C:\rd\projects\test_cuda_malloc\Release>test_cuda_malloc.exe
Elapsed: 1516.711914
0 120 240 360 480 600 720 840 960 1080