带有数组的结构的 CUDA 数组 (AoSoA)

Question

注意 4
所以代码终于修复了！原来最后一个问题是我正在将分配给每个数组的空间大小添加到 ptr，但 c 已经考虑了变量的大小，所以我基本上添加了 4 倍的空间（以字节为单位）应该有，因此只有 5 元素数组中的前两个元素会显示。 AoSoA 现在可以正常工作了。小心你的内存。如果你尝试类似的管理，我会遇到很多看似愚蠢的错误，因为我的初始代码很草率。

当心：
+ 不正确的偏移量
+ 不必要的 malloc
+ 超出范围的引用

这是工作示例代码，结果如下！

#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;
}

[终于]正确的结果：

第一年
=============

兔子
-------------
地区：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 清理了多个数组索引不一致等问题。

对于 AoSoA 中的前两个位置，结果似乎是正确的 SoA（请参阅新输出）。出于某种原因，第三点 (year 3) 的结果现在给出了错误的结果，尽管 GPU 没有错误代码。我要看看指针（h_pop[year-1].rabbits,h_pop[year-1].hyenas），看看是否有什么发现。

我对其他尝试 AoSoA 的人的唯一建议——非常小心你的索引和内存分配。当然，一般来说这是一个很好的建议，但由于所有内存都在复杂的多级数据容器（如 AoSoA）中四处飞散，如果你马虎，出错的趋势会成倍增加。感谢您的耐心等待，talonmies。

注意 2：
所以按照 talonmies 的建议，我修复了循环 #ing，包装了我的 cuda 调用 w。错误检查并通过重用dev_rabbits/dev_hyenas 压缩了我的cudaMemcpy 调用。还将第一个字母的大小写切换为小写，因为我正在考虑 [djmj][4] 对大小写的抱怨，我意识到 NVIDIA 确实将其常量中的第一个字母小写，所以 [djmj][4] 是是的，从某种意义上说，我应该这样设计我的代码以保持一致性，无论我的个人偏好/经验如何。

通常还清理了代码，因为我写的时候睡得很少，而且有点害怕@它是多么的草率。

现在我遇到了一个新问题……我的程序挂起@第一个cudaMemcpy，并且没有返回（因此talonmies 的方便包装器没有捕获任何东西）。我不太清楚这是为什么...我已经编译了几个程序，包括设备上更大/运行时间更长的程序，它们都运行良好。

此时我很困惑。如果它仍然不起作用，可能会在早上发布一些东西。

注 1
第一个答案似乎真的没有抓住重点。这只是一个玩具代码，并不代表真正的程序。它的唯一目的是尝试设置内存，向其中写入一些垃圾并读回它，以验证 AoSoA 是否正常工作。

因此，就共享内存等向我发表评论不会有成效。这不是这个线程的重点。当然，如果这是一个真正的代码，我会在我的内核中消除分支、使用共享内存、对齐我的数据、使用扭曲级别求和等。我已经在过去的代码中完成了所有这些并让它工作。

此代码是玩具，概念验证代码，仅此而已，旨在尝试使 AoSoA 正常工作。这是它的唯一目的，它不是真正的代码。这是一个概念证明。

至于 var 名称的大小写，我曾在两个不同的地方工作，它们在其编码标准中使用完全大小写的 var 名称（他们使用标签，我在 structs/typedefs 上使用 _s），所以很友好卡住了。对不起，你不喜欢它。至于缩进，我稍后会尝试修复它……Windows 和 Linux 玩得不太好。

还有一点需要注意，如果您对设备指针偏移感到困惑，请在此处查看Anycom 的答案：
Pointers in structs passed to CUDA

我编写了以下代码来测试 CUDA 中包含数组的结构数组....

编辑：固定代码——在meh 之后和hi 之前挂起，大概在cudaMemcpy...不确定原因！

...知道这里发生了什么以及如何解决它吗？

注意： 我担心cudaFrees 可能会把事情搞砸，但删除它们什么也没做。 [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 版本中都可以工作、相同，甚至存在。