如何将展平的二维数组从全局内存复制到 CUDA 中的共享内存答案

【问题标题】：How to copy a flattened 2D array from Global Memory to Shared Memory in CUDA如何将展平的二维数组从全局内存复制到 CUDA 中的共享内存
【发布时间】：2013-03-22 13:06:22
【问题描述】：

我有一个内核接收一个扁平的二维数组，我想在每次共享内存时复制一行数组，我的内核如下所示：

__global__ void searchKMP(char *test,size_t pitch_test,int ittNbr){
    int  tid = blockDim.x * blockIdx.x + threadIdx.x;
    int strideId = tid * 50;

    int m = 50;

    __shared__ char s_test[m];

    int j;
               //this loops over the number of lines in my 2D array           
                   for(int k=0; k<ittNbr; k++){

                   //this loops to store my flattened (basically threats 1 line at a time) array into shared memory     
                   if(threadIdx.x==0){
                     for(int n =0; n<50; ++n){
                    s_test[n] = *(((char*)test + k * pitch_test) + n);

                }
             }
            __syncthreads();


             j=0;

            //this is loop to process my shared memory array against another 1D array
             for(int i=strideID; i<(strideID+50); i++{
             ...dosomething...
             (increment x if a condition is met) 
             ...dosomething...
             }
             __syncthreads();
             if(x!=0)
                cache[0]+=x;

            ...dosomething...

}

虽然当我验证 x 的值时，x 的值始终在变化，或者随着线程数的变化而变化。例如，当 20 个 250 个线程块根据执行返回值 7 或 6 时，10 个 500 个线程块返回 9。我想知道问题是否来自复制到共享内存中的二维扁平数组，或者是否在这段代码中做错了什么。

【问题讨论】：

您使用此内核运行的每个块有多少线程？
20 个 500 个线程的块或 10 个 1000 个线程的块
好的，那么为什么块中的每个线程都试图将数据加载到您的共享内存数组中呢？正如所写的那样，您展示的那个非常粗略的内核代码没有什么意义，恐怕......
确实如此，我实际上改变了这一点，以确保我尝试加载的每个数据只加载一次，然后在线程之间共享

标签： c cuda parallel-processing nvidia

【解决方案1】：

看起来您在共享内存中的数组有 20 个元素：

   int m = 20;
   __shared__ char s_test[m];

但在您的内部循环中，您尝试编写 50 个元素：

   for(int n =0; n<50; ++n){
      s_test[n] = *(((char*)test + k * pitch_test) + n);

我不知道这是否是您正在寻找的具体问题，但这似乎不起作用。

【讨论】：

没错，我实际上犯了一个错误，因为我一直在更改代码以试图找出导致代码错误的原因：/但感谢您指出这一点。

【解决方案2】：

共享内存在同一块中的所有线程之间共享

不是很清楚，为什么需要共享内存以及你在做什么：

在您的代码中，块中的所有线程多次将相同的值写入您的共享内存，但这是多余的

使用共享内存的常用方法如下：

if(threadIdx.x < m)
  s_test[threadIdx.x] = *(global_mem_pointer + threadIdx.x);

__syncthreads();

块中的所有线程“同时”写入它们自己的值，在 __syncthreads(); 之后，您的内存将充满您需要的内容，并且对于块中的所有线程都是可见的

【讨论】：