0

1 <= x <= 10^7 の場合、x^2+1 の形式の数がいくつ素数であるかを判断したかったのです。CUDAで並列化して違いを確認したかっただけなので、単純な素数チェックを使用しましたが、アルゴリズムの改善には関心がありません。

グリッドを配置し、間隔を空けてスレッド化し、結果を各ブロックの共有メモリに記録し、各ブロックで GPU の削減を実行し、最後に CPU の削減を実行して最終結果を取得しました。

私の問題は、ブロック数と各ブロック内のスレッド数を変更すると、出力結果が変わることです。説明できないもう 1 つのことは、8 ブロックとブロックあたり 2048 スレッドの構成の場合、コードは 100 ミリ秒未満で実行されますが、スレッド数を 1024 に減らしてブロック数を 2 倍にすると、コードがタイムアウトを引き起こすということです。デバイスからホストへのmemcpyで!! この動作と、正確性が問題になる場所をどのように説明できますか?

GTX 480 nvidia GPU を使用しています。

私のコードは次のとおりです。 

#include <stdio.h>
static void HandleError( cudaError_t err, const char *file, int line )
{
    if (err != cudaSuccess) {
        printf( "%s in %s at line %d\n", cudaGetErrorString( err ), file, line );
        exit( EXIT_FAILURE );
    }
}

#define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))
#define N 10000000
#define BLOCKS 8
#define THREADS 2048

__device__ int isprime(int x)
{
    long long n = (long long)x*x + 1;
    for( int p=3; p<=x+1; p+=2 )
        if ( n % p == 0 ) return 0;
    return 1;
}

__global__ void solve(int n, int* result)
{
    __shared__ int ipc[THREADS];

    int tid = threadIdx.x;
    int x = blockIdx.x*blockDim.x + threadIdx.x + 2;

    // sliding grid window over interval of to-be-computed data
    int acc = 0;
    while( x <= n )
    {
        if ( isprime(x) ) acc++;
        x += blockDim.x*gridDim.x;
    }
    ipc[tid] = acc;
    __syncthreads();


    // reduction over each block in parallel
    for( int s=blockDim.x/2; s>0; s>>=1 )
    {
        if ( tid < s )
        {
            ipc[tid] += ipc[tid+s];
        }
        __syncthreads();
    }

    if ( tid == 0 ) result[blockIdx.x] = ipc[0];
}

int main()
{
    int *dev;
    int res[BLOCKS];

    int ans = 0;

    HANDLE_ERROR( cudaMalloc((void**)&dev, BLOCKS * sizeof(int)) );

    solve<<<BLOCKS, THREADS>>>(N, dev);

    HANDLE_ERROR( cudaMemcpy(res, dev, BLOCKS*sizeof(int), cudaMemcpyDeviceToHost) );

    // final reduction over results for each block
    for( int j=0; j<BLOCKS; j++ )
        ans += res[j];

    printf("ans = %d\n", ans);

    HANDLE_ERROR( cudaFree( dev ) );
    return 0;
}
4

1 に答える 1

5

現在の GPU では、ブロックごとに 2048 スレッドを実行することはできません。

#define THREADS 2048
...
solve<<<BLOCKS, THREADS>>>(N, dev);
                  ^
                  |
                2048 is illegal here

カーネル呼び出しで適切なcuda エラー チェックを行っていないため、コードはこのエラーが発生していることを通知しません。

したがって、ブロックごとに 2048 スレッドの場合、カーネルは実行さえしていません (結果は偽物になるはずです)。

スレッドを半分にカットした場合、タイムアウトはおそらくカーネルの実行に時間がかかりすぎており、Windows TDR メカニズムが作動していることが原因です。

BLOCKS= 16 およびTHREADS= 1024でコードを実行してみました

N = 100000 の場合、M2050 GPU での合計実行時間は約 1.5 秒でした。N=1000000で、実行時間は約75秒でした。あなたが持っているN = 10000000では、実行時間が非常に長くなります。

于 2013-07-09T00:07:00.777 に答える