需要CUDA设备内存交易

Question

我写了一个小型cuda代码来理解全局内存共享内存传输事务。代码如下：从设备存储器到共享存储器

#include <iostream> 
using namespace std; 

__global__ void readUChar4(uchar4* c, uchar4* o){ 
    extern __shared__ uchar4 gc[]; 
    int tid = threadIdx.x; 
    gc[tid] = c[tid]; 
    o[tid] = gc[tid]; 
} 

int main(){ 
    string a = "aaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaa\ 
aaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaa\ 
aaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaa"; 
    uchar4* c; 
    cudaError_t e1 = cudaMalloc((void**)&c, 128*sizeof(uchar4)); 
    if(e1==cudaSuccess){ 
    uchar4* o; 
    cudaError_t e11 = cudaMalloc((void**)&o, 128*sizeof(uchar4)); 

    if(e11 == cudaSuccess){ 
     cudaError_t e2 = cudaMemcpy(c, a.c_str(), 128*sizeof(uchar4), cudaMemcpyHostToDevice); 
     if(e2 == cudaSuccess){ 
     readUChar4<<<1,128, 128*sizeof(uchar4)>>>(c, o); 
     uchar4* oFromGPU = (uchar4*)malloc(128*sizeof(uchar4)); 
     cudaError_t e22 = cudaMemcpy(oFromGPU, o, 128*sizeof(uchar4), cudaMemcpyDeviceToHost); 
     if(e22 == cudaSuccess){ 
      for(int i =0; i < 128; i++){ 
      cout << oFromGPU[i].x << " "; 
      cout << oFromGPU[i].y << " "; 
      cout << oFromGPU[i].z << " "; 
      cout << oFromGPU[i].w << " " << endl; 

      } 
     } 
     else{ 
      cout << "Failed to copy from GPU" << endl; 
     } 
     } 
     else{ 
     cout << "Failed to copy" << endl; 
     } 
    } 
    else{ 
     cout << "Failed to allocate output memory" << endl; 
    } 
    } 
    else{ 
    cout << "Failed to allocate memory" << endl; 
    } 
    return 0; 
} 

此代码简单地复制数据，并返回到设备存储器。我有以下三个问题：

1. 在这种情况下，从设备内存传输到共享内存是否保证了4次内存事务？我相信这取决于cudaMalloc如何分配内存;如果内存是以随机方式分配的，这样数据分散在内存中，那么它将花费超过4次内存交易。但是，如果cudaMalloc以128字节块的形式分配内存，或者连续分配内存，则不应超过4个内存事务。
2. 上述逻辑是否也适用于将数据从共享内存写入设备内存，即传输将在4个内存事务中完成。
3. 该代码是否会导致银行冲突。我相信，如果线程按顺序分配了ID，那么这段代码不会导致银行冲突。但是，如果线程32和64计划在同一个warp中运行，那么此代码可能会导致bank冲突。

Answer 1

在您提供的代码中（在此重复），编译器将完全删除共享内存存储和加载，因为它们没有对代码执行任何必要或有益的操作。

__global__ void readUChar4(uchar4* c, uchar4* o){ 
    extern __shared__ uchar4 gc[]; 
    int tid = threadIdx.x; 
    gc[tid] = c[tid]; 
    o[tid] = gc[tid]; 
} 

假设你没有与共享内存大一些，因此没有消除，那么：

1. 从和全局内存在这段代码的加载和存储将采取每经一次（假设费米或更高版本的GPU），因为它们每个线程只有32位（uchar4 = 4 * 8位）（每个warp总共128字节）。 cudaMalloc连续分配内存。
2. 1.的答案也适用于商店，是的。
3. 此代码中没有银行冲突。变形中的线程始终是连续的，第一个线程是变形大小的倍数。所以线程32和64永远不会处于相同的翘曲状态。而且，由于您正在加载和存储32位数据类型，并且这些存储区的宽度为32位，所以不存在冲突。