Чтение из unaligned uint8_t recast как массив uint32_t - не получение всех значений

Question

Я пытаюсь использовать массив uint8_t для массива uint32_t. Однако, когда я пытаюсь это сделать, я не могу получить доступ к каждому последующему 4 байтам.

Скажем, у меня есть массив uint8_t с 8 байтами. Я хотел бы получить доступ к байту 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 имеет удобный режим для эмулирования эффекта нестандартных чтений в пределах регистров. Это может быть выставлено бит 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, 0xfffffffc;\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 является произвольным) с использованием инструкции PTX prmt.

я предлагаю решение, основанное на @ тера, который позволит вам сделать:

value = read_unaligned(&uint8Array[offset]); 

безопасно и (относительно) эффективно. Кроме того, он будет иметь только одну встроенную инструкцию по сборке PTX и «небезопасный» вариант, если вам это необходимо:

#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); 
}