Как динамически выделять массивы внутри ядра?

Ответ 1

Распределение динамической памяти поддерживается только для возможности вычисления 2.x и более нового оборудования. Вы можете использовать либо новое ключевое слово С++, либо malloc в ядре, поэтому ваш пример может стать:

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

Это выделяет память в куче времени хранения в локальной памяти, которая имеет время жизни контекста, поэтому убедитесь, что вы освобождаете память после завершения работы ядра, если ваше намерение не использовать память снова. Вы также должны заметить, что куча памяти времени выполнения не может быть доступна напрямую из API-интерфейсов хоста, поэтому вы не можете передать указатель, выделенный внутри ядра, в качестве аргумента для cudaMemcpy, например.

Ответ 2

@talonmies ответили на ваш вопрос о том, как динамически распределять память в ядре. Это предназначено в качестве дополнительного ответа, отвечающего за производительность __device__ malloc() и альтернативы, которые вы, возможно, захотите рассмотреть.

Динамическое распределение памяти в ядре может быть заманчивым, поскольку позволяет графическому процессору больше походить на код ЦП. Но это может серьезно повлиять на производительность. Я написал собственный тест и включил его ниже. Тест запускает около 2,6 миллиона потоков. Каждый поток заполняет 16 целых чисел глобальной памяти с некоторыми значениями, полученными из индекса потока, затем суммирует значения и возвращает сумму.

Тест реализует два подхода. В первом подходе используется __device__ malloc(), а во втором подходе используется память, выделенная до запуска ядра.

На моем устройстве 2.0 ядро ​​запускается в 1500 мс при использовании __device__ malloc() и 27ms при использовании предварительно выделенной памяти. Другими словами, для выполнения теста в 56x больше выполняется при динамическом распределении памяти в ядре. Время включает в себя внешний цикл cudaMalloc()/cudaFree(), который не является частью ядра. Если одно и то же ядро ​​запускается много раз с тем же количеством потоков, как это часто бывает, стоимость cudaMalloc()/cudaFree() амортизируется по всем запускам ядра. Это приносит разницу даже выше, примерно в 60 раз.

Говоря, я думаю, что поражение производительности частично вызвано неявной сериализацией. Графический процессор должен, вероятно, сериализовать все одновременные вызовы на __device__ malloc(), чтобы предоставить отдельным кускам памяти каждому вызывающему абоненту.

Версия, не использующая __device__ malloc(), распределяет всю память GPU перед запуском ядра. Указатель на память передается ядру. Каждый поток вычисляет индекс в ранее выделенную память вместо использования __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

Ответ 3

Если значение n и nn было известно до вызова ядра, то почему бы не cudaMalloc память на стороне хоста и передать указатель памяти устройства на ядро?