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

Question

我正在尝试将uint8_t数组强制转换为uint32_t数组。但是，当我尝试执行此操作时，我似乎无法访问每个连续的4个字节。

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

这些都得到相同的值*((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中不允许未对齐访问，但是prmtPTX指令具有一种方便的模式来模拟寄存器内未对齐读取的影响。这可以通过一些嵌入式PTX组件来实现。如果您可以忍受数组末尾的读取，则代码将变得非常简单：

// 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个字节。

上面的设备代码在允许不对齐访问的little-endian主机上具有与以下代码相同的效果：

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