CUDA Массив структур с массивами (AoSoA)

Question

Примечание 4
Итак, код окончательно исправлен! Оказалось, что последняя проблема заключалась в том, что я добавлял размер пространства, выделенного каждому массиву, в ptr, но c уже учитывает размер переменной, поэтому я, по сути, добавил в 4 раза больше места в байтах, чем следовало бы был, следовательно, будут отображаться только первые два элемента в массиве из 5 элементов. AoSoA теперь полностью работает. Будьте осторожны с вашей запиской управление, если вы попробуете что-то подобное, я боролся с множеством, казалось бы, глупых ошибок, потому что мой исходный код был неаккуратным.

Осторожно:
+ Неправильные смещения
+ Ненужный Маллок
+ Недопустимые ссылки

Вот рабочий пример кода, результаты следуют!

#include <stdio.h>

#define REGIONS 20
#define YEARS 5

__inline __host__ void gpuAssert(cudaError_t code, char *file, int line, 
                 bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code),
          file, line);
      if (abort) exit(code);
   }
}

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }

struct AnimalPopulationForYear_s
{
   bool isYearEven;
   int * rabbits;
   int * hyenas;
};

AnimalPopulationForYear_s * dev_pop;

__global__ void RunSim(AnimalPopulationForYear_s dev_pop[],
               int year)
{
   int idx = blockIdx.x*blockDim.x+threadIdx.x;
   int rabbits, hyenas;
   int arrEl = year-1;

   rabbits = (idx+1) * year * year; 
   hyenas = rabbits / 10;

   if ( rabbits > 100000 ) rabbits = 100000;   
   if ( hyenas < 2 ) hyenas = 2;

   if ( idx < REGIONS ) dev_pop[arrEl].rabbits[idx] = rabbits;
   if ( idx < REGIONS ) dev_pop[arrEl].hyenas[idx] = hyenas;

   if (threadIdx.x == 0 && blockIdx.x == 0)
      dev_pop[arrEl].isYearEven = (year & 0x01 == 0x0);
}

int main()
{
   //Various reused sizes...
   const size_t fullArrSz = size_t(YEARS) * size_t(REGIONS) * sizeof(int);
   const size_t structArrSz = size_t(YEARS) * sizeof(AnimalPopulationForYear_s);

   //Vars to hold struct and merged subarray memory inside it.
   AnimalPopulationForYear_s * h_pop;
   int * dev_hyenas, * dev_rabbits, * h_hyenas, * h_rabbits, arrEl;

   //Alloc. memory.
   h_pop = (AnimalPopulationForYear_s *) malloc(structArrSz);
   h_rabbits = (int *) malloc(fullArrSz);
   h_hyenas = (int *) malloc(fullArrSz);
   gpuErrchk(cudaMalloc((void **) &dev_pop,structArrSz));
   gpuErrchk(cudaMalloc((void **) &dev_rabbits,fullArrSz));
   gpuErrchk(cudaMalloc((void **) &dev_hyenas,fullArrSz));

   //Offset ptrs.
   for (int i = 0; i < YEARS; i++)
   {
      h_pop[i].rabbits = dev_rabbits+i*REGIONS;
      h_pop[i].hyenas = dev_hyenas+i*REGIONS;
   }

   //Copy host struct with dev. pointers to device.
   gpuErrchk
      (cudaMemcpy(dev_pop,h_pop, structArrSz, cudaMemcpyHostToDevice));

   //Call kernel
   for(int i=1; i < YEARS+1; i++) RunSim<<<REGIONS/128+1,128>>>(dev_pop,i);

   //Make sure nothing went wrong.
   gpuErrchk(cudaPeekAtLastError());
   gpuErrchk(cudaDeviceSynchronize());

   gpuErrchk(cudaMemcpy(h_pop,dev_pop,structArrSz, cudaMemcpyDeviceToHost));
   gpuErrchk
      (cudaMemcpy(h_rabbits, dev_rabbits,fullArrSz, cudaMemcpyDeviceToHost));
   gpuErrchk(cudaMemcpy(h_hyenas,dev_hyenas,fullArrSz, cudaMemcpyDeviceToHost));

   for(int i=0; i < YEARS; i++)
   {
      h_pop[i].rabbits = h_rabbits + i*REGIONS;
      h_pop[i].hyenas = h_hyenas + i*REGIONS;
   }

   for(int i=1; i < YEARS+1; i++)
   {
      arrEl = i-1;
      printf("\nYear %i\n=============\n\n", i);      
      printf("Rabbits\n-------------\n");
      for (int j=0; j < REGIONS; j++)
     printf("Region: %i  Pop: %i\n", j, h_pop[arrEl].rabbits[j]);;      
      printf("Hyenas\n-------------\n");
      for (int j=0; j < REGIONS; j++)
     printf("Region: %i  Pop: %i\n", j, h_pop[arrEl].hyenas[j]);
   }

   //Free on device and host
   cudaFree(dev_pop);
   cudaFree(dev_rabbits);
   cudaFree(dev_hyenas);

   free(h_pop);
   free(h_rabbits);
   free(h_hyenas);

   return 0;
}

[Наконец] правильные результаты:

Год 1
=============

Кролики
-------------
Регион: 0 Поп: 1
Регион: 1 Поп: 2
Регион: 2 Поп: 3
Регион: 3 Поп: 4
Регион: 4 Поп: 5
Регион: 5 Поп: 6
Регион: 6 Поп: 7
Регион: 7 Поп: 8
Регион: 8 Поп: 9
Регион: 9 Поп: 10
Регион: 10 Поп: 11
Регион: 11 Поп: 12
Регион: 12 Поп: 13
Регион: 13 Поп: 14
Регион: 14 Поп: 15
Регион: 15 Поп: 16
Регион: 16 Поп: 17
Регион: 17 Поп: 18
Регион: 18 Поп: 19
Регион: 19 Поп: 20
Гиены
-------------
Регион: 0 Поп: 2
Регион: 1 Поп: 2
Регион: 2 Поп: 2
Регион: 3 Поп: 2
Регион: 4 Поп: 2
Регион: 5 Поп: 2
Регион: 6 Поп: 2
Регион: 7 Поп: 2
Регион: 8 Поп: 2
Регион: 9 Поп: 2
Регион: 10 Поп: 2
Регион: 11 Поп: 2
Регион: 12 Поп: 2
Регион: 13 Поп: 2
Регион: 14 Поп: 2
Регион: 15 Поп: 2
Регион: 16 Поп: 2
Регион: 17 Поп: 2
Регион: 18 Поп: 2
Регион: 19 Поп: 2

Год 2
=============

Кролики
-------------
Регион: 0 Поп: 4
Регион: 1 Поп: 8
Регион: 2 Поп: 12
Регион: 3 Поп: 16
Регион: 4 Поп: 20
Регион: 5 Поп: 24
Регион: 6 Поп: 28
Регион: 7 Поп: 32
Регион: 8 Поп: 36
Регион: 9 Поп: 40
Регион: 10 Поп: 44
Регион: 11 Поп: 48
Регион: 12 Поп: 52
Регион: 13 Поп: 56
Регион: 14 Поп: 60
Регион: 15 Поп: 64
Регион: 16 Поп: 68
Регион: 17 Поп: 72
Регион: 18 Поп: 76
Регион: 19 Поп: 80
Гиены
-------------
Регион: 0 Поп: 2
Регион: 1 Поп: 2
Регион: 2 Поп: 2
Регион: 3 Поп: 2
Регион: 4 Поп: 2
Регион: 5 Поп: 2
Регион: 6 Поп: 2
Регион: 7 Поп: 3
Регион: 8 Поп: 3
Регион: 9 Поп: 4
Регион: 10 Поп: 4
Регион: 11 Поп: 4
Регион: 12 Поп: 5
Регион: 13 Поп: 5
Регион: 14 Поп: 6
Регион: 15 Поп: 6
Регион: 16 Поп: 6
Регион: 17 Поп: 7
Регион: 18 Поп: 7
Регион: 19 ...

Примечание 3:
После talonmies исправлено несколько несоответствий индексации массивов и т. Д. В моем коде.

Результаты выглядят как правильные SoA для первых двух точек в AoSoA (см. Новый вывод). По какой-то причине результаты с третьего места (year 3) теперь дают неверные результаты, хотя код ошибки от GPU отсутствует. Я собираюсь посмотреть на указатели (h_pop[year-1].rabbits, h_pop[year-1].hyenas) и посмотреть, покажет ли это что-нибудь.

Мой единственный совет для всех, кто пытается AoSoA - быть ОЧЕНЬ осторожным с индексацией и распределением памяти. Конечно, это хороший совет в целом, но, когда вся память летит в сложном многоуровневом контейнере данных, таком как AoSoA, тенденция к ошибкам, если вы небрежны, экспоненциально возрастает. Спасибо за ваше терпение, talonmies .

Примечание 2:
Поэтому, следуя совету talonmies , я исправил свой цикл #ing, обернул свои cuda-вызовы w. проверка ошибок и сжала мои cudaMemcpy звонки путем повторного использования dev_rabbits / dev_hyenas. Также переключил регистр на первую букву в нижнем регистре, так как я думал о жалобе [djmj] [4] на регистр, и я понял, что NVIDIA делает строчные буквы первой буквы в своих константах, поэтому [djmj] [4] Правильно, в некотором смысле, я должен был сделать свой код таким, чтобы он был согласованным, независимо от моих личных предпочтений / опыта.

Также, как правило, очищал код, так как я написал, что он не спал и был немного испуган @ как неряшливо это было.

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

На данный момент я озадачен. Если он все еще не работает, возможно, отправьте что-нибудь утром.

Примечание 1
Первый ответ, казалось, действительно упускал суть. Это всего лишь игрушечный код, он не предназначен для представления реальной программы. Его единственная цель - попытаться настроить память, записать в нее немного мусора и прочитать его обратно, чтобы убедиться, что AoSoA работает.

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

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

Что касается размещения имен переменных, я работал в двух разных местах, которые использовали полностью покрытые имена переменных в своем стандарте кодирования (они использовали теги, я делаю _s для structs / typedefs), так что застрял. Извините, вам не нравится это. Что касается отступа, я постараюсь исправить это позже ... Windows и Linux не играли хорошо.

Еще одно замечание: если вас смущает смещение указателя устройства, см. Ответ Anycom здесь:
Указатели в структурах переданы CUDA

Я написал следующий код для проверки массивов структур с массивами внутри в CUDA ....

Редактировать: Исправлен код - зависает после meh и до hi, предположительно на cudaMemcpy ... не знаю почему!

... Есть идеи, что здесь происходит и как это исправить?

Примечание: Я волновался, что cudaFree могут испортить вещи, но удаление их ничего не сделало. [4]:

Answer 1

В этом коде очень много ошибок, но основная причина "искаженных" результатов, о которых вы спрашиваете, заключается в том, что вы смотрите на неинициализированную память.dev_Pop[0].Rabbits никогда не устанавливается на что-либо в памяти устройства, поэтому вы не должны быть слишком удивлены, что его содержимое «искажено».Основная причина проблемы:

for(int i=1; i < YEARS+1; i++)
    RunSim<<<REGIONS/128+1,128>>>(dev_Pop,i);

Здесь вы начинаете с year=1, что означает, что year=0 никогда не устанавливается на что-либо, а year=YEARS - это гарантированное переполнение буфера в памяти устройства.

Позже в коде обратного копирования вы делаете это на каждой итерации:

cudaFree(h_Pop[i].Rabbits);
cudaFree(h_Pop[i].Hyenas);

, но в первую очередь вы никогда не использовали их неправильно, так что операция копирования, вероятно, также будет неудачной.Трудно сказать, как он потерпит неудачу, без компиляции и запуска кода, но я бы предположил , что среда выполнения CUDA полностью освободит dev_Rabbits и dev_Hyenas при первом вызове.Это должно привести к сбою последующих вызовов cudaMemcpy в цикле.Независимо от точной механики, я был бы невероятно удивлен, если бы ваш цикл обратного копирования успешно вернул все данные на хост.Гораздо более разумная реализация была бы похожей на код, который вы использовали для создания образа памяти устройства, например:

const size_t dsize = size_t(YEARS) * size_t(REGIONS) * sizeof(int);
int * Rabbits = (int *) malloc(dsize);
int * Hyenas = (int *) malloc(dsize);
cudaMemcpy(Rabbits, dev_Rabbits, dsize, cudaMemcpyDeviceToHost);
cudaMemcpy(Hyenas, dev_Hyenas, dsize, cudaMemcpyDeviceToHost);

for(int i=0; i < YEARS; i++)
{
    h_Pop[i].Rabbits = Rabbits + i*REGIONS;
    h_Pop[i].Hyenas = Hyenas + i*REGIONS;
}

Таким образом избавляется от большого количества избыточных устройств.-> транзакции хоста по шине PCI-e и все эти ненужные вызовы malloc на стороне хоста в цикле.

Так что я бы предположил, что в коде происходит множество точек сбоя во время выполнения, но потому чтовы забыли включить любую проверку ошибок, все происходит молча, а вы просто этого не замечаете.Чтобы исправить это, добавьте что-то вроде этого в ваш код:

inline void gpuAssert(cudaError_t code, char * file, int line, bool Abort=true)
{
    if (code != cudaSuccess) {
        fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code),file,line);
        if (Abort) exit(code);
    }       
}

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }

, а затем используйте gpuErrchk для проверки состояния возврата каждого вызова API, например:

gpuErrchk(cudaMalloc((void **) &dev_Pop,YEARS*sizeof(AnimalPopulationForYear_s)));

Для запуска вашего ядра я рекомендую сделать следующее:

RunSim<<<REGIONS/128+1,128>>>(dev_Pop,i);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());

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

---

РЕДАКТИРОВАТЬ:

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

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

#include <cstdio>
#include <cstdlib>

#define REGIONS 20
#define YEARS 5
#define POPMIN 2
#define POPMAX 100000

inline void gpuAssert(cudaError_t code, char *file, int line, 
                 bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code),
          file, line);
      if (abort) exit(code);
   }
}

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }

struct Population_s
{
   int * rabbits;
   int * hyenas;
};

__global__ void RunSim(Population_s * dev_pop, int year)
{
   int idx = blockIdx.x*blockDim.x+threadIdx.x;

   if (idx < REGIONS) {
      int rabbits, hyenas;

      rabbits = min(POPMAX, idx * year * year); 
      hyenas = max(POPMIN, rabbits / 10);

      dev_pop[year-1].rabbits[idx] = rabbits;
      dev_pop[year-1].hyenas[idx] = hyenas;
   }
}

int main()
{
   const size_t subArrSz = size_t(REGIONS) * sizeof(int);
   const size_t fullArrSz = size_t(YEARS) * subArrSz;
   const size_t structArrSz = size_t(YEARS) * sizeof(Population_s);

   Population_s * h_pop = (Population_s *) malloc(structArrSz);
   int * h_rabbits = (int *) malloc(fullArrSz);
   int * h_hyenas = (int *) malloc(fullArrSz);

   Population_s * dev_pop;
   int * dev_hyenas, * dev_rabbits;

   gpuErrchk(cudaMalloc((void **) &dev_pop,structArrSz));
   gpuErrchk(cudaMalloc((void **) &dev_hyenas,fullArrSz));
   gpuErrchk(cudaMalloc((void **) &dev_rabbits,fullArrSz));

   gpuErrchk(cudaMemset(dev_rabbits, 1, fullArrSz));
   gpuErrchk(cudaMemset(dev_hyenas, 1, fullArrSz));

   for (int i = 0; i < YEARS; i++)
   {
      h_pop[i].rabbits = dev_rabbits + i*REGIONS;
      h_pop[i].hyenas = dev_hyenas + i*REGIONS;
   }

   gpuErrchk
      (cudaMemcpy(dev_pop,h_pop, structArrSz, cudaMemcpyHostToDevice));

   for(int i = 1; i < (YEARS+1); i++) {
       RunSim<<<REGIONS/128+1,128>>>(dev_pop,i);
       gpuErrchk(cudaPeekAtLastError());
       gpuErrchk(cudaDeviceSynchronize());
   }

   gpuErrchk(cudaMemcpy(h_rabbits, dev_rabbits, fullArrSz, cudaMemcpyDeviceToHost));
   gpuErrchk(cudaMemcpy(h_hyenas, dev_hyenas, fullArrSz, cudaMemcpyDeviceToHost));

   for(int i=0; i < YEARS; i++)
   {
      h_pop[i].rabbits = h_rabbits + i*REGIONS;
      h_pop[i].hyenas = h_hyenas + i*REGIONS;
   }

   for(int i=0; i < YEARS; i++)
   {
      printf("\n=============\n");   
      printf("Year %i\n=============\n\n", i+1);   
      printf("Rabbits\n-------------\n", i);
      for (int j=0; j < REGIONS; j++)
         printf("Region: %i  Pop: %i\n", j, h_pop[i].rabbits[j]);;      
      printf("\nHyenas\n-------------\n", i);
      for (int j=0; j < REGIONS; j++)
         printf("Region: %i  Pop: %i\n", j, h_pop[i].hyenas[j]);
   }

   cudaFree(dev_pop);
   cudaFree(dev_rabbits);
   cudaFree(dev_hyenas);

   free(h_pop);
   free(h_rabbits);
   free(h_hyenas);

   return 0;
}

В качестве последнего совета - не используйте ничего из SDK cutilбиблиотека в вашем собственном коде, это не то, для чего она предназначена.Он не является официальной частью CUDA, не имеет документации, не считается готовым к работе и не гарантирует, что он будет работать, быть таким же или даже существовать в любом выпуске CUDA SDK.