0

我编写了小 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 冲突。
4

1 回答 1

2

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

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

假设您对共享内存做了一些事情，所以它没有被消除，那么：

  1. 在此代码中，从全局内存中读取和存储到全局内存将需要每个 warp 的一个事务（假设 Fermi 或更高版本的 GPU），因为它们uchar4每个线程只有 32 位（= 4 * 8 位）（每个 warp 总共 128 字节）。cudaMalloc连续分配内存。
  2. 1. 的答案也适用于商店，是的。
  3. 此代码中没有银行冲突。经线中的线程始终是连续的，第一个线程是经线大小的倍数。所以线程 32 和 64 永远不会在同一个经线中。而且由于您正在加载和存储 32 位数据类型，并且存储库是 32 位宽，因此没有冲突。
于 2012-07-12T03:19:05.230 回答