2

次のカーネル関数は、cudpp、cudaライブラリ（http://gpgpu.org/developer/cudpp）でのコンパクトな操作です。

私の質問は、なぜ開発者が書き込み部分を8回繰り返すのかということです。そして、なぜそれがパフォーマンスを向上させることができるのですか？

そして、なぜ1つのスレッドが8つの要素を処理するのか、なぜ各スレッドが1つの要素を処理しないのでしょうか。

template <class T, bool isBackward>
__global__ void compactData(T                  *d_out, 
                        size_t             *d_numValidElements,
                        const unsigned int *d_indices, // Exclusive Sum-Scan Result
                        const unsigned int *d_isValid,
                        const T            *d_in,
                        unsigned int       numElements)
{
  if (threadIdx.x == 0)
  {
        if (isBackward)
            d_numValidElements[0] = d_isValid[0] + d_indices[0];
    else
        d_numValidElements[0] = d_isValid[numElements-1] + d_indices[numElements-1];
   }

   // The index of the first element (in a set of eight) that this
   // thread is going to set the flag for. We left shift
   // blockDim.x by 3 since (multiply by 8) since each block of 
   // threads processes eight times the number of threads in that
   // block
   unsigned int iGlobal = blockIdx.x * (blockDim.x << 3) + threadIdx.x;

   // Repeat the following 8 (SCAN_ELTS_PER_THREAD) times
   // 1. Check if data in input array d_in is null
   // 2. If yes do nothing
   // 3. If not write data to output data array d_out in
   //    the position specified by d_isValid
   if (iGlobal < numElements && d_isValid[iGlobal] > 0) {
       d_out[d_indices[iGlobal]] = d_in[iGlobal];
   }
   iGlobal += blockDim.x;  
   if (iGlobal < numElements && d_isValid[iGlobal] > 0) {
       d_out[d_indices[iGlobal]] = d_in[iGlobal];       
   }
   iGlobal += blockDim.x;
   if (iGlobal < numElements && d_isValid[iGlobal] > 0) {
       d_out[d_indices[iGlobal]] = d_in[iGlobal];
   }
   iGlobal += blockDim.x;
   if (iGlobal < numElements && d_isValid[iGlobal] > 0) {
       d_out[d_indices[iGlobal]] = d_in[iGlobal];
   }
   iGlobal += blockDim.x;
   if (iGlobal < numElements && d_isValid[iGlobal] > 0) {
       d_out[d_indices[iGlobal]] = d_in[iGlobal];
   }
   iGlobal += blockDim.x;
   if (iGlobal < numElements && d_isValid[iGlobal] > 0) {
       d_out[d_indices[iGlobal]] = d_in[iGlobal];
   }
   iGlobal += blockDim.x;
   if (iGlobal < numElements && d_isValid[iGlobal] > 0) {
       d_out[d_indices[iGlobal]] = d_in[iGlobal];
   }
   iGlobal += blockDim.x;
   if (iGlobal < numElements && d_isValid[iGlobal] > 0) {
       d_out[d_indices[iGlobal]] = d_in[iGlobal];
   }
}
4

1 に答える 1

1

私の質問は、なぜ開発者が書き込み部分を8回繰り返すのかということです。そして、なぜそれがパフォーマンスを向上させることができるのですか？

@torrential_codingが述べたように、ループ展開はパフォーマンスに役立ちます。特にこのような場合、ループは非常にタイトです（ロジックがほとんどありません）。ただし、コーダーは、手動で行うのではなく、CUDAの自動ループ展開のサポートを使用する必要がありました。

そして、なぜ1つのスレッドが8つの要素を処理するのか、なぜ各スレッドが1つの要素を処理しないのでしょうか。

iGlobalの完全なインデックスを計算し、各カーネルが1つの要素のみを実行した場合に実行する必要がある操作ごとではなく、8つの操作ごとにthreadIdx.xがゼロであることを確認するだけで、パフォーマンスがわずかに向上する可能性があります。

于 2012-06-01T18:10:35.963 に答える