Правильное использование общей памяти CUDA для двумерного несмежного доступа к данным

Question

У меня есть основной вопрос, связанный с двумерным доступом к потоку. Я хочу, чтобы скопировать несмежных данных в непрерывный буфер и использование Cuda memcopy можно проиллюстрировать следующим образом:

void pack_cuda(float *dstbuf, IOV *srciov, int num_iov) 
{ 
    int i; 
    float *ptr; 
    ptr = buf; 
    for (i = 0; i < num_iov; i++) { 
    cudaMemcpy(ptr, srciov[i].bufaddr, srciov[i].len, cudaMemcpyDefault); 
    ptr = (char *)ptr + srciov[i].len; 
    } 
} 

* srciov хранит адрес начала памяти и длину каждого несмежных данных в массиве состав.

* dstbuf будет хранить упакованные смежные данные после завершения функции.

Теперь я хочу реализовать его с использованием ядер CUDA.

__global__ void pack_cuda(float *dstbuf, IOV *srciov, int num_iov) 
{ 
    int i = blockIdx.x * blockDim.x + threadIdx.x; 
    int j = blockIdx.y * blockDim.y + threadIdx.y; 
    int k; 
    extern __shared__ size_t tmpdbuflen[16*3]; //suppose num_iov is 16 

    if (j == 0){ 
    if (i < 16){ 
    tmpdbuflen[i] = (srciov[i].len); //store length to calculate presum 
    tmpdbuflen[i+16] = tmpdbuflen[i]; //store length 
    tmpdbuflen[i+32] = ((srciov+i)->bufaddr) - (srciov->bufaddr); //store addr difference 
    } 
    __syncthreads(); 


    for (k = 0; k < i; k++) 
     tmpdbuflen[i] += srciov[k].len; 
    } 

    __syncthreads(); 

    if (i < 16 && j < srciov[i].len){ //wondering whether this is correct use 
    dst[tmpdbuflen[i] + j] = *(src + tmpdbuflen[i+32] + j); 
    } 

    __syncthreads(); 
} 

Kernel призывание часть:

dim3 dimblock(16, 16); //the length of each non-contiguous data is less than 16 
dim3 dimgrid(1,1); 
const unsigned int shm_size = sizeof(size_t) * 16 * 3; 
pack_cuda<<<dimgrid, dimblock, shm_size, 0>>>(dstbuf, srciov, num_iov); 
cudaDeviceSynchronize(); 

Тем не менее, кажется, что я не могу упаковать все необходимые ДАННЫЕ Into Dst буфера. Иногда только j = 0 и 1 (с соответствующими различными i) получают упакованные. Я думаю, что основной проблемой является использование разделяемой памяти. Я использую только потоки столбцов 0 (threadIdx.y == 0) для копирования информации в общую память. Тогда все потоки (без ограничений на threadIdx.y) будут получать доступ и читать информацию в общей памяти. Как изменить код, разрешающий такой дизайн?

Буду признателен, если кто-нибудь сможет выяснить мои проблемы.

Спасибо.

Answer 1

Некоторые намеки на ваш код:

__global__ void pack_cuda(float *dstbuf, IOV *srciov, int num_iov) 
{ 
    int i = blockIdx.x * blockDim.x + threadIdx.x; 
    int j = blockIdx.y * blockDim.y + threadIdx.y; 
    int k; 
    extern __shared__ size_t tmpdbuflen[16*3]; //suppose num_iov is 16 

Этот блок здесь будет выполняться только одним потоком, благодаря охраннику j==0 который позволяет только нить bid*bdim+tid = 0*0+0, Ergo нить 0 в блоке 0, что является нежелательным для вы. Я думаю, вы хотите поставить j < 16

if (j == 0){ 
    if (i < 16){ 
    tmpdbuflen[i] = (srciov[i].len); //store length to calculate presum 
    tmpdbuflen[i+16] = tmpdbuflen[i]; //store length 
    tmpdbuflen[i+32] = ((srciov+i)->bufaddr) - (srciov->bufaddr); //store addr difference 
    } 
    __syncthreads(); 


    for (k = 0; k < i; k++) 
     tmpdbuflen[i] += srciov[k].len; 
    } 

.