6

さまざまなブロックにわたる最終的な合計のアトミック計算を使用して、倍精度配列用の従来の内積カーネルを実装しようとしています。プログラミングガイドの116ページに記載されているように、倍精度でatomicAddを使用しました。おそらく、何か間違ったことをしています。すべてのブロックのスレッド全体の部分和は正しく計算されますが、その後、アトミック操作が正しく機能していないようです。同じデータでカーネルを実行するたびに、異なる結果を受け取ります。誰かが間違いを見つけたり、別の解決策を提供したりできれば幸いです。これが私のカーネルです：

__global__ void cuda_dot_kernel(int *n,double *a, double *b, double *dot_res)
{
    __shared__ double cache[threadsPerBlock]; //thread shared memory
    int global_tid=threadIdx.x + blockIdx.x * blockDim.x;
    int i=0,cacheIndex=0;
    double temp = 0;
    cacheIndex = threadIdx.x;
    while (global_tid < (*n)) {
        temp += a[global_tid] * b[global_tid];
        global_tid += blockDim.x * gridDim.x;
    }
    cache[cacheIndex] = temp;
    __syncthreads();
    for (i=blockDim.x/2; i>0; i>>=1) {
        if (threadIdx.x < i) {
            cache[threadIdx.x] += cache[threadIdx.x + i];
        }
        __syncthreads();
    }
    __syncthreads();
    if (cacheIndex==0) {
        *dot_res=cuda_atomicAdd(dot_res,cache[0]);
    }
}

そして、これが私のデバイス関数atomicAddです：

__device__ double cuda_atomicAdd(double *address, double val)
{
    double assumed,old=*address;
    do {
        assumed=old;
        old= __longlong_as_double(atomicCAS((unsigned long long int*)address,
                    __double_as_longlong(assumed),
                    __double_as_longlong(val+assumed)));
    }while (assumed!=old);

    return old;
}
4

3 に答える 3

9

アドホックCUDAコードを使用して正しく削減するのは難しい場合があるため、CUDAツールキットに含まれているスラストアルゴリズムを使用した代替ソリューションを次に示します。

#include <thrust/inner_product.h>
#include <thrust/device_ptr.h>

double do_dot_product(int n, double *a, double *b)
{
  // wrap raw pointers to device memory with device_ptr
  thrust::device_ptr<double> d_a(a), d_b(b);

  // inner_product implements a mathematical dot product
  return thrust::inner_product(d_a, d_a + n, d_b, 0.0);
}
于 2012-02-26T02:04:35.303 に答える
4

cuda_atomicAdd関数を誤って使用しています。カーネルのこのセクション：

if (cacheIndex==0) {
    *dot_res=cuda_atomicAdd(dot_res,cache[0]);
}

犯人です。ここでは、をアトミックに追加しdot_resます。次に、アトミックに設定さdot_resれず、返される結果が返されます。この関数からの戻り結果は、アトミックに更新されている場所の以前の値であり、「情報」または呼び出し元のローカル使用のみに提供されます。アトミックに更新されたものに割り当てないでください。そもそもアトミックメモリアクセスを使用するという目的が完全に無効になります。代わりに次のようなことをしてください：

if (cacheIndex==0) {
    double result=cuda_atomicAdd(dot_res,cache[0]);
}
于 2012-02-26T08:10:54.057 に答える
-1

あなたのコードをその深さでチェックしませんでしたが、ここにいくつかのアドバイスがあります。
複雑な問題が発生した場合、人々はGPUで効率的に並列プログラミングすることを考えていないため、このような一般的なタスクにGPUのみを使用する場合にのみThrustを使用することをお勧めします。

  1. 新しい並列還元カーネルを開始して、内積を要約します。
    データはすでにデバイス上にあるため、新しいカーネルを起動してもパフォーマンスが低下することはありません。

  2. カーネルは、最新のGPUで可能なブロックの最大数にわたって拡張できないようです。もしそうなら、カーネルが数百万の値の内積を計算できるとしたら、シリアル化されたアトミック操作のためにパフォーマンスが劇的に低下します。

  3. 初心者の間違い：入力データと共有メモリのアクセス範囲はチェックされていますか？または、入力データが常にブロックサイズの倍数であることを確認しますか？そうでなければ、あなたはゴミを読むでしょう。私の間違った結果のほとんどは、この障害によるものでした。

  4. 並列削減を最適化します。私の論文または最適化マークハリス

テストされていません。メモ帳に書き留めました。

/*
 * @param inCount_s unsigned long long int Length of both input arrays
 * @param inValues1_g double* First value array
 * @param inValues2_g double* Second value array
 * @param outDots_g double* Output dots of each block, length equals the number of blocks
 */
__global__ void dotProduct(const unsigned long long int inCount_s,
    const double* inValuesA_g,
    const double* inValuesB_g,
    double* outDots_g)
{
    //get unique block index in a possible 3D Grid
    const unsigned long long int blockId = blockIdx.x //1D
            + blockIdx.y * gridDim.x //2D
            + gridDim.x * gridDim.y * blockIdx.z; //3D


    //block dimension uses only x-coordinate
    const unsigned long long int tId = blockId * blockDim.x + threadIdx.x;

    /*
     * shared value pair products array, where BLOCK_SIZE power of 2
     *
     * To improve performance increase its size by multiple of BLOCK_SIZE, so that each threads loads more then 1 element!
     * (outDots_g length decreases by same factor, and you need to range check and initialize memory)
     * -> see harris gpu optimisations / parallel reduction slides for more informations.
     */
    __shared__ double dots_s[BLOCK_SIZE];


    /*
     * initialize shared memory array and calculate dot product of two values, 
     * shared memory always needs to be initialized, its never 0 by default, else garbage is read later!
     */
    if(tId < inCount_s)
        dots_s[threadIdx.x] = inValuesA_g[tId] * inValuesB_g[tId];
    else
        dots_s[threadIdx.x] = 0;
    __syncthreads();

    //do parallel reduction on shared memory array to sum up values
    reductionAdd(dots_s, dots_s[0]) //see my thesis link

    //output value
    if(threadIdx.x == 0)
        outDots_g[0] = dots_s[0];

    //start new parallel reduction kernel to sum up outDots_g!
}

編集：不要なポイントを削除しました。

于 2012-02-26T04:08:30.450 に答える