CUDA 中从 float* 到 float3* 的转换安全吗？

Question

我刚刚开始深入研究 CUDA 代码，它与过去相比有点爆炸，大量的指针访问和类型转换通过使用 reinterpret_cast 的指针进行。我有一个想要检查的特定案例，我在代码中看到了以下类型双关语的实例：

__device__ void func(__restrict__ float* const points, size_t size, __restrict__ float* outputPoints) {

    for (size_t index = 0; index < size; index += 3) {
        float3* const point = reinterpret_cast<float3* const>(points + index);
        float3* const output = reinterpret_cast<float3* const>(outputPoints + index);
        // operations using point;
    }
}

在 CUDA 中为您提供了一个结构 float3，如下所示：

struct float3 {
    float x, y, z
}

这种行为是否保证安全？这显然是某种双关语，但我很担心可能会有一些填充或对齐，或者会以这种方式破坏访问的东西。如果有人能够进一步了解 cuda 编译器将如何处理这个问题，因为我知道它也做了一些非常重的优化。这些会导致问题吗？

Answer 1

CUDA 保证这些内置类型的大小将在主机和设备之间保持一致，而无需填充干预（用户定义的结构和类不存在此类保证）。

设备上有对齐的基本要求，比如你读取的存储空间必须和读取的大小对齐。因此，您无法从任意字节边界读取float3，但您可以安全地从 32 位对齐的边界读取，并且 CUDA 在主机和设备上公开的内存分配 API 保证了必要的对齐以使代码您发布的内容是安全的。

您发布的代码（经过修改以消除死代码删除）基本上只是发出三个 32 位加载和三个 32 位存储。 CUDA 只有有限数量的本机事务大小，并且它们不会映射到每个线程请求的 96 位，因此这样做绝对没有优化：

__device__ void func(float* const points, size_t size, float* outputPoints) {

    for (size_t index = 0; index < size; index += 3) {
        float3* point = reinterpret_cast<float3*>(points + index);
        float3* output = reinterpret_cast<float3*>(outputPoints + index);

    float3 val = *point;
    val.x += 1.f; val.y += 2.f; val.z += 3.f;
    *output = val;
    }
}

这是做什么的：

$ nvcc -arch=sm_75 -std=c++11 -dc -ptx fffloat3.cu 
$ tail -40 fffloat3.ptx 
    // .globl   _Z4funcPfmS_
.visible .func _Z4funcPfmS_(
    .param .b64 _Z4funcPfmS__param_0,
    .param .b64 _Z4funcPfmS__param_1,
    .param .b64 _Z4funcPfmS__param_2
)
{
    .reg .pred  %p<3>;
    .reg .f32   %f<7>;
    .reg .b64   %rd<14>;


    ld.param.u64    %rd11, [_Z4funcPfmS__param_0];
    ld.param.u64    %rd8, [_Z4funcPfmS__param_1];
    ld.param.u64    %rd12, [_Z4funcPfmS__param_2];
    setp.eq.s64 %p1, %rd8, 0;
    mov.u64     %rd13, 0;
    @%p1 bra    BB6_2;

BB6_1:
    ld.f32  %f1, [%rd11];
    ld.f32  %f2, [%rd11+4];
    ld.f32  %f3, [%rd11+8];
    add.f32     %f4, %f1, 0f3F800000;
    add.f32     %f5, %f2, 0f40000000;
    add.f32     %f6, %f3, 0f40400000;
    st.f32  [%rd12], %f4;
    st.f32  [%rd12+4], %f5;
    st.f32  [%rd12+8], %f6;
    add.s64     %rd12, %rd12, 12;
    add.s64     %rd11, %rd11, 12;
    add.s64     %rd13, %rd13, 3;
    setp.lt.u64 %p2, %rd13, %rd8;
    @%p2 bra    BB6_1;

BB6_2:
    ret;
}

即所有这些转换在语法上都是虚假且毫无意义的。

如果您要更改为float2，这是每个线程的 64 位请求，并且可以进行矢量化，那么得到这个：

.visible .func _Z4funcPfmS_(
    .param .b64 _Z4funcPfmS__param_0,
    .param .b64 _Z4funcPfmS__param_1,
    .param .b64 _Z4funcPfmS__param_2
)
{
    .reg .pred  %p<3>;
    .reg .f32   %f<7>;
    .reg .b64   %rd<14>;


    ld.param.u64    %rd12, [_Z4funcPfmS__param_0];
    ld.param.u64    %rd8, [_Z4funcPfmS__param_1];
    ld.param.u64    %rd11, [_Z4funcPfmS__param_2];
    setp.eq.s64 %p1, %rd8, 0;
    mov.u64     %rd13, 0;
    @%p1 bra    BB6_2;

BB6_1:
    ld.v2.f32   {%f1, %f2}, [%rd12];
    add.f32     %f5, %f2, 0f40000000;
    add.f32     %f6, %f1, 0f3F800000;
    st.v2.f32   [%rd11], {%f6, %f5};
    add.s64     %rd12, %rd12, 8;
    add.s64     %rd11, %rd11, 8;
    add.s64     %rd13, %rd13, 2;
    setp.lt.u64 %p2, %rd13, %rd8;
    @%p2 bra    BB6_1;

BB6_2:
    ret;
}

请注意，加载和存储现在使用指令的矢量化版本。与float4相同：

    // .globl   _Z4funcPfmS_
.visible .func _Z4funcPfmS_(
    .param .b64 _Z4funcPfmS__param_0,
    .param .b64 _Z4funcPfmS__param_1,
    .param .b64 _Z4funcPfmS__param_2
)
{
    .reg .pred  %p<3>;
    .reg .f32   %f<12>;
    .reg .b64   %rd<14>;


    ld.param.u64    %rd12, [_Z4funcPfmS__param_0];
    ld.param.u64    %rd8, [_Z4funcPfmS__param_1];
    ld.param.u64    %rd11, [_Z4funcPfmS__param_2];
    setp.eq.s64 %p1, %rd8, 0;
    mov.u64     %rd13, 0;
    @%p1 bra    BB6_2;

BB6_1:
    ld.v4.f32   {%f1, %f2, %f3, %f4}, [%rd12];
    add.f32     %f9, %f3, 0f40400000;
    add.f32     %f10, %f2, 0f40000000;
    add.f32     %f11, %f1, 0f3F800000;
    st.v4.f32   [%rd11], {%f11, %f10, %f9, %f4};
    add.s64     %rd12, %rd12, 8;
    add.s64     %rd11, %rd11, 8;
    add.s64     %rd13, %rd13, 2;
    setp.lt.u64 %p2, %rd13, %rd8;
    @%p2 bra    BB6_1;

BB6_2:
    ret;
}

TLDR：您的担忧是有道理的，但是 API 和编译器会明智地处理合理的情况，但是在尝试编写“最佳代码”之前，您应该非常熟悉对齐和硬件限制，因为可以编写除非您确切地知道自己在做什么，否则很多毫无意义的废话。