0

我写了一个简单的CUDA内核如下：

    __global__ void cudaDoSomethingInSharedMemory(float* globalArray, pitch){

      __shared__ float sharedInputArray[1088];
      __shared__ float sharedOutputArray[1088];

      int tid = threadIdx.x //Use 1D block
      int rowIdx = blockIdx.x //Use 1D grid

      int rowOffset = pitch/sizeof(float);//Offset in elements (not in bytes)

       //Copy data from global memory to shared memory (checked)
       while(tid < 1088){
           sharedInputArray[tid] = *(((float*) globalArray) + rowIdx*rowOffset + tid);
           tid += blockDim.x;
           __syncthreads();
       }
       __syncthreads();

       //Do something (already simplified and the problem still exists)
       tid = threadIdx.x;
       while(tid < 1088){
           if(tid%2==1){
              if(tid == 1087){
                 sharedOutputArray[tid/2 + 544] = 321;
              }
              else{
                  sharedOutputArray[tid/2 + 544] = 321;
              }
           }
           tid += blockDim.x;
           __syncthreads();
       }

       tid = threadIdx.x;
       while(tid < 1088){
           if(tid%2==0){
               if(tid==0){
                    sharedOutputArray[tid/2] = 123;
               }
               else{
                    sharedOutputArray[tid/2] = 123;
               }

           }
           tid += blockDim.x;
           __syncthreads();
       }
       __syncthreads();

       //Copy data from shared memory back to global memory (and add read-back for test)
       float temp = -456;
       tid = threadIdx.x;
       while(tid < 1088){
           *(((float*) globalArray) + rowIdx*rowOffset + tid) = sharedOutputArray[tid];
            temp = *(((float*) globalArray) + rowIdx*rowOffset + tid);//(1*) Errors are found.
            __syncthreads();
            tid += blockDim.x;
       }
       __syncthreads();
    }

代码是将“sharedOutputArray”从“interlaced”更改为“clustered”：“123 321 123 321 ... 123 321”更改为“123 123 123.. 123 321 321 321...321”并输出clustered结果到全局内存数组“globalArray”。“globalArray”由“cudaMallocPitch()”分配

该内核用于处理二维数组。这个想法很简单：一行一个块（所以一维网格，块数等于行数），每行有 N 个线程。行号是 1920，列号是 1088。所以有 1920 个块。

问题是：当 N（一个块中的线程数）为 64、128 或 256 时，一切正常（至少看起来正常）。但是，当 N 为 512 时（我使用的是 GTX570，CUDA 计算能力为 2.0，一个块的每个维度的最大大小为 1024），错误发生了。

错误是：全局内存中从位置 256 到 287（索引从 0 开始，错误条长度为 32 个元素，128 位）的一行中的元素（每个元素是一个 4 字节的浮点数）是 0 而不是 123 . 它看起来像“123 123 123 ... 0 0 0 0 0 ... 0 123 123 ...”。我检查了上面的行 (1*)，这些元素在“sharedOutputArray”中为 123，当在 (1*) 中读取元素（例如 tid==270）时，“temp”显示为 0。我试图查看“tid” ==255" 和 "tid==288" 元素为 123（正确）。这种类型的错误几乎发生在所有 1920 行中。

我试图“同步”（可能已经过度同步）线程，但它不起作用。让我感到困惑的是为什么 64,128 或 256 个线程工作正常，但 512 没有工作。我知道使用 512 线程可能不会针对性能进行优化，我只想知道我在哪里犯了错误。

先感谢您。

4

1 回答 1

1

您正在使用__syncthreads()内部条件代码，其中条件不会在块的线程之间统一评估。不要那样做。

在您的情况下，您可以简单地删除循环__syncthreads()内部while，因为它没有任何作用。

于 2012-10-18T10:55:45.363 回答