从未对齐的 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

如果您想要字节 2..6，则必须组合多个对齐的负载才能获得您想要的内容。

uint32_t *ptr = ...;
uint32_t value = (ptr[0] >> 16) | (ptr[1] << 16);

从技术上讲，这也是通常用 C 语言做事的可移植方式，但我们都被宠坏了，因为您不必在 x86 上做额外的工作、ARM、Power 或其他常见架构。

Answer 2

虽然 CUDA 中不允许未对齐的访问，但 prmt PTX instruction has a handy mode to emulate the effect of unaligned reads within registers. This can be exposed with a bit of 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;
}

Answer 3

正如@DietrichEpp 所建议的那样，您必须加载两次；并且正如@tera所建议的那样，即使事先不知道未对齐（即当 uint8Array 的初始地址是任意的），您也可以使用 prmt PTX 指令以便宜的方式组合这两个负载。

我将提供一个基于@tera 的解决方案，您可以这样做：

value = read_unaligned(&uint8Array[offset]);

安全且（相对）高效。此外，它只有一个内联 PTX 汇编指令，如果需要，还有一个 "unsafe" 变体：

#include <cstdint>
#include <cuda_runtime_api.h>

__device__ __forceinline__ uint32_t prmt_forward_4_extract(
    uint32_t first_word,
    uint32_t second_word, 
    uint32_t control_bits)
{
    uint32_t result;
    asm("prmt.b32.f4e %0, %1, %2, %3;"
        : "=r"(result)
        : "r"(first_word), "r"(second_word), "r"(control_bits) );
    return result;
}

/*
 * This unsafe, faster variant may read past the 32-bit naturally-aligned
 * word containing the last relevant byte
 */
__device__ inline uint32_t read_unaligned_unsafe(const uint32_t* __restrict__ ptr)
{
    /*
     *  Clear the bottom 2 bits of the address, making the result aligned 
     *  for the purposes of reading a 32-bit (= 4-byte) value
     */
    auto aligned_ptr  = (uint32_t*) ((uint64_t) ptr & ~((uint64_t) 0x3));
    auto first_value  = *aligned_ptr;
    auto second_value = *(aligned_ptr + 1);

    auto lower_word_of_ptr = (uint32_t)((uint64_t)(ptr));

    return prmt_forward_4_extract(first_value, second_value, lower_word_of_ptr);
}

__device__ inline uint32_t read_unaligned(const uint32_t* __restrict__ ptr)
{
    auto ptr_is_already_aligned = ((uint64_t)(ptr) & 0x3 == 0);
    if (ptr_is_already_aligned) { return *ptr; }
    return read_unaligned_unsafe(ptr);
}

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

Reading from an unaligned uint8_t recast as a uint32_t array - not getting all values

cuda

alignment

memory-alignment