2

共有メモリを使用して、cuda c で内積に関する簡単なチュートリアルを実行しようとしています。コードは非常に単純で、基本的に 2 つの配列の要素間の積を計算し、各ブロックの結果を合計します。

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <cuda.h>

#define imin(a,b) (a<b?a:b)

const int N = 33*1024;
const int threadsPerBlock = 256;
const int blocksPerGrid = imin(32 , (N+threadsPerBlock-1)/threadsPerBlock);

__global__ void dot(float *a, float *b, float *c){

__shared__ float cache[threadsPerBlock];
int tid = threadIdx.x + blockIdx.x*blockDim.x; 
int cacheIndex = threadIdx.x; 
float temp = 0; 
while (tid < N){
    temp += a[tid]*b[tid];
    tid += blockDim.x*gridDim.x; /* Aggiorno l'indice per l'evenutale overshoot. */
}
cache[cacheIndex] = temp; 
__syncthreads(); 

int i = blockDim.x/2; 
while(i != 0){ /
    if(cacheIndex < i){ 
        cache[cacheIndex] += cache[cacheIndex + i];
        __syncthreads(); 
        i /= 2; 
    }      
}
if(cacheIndex == 0){ 
    c[blockIdx.x] = cache[0]; 
}
}

int main(void){

cudaError_t err = cudaSuccess; 

float a[N], b[N], c[blocksPerGrid];
float *d_a, *d_b, *d_c;   
int i;
for(i=0;i<N;i++){
    a[i] = i;
    b[i] = i*2;  
}
for(i=0; i<blocksPerGrid;i++){
    c[i] = 0;
}

err = cudaMalloc((void**)&d_a, N*sizeof(float));
if (err != cudaSuccess){fprintf(stderr, "Failed to allocate device vector a (error code %s)! \n", cudaGetErrorString(err));exit(EXIT_FAILURE);}
err = cudaMalloc((void**)&d_b, N*sizeof(float));
if (err != cudaSuccess){fprintf(stderr, "Failed to allocate device vector b (error code %s)! \n", cudaGetErrorString(err));exit(EXIT_FAILURE);}
err = cudaMalloc((void**)&d_c, blocksPerGrid*sizeof(float));
if (err != cudaSuccess){fprintf(stderr, "Failed to allocate device vector c (error code %s)! \n", cudaGetErrorString(err));exit(EXIT_FAILURE);}
/* Copio i valori dei vettori a e b nello spazio di memoria allocato precedentemente nel device. */
err = cudaMemcpy(d_a, a, N*sizeof(float), cudaMemcpyHostToDevice);
if (err != cudaSuccess){fprintf(stderr, "Failed to copy vector a from host to device (error code %s)!\n", cudaGetErrorString(err));exit(EXIT_FAILURE);}
err = cudaMemcpy(d_b, b, N*sizeof(float), cudaMemcpyHostToDevice);
if (err != cudaSuccess){fprintf(stderr, "Failed to copy vector b from host to device (error code %s)!\n", cudaGetErrorString(err));exit(EXIT_FAILURE);}
err = cudaMemcpy(d_c, c, blocksPerGrid*sizeof(float), cudaMemcpyHostToDevice);
if (err != cudaSuccess){fprintf(stderr, "Failed to copy vector c from host to device (error code %s)!\n", cudaGetErrorString(err));exit(EXIT_FAILURE);}


dot<<<blocksPerGrid,threadsPerBlock>>>(d_a, d_b, d_c); err = cudaGetLastError();

err = cudaMemcpy(c, d_c, blocksPerGrid*sizeof(float), cudaMemcpyDeviceToHost);
if (err != cudaSuccess){fprintf(stderr, "Failed to copy vector c from device to host (error code %s)!\n", cudaGetErrorString(err));exit(EXIT_FAILURE);}

err = cudaFree(d_a); 
if (err != cudaSuccess){fprintf(stderr, "Failed to free device vector a (error code %s)!\n", cudaGetErrorString(err));exit(EXIT_FAILURE);}
err = cudaFree(d_b); 
if (err != cudaSuccess){fprintf(stderr, "Failed to free device vector b (error code %s)!\n", cudaGetErrorString(err));exit(EXIT_FAILURE);}
err = cudaFree(d_c); 
if (err != cudaSuccess){fprintf(stderr, "Failed to free device vector c (error code %s)!\n", cudaGetErrorString(err));exit(EXIT_FAILURE);}


float result = 0;
for(i=0;i<blocksPerGrid;i++){
    result += c[i];
}

printf("il risultato finale è: %.2f\n", result);

return 0;
}

このコードは、Cuda by Example book で提示されているものと同じで、唯一の違いはベクトル a、b、および c の定義にあります (定義方法は問題にならないはずです。 ）。

問題は次のとおりです。プログラムを実行しようとすると、クラッシュします。端末は、問題は次のとおりです。Failed to copy vector c from device to host (error code the launch timed out and was terminated)!

ベクトルcを適切な方法で割り当てたと思うので、それは奇妙です...誰かが私が間違っていることを知っていますか? 何か問題があるのはグローバル関数ですか、それともメインですか?

4

1 に答える 1

2

カーネルには無限ループがあります。エラーが発生する理由は、ウォッチドッグ タイムアウトのあるプラットフォームを使用していて、ウォッチドッグがカーネルの実行を強制終了しているためです。

このコードを検討してください：

int i = blockDim.x/2; 
while(i != 0){
    if(cacheIndex < i){ 
        cache[cacheIndex] += cache[cacheIndex + i];
        __syncthreads(); 
        i /= 2; 
    }      
}

が より小さいi場合、ループ インデックス ( ) を 2で除算するだけです。他のスレッドの場合、そのスレッドが if ステートメントから脱落すると、常に同じ値のままになります。これらのスレッドの場合、while ループは終了しません (がゼロになることはありません)。すべてのスレッドの変数を分割したいとします。そのようです：cacheIndexiiii

int i = blockDim.x/2; 
while(i != 0){
    if(cacheIndex < i){ 
        cache[cacheIndex] += cache[cacheIndex + i];
    }   
    __syncthreads(); 
    i /= 2;    
}

__syncthreads()if ステートメントの外にも移動したことに注意してください。これは問題を解決するために必要ではないかもしれませんが、通常はすべてのスレッドが__syncthreads()ステートメントに参加する必要があるため、技術的には正しくありません。条件がすべてのスレッドで同じと評価される場合にのみ、条件付きコードで許可されます。これは、プログラミング ガイドに記載されています。

そして、この点であなたのコードを第 5 章の dot.cu の例のソースコードで cuda にあるものと比較すると、それらが同一ではないことがわかると思います。

于 2014-11-10T22:15:09.077 に答える