正確使用CUDA共享內存進行2D非連續數據訪問

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(); 
} 

內核調用部分：

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(); 

然而，似乎我無法收拾所有需要的DATAS到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; 
    } 

。