0

これは、CUDA プログラムと並行開発での 2 回目の試みです。

内積が2つの配列になる理由を見つけようとしています。

例：

与えられた 

A = [1 2 3] 
B = [4 5 6]

C[0] = (1)(4) + (1)(5) + (1)(6)  
C[1] = (2)(4) + (2)(5) + (2)(6)  
C[2] = (3)(4) + (3)(5) + (3)(6)

2 つの配列 A と B を初期化し、範囲内の要素をランダムに埋めます。次に、各要素をそれぞれの要素で乗算しA、B積の合計を として識別される 3 番目の配列に格納しCます。A、B、C の配列のサイズを 100 に設定しました。

これにより、100 個のブロックと 128 個のスレッドを使用して並列化した 10,000 個の乗算が得られます (ワープ サイズのため)。

ここに私のカーネル関数があります:

__global__ void kernel(float *a, float *b, float *c, const int N) {
    if( threadIdx.x < N ) 
        c[blockIdx.x] += a[blockIdx.x] * b[threadIdx.x];
}

これは私の推論です。Cピボット インデックスと同じインデックスを持つ集計を蓄積する必要があるため、正しく動作するはずの をA再利用できます。blockidx.xしかし、そうではありません。

私の疑いはC、スレッドが変更されたときにインデックスがクリアされるか、共有されないということですが、アドバイスを求める理由が本当にわかりません。

これが完全なコードです。HANDLE_ERROR短くするために関数ラッパーを明示的に避けました

#include <stdio.h>
#include <cuda.h>
#include <time.h>

#define M 100

__global__ void kernel(float *a, float *b, float *c, const int N) {
    if(threadIdx.x < N) 
        c[blockIdx.x] += a[blockIdx.x] * b[threadIdx.x];
}

void init_array(float*, const int);
void fill_array(float*, const int, const float); 
void print_array(float*, const int, *char);

int main (void) {
    srand( time(NULL) );

    float a[M], b[M], c[M] = { 0.0 };
    float *dev_a, *dev_b, *dev_c;
    const int S = sizeof(float) * M;

    init_array(a, M);
    init_array(b, M);

    print_array(a, M, "a");
    print_array(b, M, "b");
    print_array(c, M, "c");

    cudaMalloc((void**)&dev_a, S);
    cudaMalloc((void**)&dev_b, S);
    cudaMalloc((void**)&dev_c, S);

    cudaMemcpy(dev_a, a, S, cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, S, cudaMemcpyHostToDevice);
    cudaMemcpy(dev_c, c, S, cudaMemcpyHostToDevice);

    kernel<<<M, M + 28>>>(dev_a, dev_b, dev_c, M);

    cudaMemcpy(c, dev_c, S, cudaMemcpyDeviceToHost);

    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);

    print_array(c, M, "c");

    return 0;
}

void init_array(float *a, const int N) {
   int i;  
   for(i=0; i<N; i++)
       a[i] = rand() % M + 1;
}

void fill_array(float *a, const int N, const float v) {
   int i;  
   for(i=0; i<N; i++)
       a[i] = v;
}

void print_array(float *a, const int N, char *d) {
   int i;  
   for(i=0; i<N; i++)
       printf("\n%s[%d]: %f",d, i, a[i]);
}
4

3 に答える 3

1

cuBLAS を使用する方がはるかに簡単です。このcublasSdot()ルーチンは、探していることを実行します (つまり、ドット積 2 ベクトル)。

これは、並列コードの記述方法や CUDA での作業方法を学ぶのには役立ちませんが、優れたパフォーマンスが得られ、さまざまな GPU 向けに最適化されます。ベスト プラクティスは、正当な理由がない限り、可能な限りライブラリを使用することです。

別の回答では、競合を回避するにはアトミックまたはその他の方法を使用する必要があると指摘されています。より効率的な方法は、各スレッドに部分的な結果を計算させ、ブロック単位のリダクションを実行させ (この例については SDK を参照)、最終的に各ブロックの 1 つのスレッドでグローバル メモリにアトミックな追加を実行させて、異なるブロックからの結果。

于 2012-09-18T09:16:06.033 に答える
1

結果を C に蓄積すると、相互排除が発生します。複数のスレッドで同じ配列インデックスを更新することはできません。それを修正する 1 つの方法は、atomicAdd(..) の場合はアトミック命令を使用することです。

これが機能しない理由は、ブロック 0 のスレッド 0 と 1 が C 配列の同じ位置を更新するためです。競合状態になります。

于 2012-09-18T09:03:11.177 に答える
-2

コードは、グローバル メモリ内の同じ場所に複数の値を書き込みますが、同期や考慮は行われません。ある種のクリティカルセクションを使用して修正できます。これがあなたが望むことをしているコードです：

#include <stdio.h>
#include <cuda.h>
#include <time.h>

#define M 5

__global__ void kernel(float *a, float *b, float *c, const int N) {
    __shared__ int lock;
    lock=0u;
    if(threadIdx.x < N) {
        __syncthreads();
        float mult = a[blockIdx.x] * b[threadIdx.x];

        int leaveLoop = 0;
        while (leaveLoop==0) {
            if (atomicExch(&lock, 1u) == 0u) {
                //critical section
                c[blockIdx.x] += mult;
                leaveLoop = 1;
                atomicExch(&lock,0u);
            }
        }
        __syncthreads();
    }
}

void init_array(float *a, const int N) {
    int i;  
    for(i=0; i<N; i++)
        a[i] = rand() % M + 1;
}

void fill_array(float *a, const int N, const float v) {
    int i;  
    for(i=0; i<N; i++)
        a[i] = v;
}

void print_array(float *a, const int N, char *d) {
    int i;  
    for(i=0; i<N; i++)
        printf("\n%s[%d]: %f",d, i, a[i]);
}

int main (void) {
    srand( time(NULL) );

    float a[M], b[M], c[M] = { 0.0 };
    float *dev_a, *dev_b, *dev_c;
    const int S = sizeof(float) * M;

    init_array(a, M);
    init_array(b, M);

    print_array(a, M, "a");
    print_array(b, M, "b");
    print_array(c, M, "c");

    cudaMalloc((void**)&dev_a, S);
    cudaMalloc((void**)&dev_b, S);
    cudaMalloc((void**)&dev_c, S);

    cudaMemcpy(dev_a, a, S, cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, S, cudaMemcpyHostToDevice);
    cudaMemcpy(dev_c, c, S, cudaMemcpyHostToDevice);

    kernel<<<M, M + 28>>>(dev_a, dev_b, dev_c, M);

    cudaMemcpy(c, dev_c, S, cudaMemcpyDeviceToHost);

    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);

    print_array(c, M, "c");

    return 0;
}

配列の乗算の合計を集計するという目標を達成するには、可能であれば @Tom によって提案されたものを使用することをお勧めします。外部ライブラリの代わりにコードを使用する方が簡単な場合があるため、このコードを記述します。

于 2012-09-18T09:49:13.010 に答える