Cuda：使用全局内存来存储连续的各种大小的数据答案

【问题标题】：Cuda : use the global memory to store contiguous various size datasCuda：使用全局内存来存储连续的各种大小的数据
【发布时间】：2012-03-10 17:25:12
【问题描述】：

我在使用全局内存中的字节缓冲区来存储各种大小的整数（8 位、16 位、32 位、64 位）时遇到问题。

如果我将一个整数存储在一个非 4 字节的倍数的指针值处（例如，因为我只存储了一个 8 位整数），则地址会向下舍入，删除之前的数据。

__global__ void kernel(char* pointer)
{
    *(int*)(pointer+3)=3300000;
}

在此示例代码中，考虑到指针是 4 的倍数，使用 (pointer)、(pointer+1)、(pointer+2)、(pointer+3) 中的任何一个，整数存储在 (pointer) 处。

cuda 内存是否在硬件级别以 32 位块的形式组织？有什么办法可以做到这一点？

【问题讨论】：

为什么不直接使用int 指针并让编译器 决定最佳对齐方式？您知道，数百名专家设计的一项技术可以尽可能地与硬件配合使用吗？
因为我必须存储各种大小的数据：char、int、long long int 等。在我的项目中，我无法使用 int 指针，每次需要时都会丢失 24 位内存存储一个字符。（等效代码在 C 中工作得非常好，所以我似乎是 Cuda 特异性）。
Cuda 需要字大小对齐。这是不可协商的。您不能将非 32 位对齐的地址转换为 32 位类型。
除了某些硬件根本不会允许损坏对齐（直至并包括硬件陷阱和崩溃），如果您使用在 x86 等更宽松的硬件上强制执行这样的操作，这样即使这样可以正常工作，也很可能会破坏您从 GPU 硬件中获得的任何好处。
好的，谢谢你们的cmets，他们很有帮助。我承认这不是一个好主意，但在这种情况下，我非常担心内存使用优化。无论如何，我会设法使我的代码适应这个约束！

标签： pointers memory-management cuda gpgpu

【解决方案1】：

字长对齐在 CUDA 中是不可协商的。但是，如果您出于某种原因愿意承受性能损失，您可以将数据打包到 char * 中，然后编写自己的自定义存储函数，例如

__inline __device__ void Assign(int val, char * arr, int len)
{
   for (int idx = 0; idx < len; idx++)
      *(arr+idx)=(val & (0xFF<<(idx<<8))
}

__inline __device__ int Get(char * arr, int idx, int len)
{
   int val;
   for (int idx = 0; idx < len; idx++)
      val=(int)(*arr[idx+len*idx]<<(idx<<8)));
   return val;
}

希望有帮助！

【讨论】：

谢谢，很有帮助！我将尝试您的解决方案，为我的应用程序找到性能和内存使用优化之间的最佳平衡。