CUDA：文档中的共享内存对齐

Question

我们可以在 official Nvidia website 上看到，要使用多个未知大小的共享内存数组，我们可以在内核中使用该代码：

__global__ void myKernel() {
  extern __shared__ int s[];
  int *integerData = s;                        // nI ints
  float *floatData = (float*)&integerData[nI]; // nF floats
  char *charData = (char*)&floatData[nF];      // nC chars
}

// Kernel launch
myKernel<<<gridSize, blockSize, nI*sizeof(int)+nF*sizeof(float)+nC*sizeof(char)>>>(...);

但是，我们知道内存也要对齐。如果例如 sizeof(float) == 8 和 sizeof(int) == 4 ，这段代码不是未定义的行为吗？这可能是一个令人费解的大小，但如果我们只是将 char 放在上面，问题会更加明显：

__global__ void myKernel() {
  extern __shared__ char s[];
  char *charData = s;      // nC chars
  int *integerData = (int*)&s[nC];             // nI ints
  float *floatData = (float*)&integerData[nI]; // nF floats
}

如果 nC 不是 sizeof(int) 的倍数，则指针未对齐。所以我的问题是，这是否适用于这个特定示例，还是我应该担心对齐问题？如果是，常见模式是什么或如何在代码中管理它？

Answer 1

是的，您必须担心对齐问题。

解决这个问题的最简单方法是按照从大到小的对齐方式对数组进行排序，这意味着 double + long long -> float + int -> …

由于所有整数类型的大小和对齐方式都是 2 的幂，因此当您像这样打包数组时不会浪费任何 space。需要特别注意特定于平台的类型，例如 size_t 或 wchar_t。矢量类型也是如此。 float3 的对齐方式低于 float2。

在普通的 C++ 代码中，我们可以使用 std::align 来计算我们的指针。但这不适用于 CUDA，因此如果我们需要通用解决方案，我们会推出自己的版本。

#include <cstdint>

/**
 * in normal code we could use std::align but that doesn't work with CUDA
 *
 * Behavior is undefined if alignment is not a power of 2
 * (same as everywhere else)
 *
 * \tparam T array element type
 * \param n_elements number of array entries
 * \param ptr in-out parameter. On entry, points at first usable location.
 *     On exit, will point at first location after the end of the array
 * \param space if not null, the used space (including padding for alignment)
 *    will be added to the value currently stored in here
 * \return properly aligned pointer to beginning of array
 */
template<class T>
__host__ __device__
T* align_array(std::size_t n_elements, void*& ptr,
      std::size_t* space=nullptr) noexcept
{
    const std::size_t alignment = alignof(T);
    const std::uintptr_t intptr = reinterpret_cast<uintptr_t>(ptr);
    const std::uintptr_t aligned = (intptr + alignment - 1) & -alignment;
    const std::uintptr_t end = aligned + n_elements * sizeof(T);
    if(space)
        *space += static_cast<std::size_t>(end - intptr);
    ptr = reinterpret_cast<void*>(end);
    return reinterpret_cast<T*>(aligned);
}

__global__ void myKernel(int nI, int nF, int nC)
{
  extern __shared__ char s[];
  void* sptr = s;
  volatile char* charData = align_array<char>(nC, sptr);
  volatile int* integerData = align_array<int>(nI, sptr);
  volatile float* floatData = align_array<float>(nF, sptr);
  floatData[0] = charData[0] + integerData[0];
}

void callKernel(int nI, int nF, int nC)
{
  std::size_t shared_size = 0;
  void* sptr = nullptr;
  align_array<char>(nC, sptr, &shared_size);
  align_array<int>(nI, sptr, &shared_size);
  align_array<float>(nF, sptr, &shared_size);
  myKernel<<<gridsize, blocksize, sharedsize>>>(nI, nF, nC);
}

CUDA：文档中的共享内存对齐

CUDA : Shared memory alignement in documentation

c++

cuda