#cuda #alignment #memory-alignment
#cuda #выравнивание #выравнивание по памяти
Вопрос:
Я пытаюсь преобразовать массив 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 последовательным байтам с любого адреса?
Есть ли какой-либо способ, которым я могу это сделать?
Комментарии:
1. Вы не можете этого сделать. Вы можете получить доступ только к правильно выровненным словам. (Некоторые процессоры могут позволить вам обойтись без этого, но это дополнительная работа над работой процессора, а некоторые процессоры этого не делают, и язык C не требует от них этого.) Если вы хотите выполнить переносимый доступ без выравнивания, вы должны сделать это «вручную», как в ответе @DietrichEpp .
2. CUDA имеет очень хорошо документированные требования к выравниванию, и все транзакции в памяти должны быть выровнены по размеру транзакции
Ответ №1:
Хотя в 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;
}
Комментарии:
1. Разве вы не предполагаете, что адрес меньше 2 ^ 32 — 1?
2. Также, пожалуйста, взгляните на мой ответ.
3. Как вы думаете, почему код завершится ошибкой для адресов > = 2 ^ 32? Обратите внимание, что релевантны только два младших бита
alignment
, поэтому 32-разрядного регистра более чем достаточно.4. Также этот код очень консервативен, явно маскируя два младших бита адреса. Я не думаю, что какое-либо существующее в настоящее время оборудование Nvidia требует этого.
Ответ №2:
Если вам нужны байты 2,6, вам придется объединить несколько выровненных загрузок, чтобы получить то, что вы хотите.
uint32_t *ptr = ...;
uint32_t value = (ptr[0] >> 16) | (ptr[1] << 16);
Технически, это также переносимый способ выполнения чего-либо на C в целом, но мы все избалованы, потому что вам не нужно выполнять дополнительную работу на x86, ARM, Power или других распространенных архитектурах.
Комментарии:
1. Переосмысление типа uint8_t как несовместимого типа является неопределенным поведением.
2. @2501: Если бы это был C, это было бы в том случае, если вы используете очень узкое прочтение стандарта C, но это не C. Это очень похоже на C, поэтому я могу понять, почему вы так думаете.
3. Если cuda не использует язык c, какой язык он использует и где я могу найти спецификацию?
4. @2501: Я не знаю, где найти спецификацию CUDA.
5. Тогда откуда вы знаете, что это не соответствует строгому сглаживанию; откуда вы берете эту информацию?
Ответ №3:
Как предлагает @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);
}
Комментарии:
1. Мне это кажется неправильным — вам нужно передать (два младших бита) невыровненный указатель для управляющих битов. Почему бы вам не протестировать свой код, когда вы его уже скомпилировали?
2. @tera: Действительно, мне нужно было передать младшие биты невыровненного указателя. Исправлено и протестировано.