在 cuda 的 1 个事务中每个线程读取 4 个字符

Question

我最近在学习 CUDA。我有一个关于内存事务的问题。 我的理解是，在每个事务中，32 个连续的线程（在同一个块中）可以同时访问连续的 128 个字节（32 个单精度字）的内存，这称为 warp。 但在示例中，每个线程始终将（4 字节）字作为 1 个完整变量进行访问。所以我的问题是，如果我在全局内存中的数组是用char类型定义的，那么32个线程是否可以全部访问这块内存，同时分别读取4个连续的char？

因此，例如，如果我编写代码：

__global__
void kernel(char *d_mask)
{
    extern __shared__ char s_tmp[];
    const unsigned int thId = threadIdx.x;
    const unsigned int elementId = 4 * (threadIdx.x + blockDim.x * blockIdx.x);

    s_tmp[thId_x] = d_mask[elementId];
    s_tmp[1 + thId_x] = d_mask[elementId + 1];
    s_tmp[2 + thId_x] = d_mask[elementId + 2];
    s_tmp[3 + thId_x] = d_mask[elementId + 3];
    __syncthreads();

    /* calculation */
}

那么，每个线程会同时读取 4 个字节吗？如果没有，我怎么能做到呢？我应该使用 memcpy 之类的 API 吗？

Answer 1

为了获得适当有效的读取，有必要将正在读取的字节组合到单个事务中；我们通常不能通过将事情分解成几行代码来做到这一点。

为了将事物组合成一个事务，有向量类型将多个元素组合成一个类型。只要我们pay attention to proper alignment，我们就可以将char 或unsigned char 数组视为例如数组。 uchar4 是一种向量类型，将四个字符组合成一个（32 位）类型。您可以在 cuda 头文件 vector_types.h 和 vector_functions.h 中找到更多好东西。

无论如何，我们可以像这样重写您的示例，以利用“矢量负载”：

__global__
void kernel(char *d_mask)
{
    extern __shared__ char s_tmp[];
    const unsigned int thId = threadIdx.x;
    const unsigned int elementId = threadIdx.x + blockDim.x * blockIdx.x;

    uchar4 *s_tmp_v  = reinterpret_cast<uchar4 *>(s_tmp);
    uchar4 *d_mask_v = reinterpret_cast<uchar4 *>(d_mask);
    s_tmp_v[thId] = d_mask_v[elementId];
    __syncthreads();

    /* calculation */
}