Я пытаюсь разбить и изменить структуру массива асинхронно с использованием ядра CUDA. memcpy() не работает внутри ядра, а также cudaMemcpy() *; Я в недоумении.
Может ли кто-нибудь сказать мне предпочтительный метод для копирования памяти из ядра CUDA?
Стоит отметить, cudaMemcpy(void *to, void *from, size, cudaMemcpyDeviceToDevice) НЕ будет работать для того, что я пытаюсь сделать, потому что он может быть вызван только из-за ядра и не выполняется асинхронно.
Да, есть эквивалент memcpy, который работает внутри ядер cuda. Это называется memcpy. Как пример:
__global__ void kernel(int **in, int **out, int len, int N)
{
int idx = threadIdx.x + blockIdx.x*blockDim.x;
for(; idx<N; idx+=gridDim.x*blockDim.x)
memcpy(out[idx], in[idx], sizeof(int)*len);
}
который без ошибок компилируется следующим образом:
$ nvcc -Xptxas="-v" -arch=sm_20 -c memcpy.cu
ptxas info : Compiling entry function '_Z6kernelPPiS0_ii' for 'sm_20'
ptxas info : Function properties for _Z6kernelPPiS0_ii
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 11 registers, 48 bytes cmem[0]
и испускает PTX:
.version 3.0
.target sm_20
.address_size 32
.file 1 "/tmp/tmpxft_00000407_00000000-9_memcpy.cpp3.i"
.file 2 "memcpy.cu"
.file 3 "/usr/local/cuda/nvvm/ci_include.h"
.entry _Z6kernelPPiS0_ii(
.param .u32 _Z6kernelPPiS0_ii_param_0,
.param .u32 _Z6kernelPPiS0_ii_param_1,
.param .u32 _Z6kernelPPiS0_ii_param_2,
.param .u32 _Z6kernelPPiS0_ii_param_3
)
{
.reg .pred %p<4>;
.reg .s32 %r<32>;
.reg .s16 %rc<2>;
ld.param.u32 %r15, [_Z6kernelPPiS0_ii_param_0];
ld.param.u32 %r16, [_Z6kernelPPiS0_ii_param_1];
ld.param.u32 %r2, [_Z6kernelPPiS0_ii_param_3];
cvta.to.global.u32 %r3, %r15;
cvta.to.global.u32 %r4, %r16;
.loc 2 4 1
mov.u32 %r5, %ntid.x;
mov.u32 %r17, %ctaid.x;
mov.u32 %r18, %tid.x;
mad.lo.s32 %r30, %r5, %r17, %r18;
.loc 2 6 1
setp.ge.s32 %p1, %r30, %r2;
@%p1 bra BB0_5;
ld.param.u32 %r26, [_Z6kernelPPiS0_ii_param_2];
shl.b32 %r7, %r26, 2;
.loc 2 6 54
mov.u32 %r19, %nctaid.x;
.loc 2 4 1
mov.u32 %r29, %ntid.x;
.loc 2 6 54
mul.lo.s32 %r8, %r29, %r19;
BB0_2:
.loc 2 7 1
shl.b32 %r21, %r30, 2;
add.s32 %r22, %r4, %r21;
ld.global.u32 %r11, [%r22];
add.s32 %r23, %r3, %r21;
ld.global.u32 %r10, [%r23];
mov.u32 %r31, 0;
BB0_3:
add.s32 %r24, %r10, %r31;
ld.u8 %rc1, [%r24];
add.s32 %r25, %r11, %r31;
st.u8 [%r25], %rc1;
add.s32 %r31, %r31, 1;
setp.lt.u32 %p2, %r31, %r7;
@%p2 bra BB0_3;
.loc 2 6 54
add.s32 %r30, %r8, %r30;
ld.param.u32 %r27, [_Z6kernelPPiS0_ii_param_3];
.loc 2 6 1
setp.lt.s32 %p3, %r30, %r27;
@%p3 bra BB0_2;
BB0_5:
.loc 2 9 2
ret;
}
Блок кода в BB0_3 представляет собой цикл memcpy байтового размера, автоматически генерируемый компилятором. Возможно, с точки зрения производительности не очень хорошая идея использовать его, но он полностью поддерживается (и долгое время использовался на всех архитектурах).
Отредактировано четыре года спустя, чтобы добавить, что, поскольку API времени выполнения на стороне устройства был выпущен как часть цикла выпуска CUDA 6, также можно напрямую вызывать что-то вроде
cudaMemcpyAsync(void *to, void *from, size, cudaMemcpyDeviceToDevice)
в коде устройства для всех архитектур, которые его поддерживают (Compute Capability 3.5 и более новое оборудование, использующее отдельную компиляцию и связывание устройств).
В моем тестировании лучший ответ - написать свою собственную процедуру циклического копирования. В моем случае:
__device__
void devCpyCplx(const thrust::complex<float> *in, thrust::complex<float> *out, int len){
// Casting for improved loads and stores
for (int i=0; i<len/2; ++i) {
((float4*) out)[i] = ((float4*) out)[i];
}
if (len%2) {
((float2*) out)[len-1] = ((float2*) in)[len-1];
}
}
memcpy работает в ядре, но может быть намного медленнее. cudaMemcpyAsync с хоста является допустимым вариантом.
Мне нужно было разделить 800 смежных векторов длиной ~ 33 000 и длиной 16 500 в другом буфере с 1600 вызовами копирования. Время с nvvp:
@talonmies сообщает, что memcpy копирует побайтно, что неэффективно при загрузке и хранении. Я все еще ориентируюсь на compute 3.0, поэтому не могу протестировать cudaMemcpy на устройстве.
Изменить: Протестировано на более новом устройстве. время выполнения устройства cudaMemcpyAsync(out, in, bytes, cudaMemcpyDeviceToDevice, 0) сопоставимо с хорошим циклом копирования и лучше, чем с плохим циклом копирования. Обратите внимание, что для использования API времени выполнения устройства могут потребоваться изменения компиляции (sm> = 3.5, отдельная компиляция). Для компиляции обратитесь к руководству по программированию и nvcc.
Устройство memcpy плохое. Хозяин cudaMemcpyAsync хорошо. Устройство cudaMemcpyAsync хорошо.
cudaMemcpy() действительно выполняется асинхронно, но вы правы, он не может быть выполнен изнутри ядра.
Определена ли новая форма массива на основе какого-то вычисления? Затем вы обычно запускаете то же количество потоков, что и записи в вашем массиве. Каждый поток запускает вычисление, чтобы определить источник и назначение одной записи в массиве, а затем скопировать ее там с одним присваиванием. (dst[i] = src[j]). Если новая форма массива не основана на вычислениях, может быть более эффективным запустить серию cudaMemcpy() с cudaMemCpyDeviceToDevice с хоста.