CUDA中float3的合并访存问题

缘起

使用CUDA做三维点云计算的时候，需要在单线程中读取float3类型的三维点，做一系列计算，再将计算后的结果写入。写完之后，Nsight Compute总是提醒有 uncoalesced global memmory，对应到行后发现，这个问题出现在float3的读入和写入位置。
众所周知，CUDA支持单次读取1byte, 2bytes,4 bytes, 8 bytes 或16 bytes，但是唯独没有float3类型的12 bytes，想来想要float3类型的数据合并访存是比较麻烦的。在stack overflow上找到一篇不错的文章，链接，本文对该文章做简要翻译。

CUDA架构对合并访存的支持

当前的CUDA 架构支持 1 字节、2 字节、4 字节、8 字节和 16 字节全局内存加载和存储。但不存在一个可用的，12 字节的float3数据的直接加载或存储机制。
float3是存在于 CUDA C++ 语言级别的抽象，硬件实现中没有float3这个概念。CUDA C++中的float3 由三个浮点数组成，因此，访问float3通常会映射到 4 字节加载/存储，它通常由三个 4 字节加载/存储操作完成。线程与线程间的访问（一次float）地址不连续，因此，无法直接合并访存。
例如，对于如下语句：

__global__ void test(float3* dest)
{
    dest[threadIdx.x] = { 1.0f, 2.0f, 3.0f };
}

数据的存储顺序为：

1st store:  xx xx t1 xx xx t2 xx xx t3 xx xx t4 xx xx t5 xx xx t6 …
2nd store:  xx t1 xx xx t2 xx xx t3 xx xx t4 xx xx t5 xx xx t6 xx …
3rd store:  t1 xx xx t2 xx xx t3 xx xx t4 xx xx t5 xx xx t6 xx xx …

其中 t i是warp 的第 ixx个线程，表示跳过的 4 字节地址。如您所见，线程执行的存储之间存在 8 字节间隙。

将float3拓展为float4存储

例如，如下方案添加了一个辅助位，且使用alignas内存对齐

struct alignas(16) Stuff
{
    float3 p;
    int blub;
};

__global__ void test(Stuff* dest)
{
    dest[threadIdx.x].p = { 1.0f, 2.0f, 3.0f };
    dest[threadIdx.x].blub = 42;
}

这样，新的Stuff数据将会通过一次写入指令存储，能减少延时，但问题在于，会增加1/3的内存消耗。这种方式做写入数据的排列为：

t1 t1 t1 t1 t2 t2 t2 t2 t3 t3 t3 t3 t4 t4 t4 t4 …

使用3个向量线程写入float3

例如，使用如下操作：

__global__ void test(float3* dest)
{
    auto i = threadIdx.x % 3;
    auto m = i == 0 ? &float3::x : i == 1 ? &float3::y : &float3::z;
    dest[threadIdx.x / 3].*m = i;
}

此时，每个线程都会对 afloat3的成员之一执行一次存储，并且连续的线程将存储到连续的 4 字节地址，从而实现完美合并的内存访问，存储结果为：

t1 t2 t3 t4 t5 t6 t7 t8 t9 t10 t11 t12 t13 t14 t15 …

使用这种方法的唯一顾虑在于，与最终写入值相关的计算消耗相比，合并访存带来的收益是否更大