Распределение памяти на GPU для динамического массива структур

У меня проблема с передачей массива структуры в ядро ​​GPU. Я основывался на этой теме - ошибка сегментации cudaMemcpy, и я написал что-то вроде этого:

#include <stdio.h>
#include <stdlib.h>

struct Test {
    char *array;
};

__global__ void kernel(Test *dev_test) {
    for(int i=0; i < 5; i++) {
        printf("Kernel[0][i]: %c \n", dev_test[0].array[i]);
    }
}

int main(void) {

    int n = 4, size = 5;
    Test *dev_test, *test;

    test = (Test*)malloc(sizeof(Test)*n);
    for(int i = 0; i < n; i++)
        test[i].array = (char*)malloc(size * sizeof(char));

    for(int i=0; i < n; i++) {
        char temp[] = { 'a', 'b', 'c', 'd' , 'e' };
        memcpy(test[i].array, temp, size * sizeof(char));
    }

    cudaMalloc((void**)&dev_test, n * sizeof(Test));
    cudaMemcpy(dev_test, test, n * sizeof(Test), cudaMemcpyHostToDevice);
    for(int i=0; i < n; i++) {
        cudaMalloc((void**)&(test[i].array), size * sizeof(char));
        cudaMemcpy(&(dev_test[i].array), &(test[i].array), size * sizeof(char), cudaMemcpyHostToDevice);
    }

    kernel<<<1, 1>>>(dev_test);
    cudaDeviceSynchronize();

    //  memory free
    return 0;
}

Ошибка отсутствует, но отображаемые значения в ядре неверны. Что я делаю не так? Заранее благодарю за любую помощь.

Источник

2 ответа

Решение

  1. Это выделяет новый указатель на память хоста:

    test[i].array = (char*)malloc(size * sizeof(char));

  2. Это копирование данных в этот регион в памяти хоста:

    memcpy(test[i].array, temp, size * sizeof(char));

  3. Это перезаписывает ранее выделенный указатель на память хоста (с шага 1 выше) новым указателем на память устройства:

    cudaMalloc((void**)&(test[i].array), size * sizeof(char));

После шага 3 данные, настроенные на шаге 2, полностью теряются и больше не доступны. Ссылаясь на шаги 3 и 4 в вопросе / ответе, который вы связали:

3. Создайте отдельный указатель int на хост, давайте назовем его myhostptr

4.cudaMalloc int хранилище на устройстве для myhostptr

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

Вот модифицированная версия вашего кода, которая правильно реализует пропущенные шаги 3 и 4 (и 5), которые вы не выполнили правильно в соответствии с вопросом / ответом, который вы связали: (см. Комментарии, описывающие шаги 3,4,5)

$ cat t755.cu
#include <stdio.h>
#include <stdlib.h>

struct Test {
    char *array;
};

__global__ void kernel(Test *dev_test) {
    for(int i=0; i < 5; i++) {
        printf("Kernel[0][i]: %c \n", dev_test[0].array[i]);
    }
}

int main(void) {

    int n = 4, size = 5;
    Test *dev_test, *test;

    test = (Test*)malloc(sizeof(Test)*n);
    for(int i = 0; i < n; i++)
        test[i].array = (char*)malloc(size * sizeof(char));

    for(int i=0; i < n; i++) {
        char temp[] = { 'a', 'b', 'c', 'd' , 'e' };
        memcpy(test[i].array, temp, size * sizeof(char));
    }

    cudaMalloc((void**)&dev_test, n * sizeof(Test));
    cudaMemcpy(dev_test, test, n * sizeof(Test), cudaMemcpyHostToDevice);

    // Step 3:
    char *temp_data[n];
    // Step 4:
    for (int i=0; i < n; i++)
      cudaMalloc(&(temp_data[i]), size*sizeof(char));
    // Step 5:
    for (int i=0; i < n; i++)
      cudaMemcpy(&(dev_test[i].array), &(temp_data[i]), sizeof(char *), cudaMemcpyHostToDevice);
    // now copy the embedded data:
    for (int i=0; i < n; i++)
      cudaMemcpy(temp_data[i], test[i].array, size*sizeof(char), cudaMemcpyHostToDevice);

    kernel<<<1, 1>>>(dev_test);
    cudaDeviceSynchronize();

    //  memory free
    return 0;
}

$ nvcc -o t755 t755.cu
$ cuda-memcheck ./t755
========= CUDA-MEMCHECK
Kernel[0][i]: a
Kernel[0][i]: b
Kernel[0][i]: c
Kernel[0][i]: d
Kernel[0][i]: e
========= ERROR SUMMARY: 0 errors
$

Поскольку приведенная выше методология может быть сложной для начинающих, обычный совет не делать этого, а вместо этого выравнивать структуры данных. Сглаживание обычно означает изменение порядка хранения данных, чтобы удалить встроенные указатели, которые должны быть выделены отдельно.

Тривиальным примером выравнивания этой структуры данных будет использование этого вместо этого:

struct Test {
    char array[5];
};

Конечно, признается, что этот конкретный подход не будет служить всем целям, но он должен иллюстрировать общую идею / намерение. С этой модификацией, например, код становится намного проще:

$ cat t755.cu
#include <stdio.h>
#include <stdlib.h>

struct Test {
    char array[5];
};

__global__ void kernel(Test *dev_test) {
    for(int i=0; i < 5; i++) {
        printf("Kernel[0][i]: %c \n", dev_test[0].array[i]);
    }
}

int main(void) {

    int n = 4, size = 5;
    Test *dev_test, *test;

    test = (Test*)malloc(sizeof(Test)*n);

    for(int i=0; i < n; i++) {
        char temp[] = { 'a', 'b', 'c', 'd' , 'e' };
        memcpy(test[i].array, temp, size * sizeof(char));
    }

    cudaMalloc((void**)&dev_test, n * sizeof(Test));
    cudaMemcpy(dev_test, test, n * sizeof(Test), cudaMemcpyHostToDevice);

    kernel<<<1, 1>>>(dev_test);
    cudaDeviceSynchronize();

    //  memory free
    return 0;
}
$ nvcc -o t755 t755.cu
$ cuda-memcheck ./t755
========= CUDA-MEMCHECK
Kernel[0][i]: a
Kernel[0][i]: b
Kernel[0][i]: c
Kernel[0][i]: d
Kernel[0][i]: e
========= ERROR SUMMARY: 0 errors
$
Источник

Спасибо @Robert Crovella за ответ. Приведенный выше ответ был для меня очень полезен. Я обновил код с использования массива символов на использование целочисленного массива структур для справки. В ядре значения обновляются и возвращаются обратно на хост.

        #include <stdio.h>
  #include <stdlib.h>

  struct Test {
  int *array;
  };

  __global__ void kernel(Test *dev_test) {
  //    for(int i=0; i < 4; i++)
  //      for(int j=0; j < 5; j++) {
  //        printf("Kernel[X][i]: %d \n", dev_test[i].array[j]);
  //    }
  int i =threadIdx.x + blockIdx.x*blockDim.x;
  int j = threadIdx.y + blockIdx.y*blockDim.y;

  if(i<4 && j<5)
      dev_test[i].array[j] *= 2; 
  }

  int main(void) {

  int n = 4, size = 5;
  Test *dev_test, *test;

  test = (Test*)malloc(sizeof(Test)*n);
  for(int i = 0; i < n; i++)
      test[i].array = (int*)malloc(size * sizeof(int));

  for(int i=0;i<n;i++)
      for(int j=0; j<size; j++)
          test[i].array[j] = i*n+j;

  cudaMalloc((void**)&dev_test, n * sizeof(Test));
  cudaMemcpy(dev_test, test, n * sizeof(Test), cudaMemcpyHostToDevice);

  // Step 3:
  int *temp_data[n];
  // Step 4:
  for (int i=0; i < n; i++)
      cudaMalloc(&(temp_data[i]), size*sizeof(int));
  // Step 5:
  for (int i=0; i < n; i++)
      cudaMemcpy(&(dev_test[i].array), &(temp_data[i]), sizeof(int *), cudaMemcpyHostToDevice);
  // now copy the embedded data:
  for (int i=0; i < n; i++)
      cudaMemcpy(temp_data[i], test[i].array, size*sizeof(int), cudaMemcpyHostToDevice);

  dim3 threads(4,5);
  kernel<<<1, threads>>>(dev_test);
  cudaDeviceSynchronize();

  for(int i=0;i<n;i++)
      cudaMemcpy(test[i].array, temp_data[i], size*sizeof(int), cudaMemcpyDeviceToHost);

  for(int i=0;i <n;i++){
      for(int j=0;j<size;j++)
          printf("%d ",test[i].array[j]);
  printf("\n");
  }
  //  memory free
  return 0;
  }
Источник
Другие вопросы по тегам