c - 带有数组的结构的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;
}

[最后]正确的结果：


  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清除了我的代码中的多个数组索引不一致等情况。

对于AoSoA中的前两个位置，结果看起来似乎是正确的SoA（请参阅新输出）。由于某种原因，尽管GPU没有错误代码，但第三点（ year 3）的结果现在给出了错误的结果。我将偷看指针（ 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正常工作。那只是它的唯一目的，不是真正的代码。这是概念的证明。

至于var名称的大小写，我在两个不同的地方工作过，它们在其编码标准中使用了全大小写的var名称（它们使用的标记，我在structs / typedefs上使用 _s），因此有点卡住了。抱歉，您不喜欢它。至于缩进，我稍后会尝试解决。Windows和Linux的表现不佳。

还有一点需要注意的是，如果您对设备指针偏移感到困惑，请参见 Anycom的答案：
Pointers in structs passed to 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版本中都能正常工作，相同甚至存在。