在 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 和编译器会明智地处理合理的情况，但是在尝试编写 "optimal code" 之前您应该非常熟悉对齐和硬件限制，因为它是除非你确切地知道你在做什么，否则可能会写很多毫无意义的废话。

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

Are conversions from float* to float3* in CUDA safe?

c++

cuda

type-conversion

type-safety