Оптимизация копирования по потокам и заполнение нулями данных в регистрах / L1

Я пишу ядро, в котором, помимо прочего, каждый поток заполняет одну переменную данными, составляющими младшие байты, и дополняет остальные (при условии, что они имеют порядок байтов). Это выполняется неоднократно и неравномерно по потокам, поскольку некоторые потоки могут иметь больше байтов для копирования в свои переменные и меньше заполнителей, а некоторые потоки — меньше и больше заполнителей. И результат, и незаполненные данные находятся либо в регистре (для меньших размеров), в разделяемой памяти или в локальной памяти (которая должна быть покрыта L1).

Другими словами, предположим, что каждый поток выполняет:

T padded;
pad_high_bytes(padded, my_low_bytes, my_num_low_bytes);
do_stuff(padded);

где мы имеем:

template <typename T>
__device__ __forceinline__
void pad_high_bytes(
T&                                result,
const unsigned char* __restrict__ low_bytes,
unsigned                          num_low_bytes);

Если T большой (скажем, struct { int data[50]; }) тогда я думаю, что я должен просто использовать код устройства CUDA memcpy(), Однако обычно это не так — размер T обычно равен 4 или 8, а число младших байтов обычно составляет от 1 до 3, и даже 0 не редкость.

Я, очевидно, могу перебрать байты и надеяться на лучшее. Я также могу «перебрать все целые», если осталось скопировать более 4 байтов, а затем перебрать оставшиеся байты. Но:

1. Могу ли я сделать лучше?
2. Должен ли я как-то учитывать выравнивание? Или это слишком много хлопот?
3. Должны ли копирование и заполнение полностью разделяться в коде или я должен как-то их объединить?
4. Что было бы разумным значением для sizeof (T), при котором нужно переключиться на memcpy () — если вообще?

Помните, что функция основана на типе результата, поэтому, если вам есть, что сказать в отношении меньших / больших Т, это полезно.