6

私はCUDAを初めて使用します。次のコードを並列化しようとしています。現在、カーネル上にありますが、スレッドをまったく使用していないため、低速です。私はこの答えを使おうとしましたが、今のところ役に立ちません。

カーネルは最初のn個の素数を生成し、それらをdevice_primes配列に入れ、この配列は後でホストからアクセスされることになっています。コードは正しく、シリアルバージョンでは正常に動作しますが、おそらく共有メモリを使用して、コードを高速化する必要があります。

//CUDA kernel code
__global__ void generatePrimes(int* device_primes, int n) 
{
//int i = blockIdx.x * blockDim.x + threadIdx.x;
//int j = blockIdx.y * blockDim.y + threadIdx.y;

int counter = 0;
int c = 0;

for (int num = 2; counter < n; num++)
{       
    for (c = 2; c <= num - 1; c++)
    { 
        if (num % c == 0) //not prime
        {
            break;
        }
    }
    if (c == num) //prime
    {
        device_primes[counter] = num;
        counter++;
    }
}
}

これを並列化するための私の現在の、予備的な、そして間違いなく間違った試みは、次のようになります。

//CUDA kernel code
__global__ void generatePrimes(int* device_primes, int n) 
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    int num = i + 2; 
    int c = j + 2;
    int counter = 0;

    if ((counter >= n) || (c > num - 1))
    {
        return;
    }
    if (num % c == 0) //not prime
    {
    
    }
    if (c == num) //prime
    {
       device_primes[counter] = num;
       counter++;
    }
    num++;
    c++;
}

ただし、このコードは、意味のないデータを配列に入力します。さらに、多くの値はゼロです。助けてくれてありがとう、それはありがたいです。

4

1 に答える 1

4

コードにいくつかの問題があります。次に例を示します。

int num = i + 2;

thread 0この式は、相互作用2、thread 1反復3などに割り当てられます。問題は、スレッドが計算する次の反復が式に基づいていることnum++;です。したがって、thread 0は、によってすでに計算されている反復3を次に計算しthread 1ます。したがって、冗長な計算につながります。さらに、この問題では、2次元ではなく1次元のみを使用する方が簡単だと思います(x,y)。したがって、これを念頭に置いて、次のように変更num++する必要があります。

num += blockDim.x * gridDim.x;

counterもう1つの問題は、変数をスレッド間で共有する必要があることを考慮していなかったことです。それ以外の場合、各スレッドは「n」個の素数を見つけようとし、それらすべてが配列全体に入力されます。int counter = 0;したがって、共有変数またはグローバル変数に変更する必要があります。グローバル変数を使用して、すべてのブロックのすべてのスレッドで表示できるようにします。配列の位置0を使用してdevice_primes、変数を保持できますcounter。

また、この値を初期化する必要があります。このジョブを1つのスレッド、つまり `id = 0のスレッドにのみ割り当てましょう。つまり、次のようになります。

if (thread_id == 0) device_primes[0] = 1;

ただし、この変数はグローバルであり、すべてのスレッドによって書き込まれます。したがって、すべてのスレッドがそのグローバル変数に書き込む前に、変数counterが1（device_primes素数の最初の位置、ゼロはcounter）であることを確認する必要があるため、最後にバリアも追加する必要があります。

if (thread_id == 0) 
    device_primes[0] = 1;
__syncthreads()

したがって、考えられる解決策（ただし、非効率的な解決策）：

__global__ void getPrimes(int *device_primes,int n)
{ 
    int c = 0;
    int thread_id = blockIdx.x * blockDim.x + threadIdx.x;
    int num = thread_id;

    if (thread_id == 0) device_primes[0] = 1;
    __syncthreads();

    while(device_primes[0] < n)
    {

        for (c = 2; c <= num - 1; c++)
        { 
            if (num % c == 0) //not prime
            {
                break;
            }
        }

        if (c == num) //prime
        {
            int pos = atomicAdd(&device_primes[0],1);
            device_primes[pos] = num;

        }

        num += blockDim.x * gridDim.x; // Next number for this thread       
    }
}

次の行atomicAdd(&device_primes[0], 1);は基本的に実行されますdevice_primes[0]++;。変数counterがグローバルであり、相互排除を保証する必要があるため、アトミック操作を使用しています。を使用してコンパイルする必要がある場合があることに注意してくださいflag -arch sm_20。

最適化：コードに関しては、同期が少ない/ないアプローチを使用する方がよいでしょう。さらに、 http：//en.wikipedia.org/wiki/Sieve_of_Eratosthenesで示されているように、素数のいくつかの特性を考慮に入れることで、計算の数を減らすこともできます。

于 2012-11-04T19:35:49.723 に答える