cuda - 从未对齐的 uint8_t 重铸为 uint32_t 数组读取 - 未获取所有值

Question

我正在尝试将 uint8_t 数组转换为 uint32_t 数组。但是，当我尝试这样做时，我似乎无法访问每一个连续的 4 个字节。

假设我有一个 8 字节的 uint8_t 数组。我想访问字节 2 -> 6 作为一个 uint32_t。

这些都得到相同的值 *((uint32_t*)&uint8Array[0]), *((uint32_t*)&uint8Array[1]), *((uint32_t*)&uint8Array[2]), *((uint32_t*)&uint8Array[3])

虽然 *((uint32_t*)&uint8Array[4]) 按预期获取字节 4 -> 8。

所以我似乎无法从任何地址访问 4 个连续字节？

有什么办法可以做到这一点吗？

Answer 1

虽然 CUDA 中不允许未对齐的访问，但 prmt PTX instruction有一个方便的模式来模拟寄存器内未对齐读取的影响。这可以通过一点 inline PTX assembly 来暴露。 .如果您可以容忍读取超过数组末尾，代码将变得非常简单:

// WARNING! Reads past ptr!
__device__ uint32_t read_unaligned(void* ptr)
{
    uint32_t result;
    asm("{\n\t"
        "   .reg .b64    aligned_ptr;\n\t"
        "   .reg .b32    low, high, alignment;\n\t"
        "   and.b64      aligned_ptr, %1, 0xfffffffffffffffc;\n\t"
        "   ld.u32       low, [aligned_ptr];\n\t"
        "   ld.u32       high, [aligned_ptr+4];\n\t"
        "   cvt.u32.u64  alignment, %1;\n\t"
        "   prmt.b32.f4e %0, low, high, alignment;\n\t"
        "}"
        : "=r"(result) : "l"(ptr));
    return result;
}

为确保超出数组末尾的访问保持无害，将分配的字节数四舍五入为 4 的倍数，然后再添加 4 个字节。

以上设备代码与以下代码在容忍未对齐访问的小端主机上具有相同的效果:

__host__ uint32_t read_unaligned_host(void* ptr)
{
    return *(uint32_t*)ptr;
}