Блог

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

Допустим, у меня есть массив uint8_t с 8 байтами. Я хотел бы получить доступ к байту 2 -> 6 как к одному uint32_t.

Все они получают одно и то же значение *((uint32_t*)amp;uint8Array[0]) , *((uint32_t*)amp;uint8Array[1]) , *((uint32_t*)amp;uint8Array[2]) , *((uint32_t*)amp;uint8Array[3])

While *((uint32_t*)amp;uint8Array[4]) получает байты 4 -> 8, как и ожидалось.

Похоже, я не могу получить доступ к 4 последовательным байтам с любого адреса?

Есть ли какой-либо способ, которым я могу это сделать?

Хотя в CUDA не разрешен невыровненный доступ, prmt инструкция PTX имеет удобный режим для эмуляции эффекта невыровненных чтений внутри регистров. Это может быть представлено с помощью немного встроенной PTX-сборки. Если вы можете допустить чтение после конца массива, код становится довольно простым:

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

Чтобы гарантировать, что доступ после конца массива остается безопасным, округлите количество выделенных байт до кратного 4 и добавьте еще 4 байта.

Приведенный выше код устройства имеет тот же эффект, что и следующий код на хосте младшего уровня, который допускает невыровненные обращения:

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

Если вам нужны байты 2,6, вам придется объединить несколько выровненных загрузок, чтобы получить то, что вы хотите.

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

Технически, это также переносимый способ выполнения чего-либо на C в целом, но мы все избалованы, потому что вам не нужно выполнять дополнительную работу на x86, ARM, Power или других распространенных архитектурах.

Как предлагает @DietrichEpp, вам придется выполнить две загрузки; и, как предлагает @tera, вы можете объединить эти две загрузки в общем виде по дешевке, даже если смещение не известно заранее (т. Е. Когда начальный адрес uint8Array произвольный), используя prmt инструкцию PTX.

Я предложу решение, основанное на @tera, которое позволит вам сделать:

 value = read_unaligned(amp;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 resu<
    asm("prmt.b32.f4e %0, %1, %2, %3;"
        : "=r"(result)
        : "r"(first_word), "r"(second_word), "r"(control_bits) );
    return resu<
}

/*
 * 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 amp; ~((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) amp; 0x3 == 0);
    if (ptr_is_already_aligned) { return *ptr; }
    return read_unaligned_unsafe(ptr);
}