1

私の以前の回答済みの質問: My Previous Questionを見てください。ちなみに、これはRobert Crovellaによって適切に回答されました。

ポイントへのランダムなステップを計算し (前の質問と同じ RNG を使用して)、そのポイントの以前の位置 (座標) に対するエネルギーの差を計算する別のカーネルを思いつきました。これはカーネルです：

__global__
void DeltaE(float *X, float *Y,float *Xn, float *Yn, float *step,float *DELTA, curandState *state, const int N,const int n){

    int tIdx = blockIdx.x*blockDim.x + threadIdx.x;
    int bIdx = blockIdx.x;
    //int sIdx = blockIdx.x*blockDim.x;

    float x , y;
    float rdmn1, rdmn2;

    float dIntE = 0.0e0f, dConfE = 0.0e0f, dTotE = 0.0e0f;
    if(tIdx < N){

        if(tIdx == n){
            step[tIdx] = 0.2;
            rdmn1 = curand_uniform(&state[tIdx]);
            rdmn2 = curand_uniform(&state[tIdx]);

            Xn[tIdx] = X[tIdx] + step[tIdx]*(2.0e0f*rdmn1 - 1.0e0f);
            Yn[tIdx] = Y[tIdx] + step[tIdx]*(2.0e0f*rdmn2 - 1.0e0f);
            dConfE = - (X[tIdx]*X[tIdx] + Y[tIdx]*Y[tIdx]);
            dConfE += Xn[tIdx]*Xn[tIdx] + Yn[tIdx]*Yn[tIdx];

        }
        else{
            x = X[tIdx] - X[n];
            y = Y[tIdx] - Y[n];

            dIntE += -1.0e0f/sqrt(x*x + y*y);
        }
        __syncthreads();
        if(tIdx != n){
            x = X[tIdx] - Xn[n];
            y = Y[tIdx] - Yn[n];

            dIntE += 1.0e0f/sqrt(x*x + y*y);
        }       
        dTotE = dConfE + dIntE;
        dTotE = ReduceSum2(dTotE);
        if(threadIdx.x == 0)DELTA[bIdx] = dTotE;


    }
}

次に、CPU で最終的な合計を計算します。

cudaMemcpy(&delta[0],&d_delta[0],blocks.x*sizeof(float),cudaMemcpyDeviceToHost);
float dE = 0;
for(int i = 0; i < blocks.x; i++){
    dE += delta[i];
}

私のカーネルは、次の構成で起動されます。

dim3 threads(BlockSize,BlockSize);
dim3 blocks(ceil(Np/threads.x),ceil(Np/threads.y));
DeltaE<<<blocks.x,threads.x,threads.x*sizeof(float)>>>(d_rx,d_ry,d_rxn,d_ryn,d_step,d_delta,d_state,Np,nn);

ここで、Np はポイントの数です (私は 1k - 4k を使用しました)。私は GeForce 9500 GT を持っていますが、これは倍増をサポートしていません。そして、フラグなし/オプションなしを使用してコンパイルします。

たとえば、Np = 1k を取ります。コンパイルして実行すると、結果は dE = 6.557993 になります。2 回目、3 回目、4 回目、いつ実行しても、dE = -0.3515406 です。これがどこから来たのか誰か知っていますか？

PS: 言い忘れましたが、前の質問にあるのと同じカーネル AvgDistance が DeltaE の直前に呼び出されます。これが何か関係があるかどうかはわかりませんが、言及する価値があると思いました。

PS2: nn は任意の選択されたポイント (粒子) です。

4

1 に答える 1

2

上記のコメントでRobert Crovellaによって指摘されたように、おそらく起こっていたのは whiletIdx = nが計算Xn[n]していてYn[n]、他のスレッドがこの値を使用していて、まだ計算されていない可能性があるということです。その場合、他の実行 (最初の実行以外) が同じ (正しい) 値を取得する唯一の理由は、Xn と Yn が指すメモリが既に正しい値で占有されており、その同期でも使用されているためです。アプリケーションが正しい値を返す問題。

いずれにせよ、 Robert Crovellaからコメントでアドバイスされたように、カーネルを 2 つに分割することで同期の問題を回避しました。

__global__
void DeltaE1(float *X, float *Y,float *Xn, float *Yn, float *step,float *DELTA, curandState *state, const int N,const int n){

    int tIdx = blockIdx.x*blockDim.x + threadIdx.x;
    float x , y;
    float rdmn1, rdmn2;

    if(tIdx < N){
        DELTA[tIdx] = 0.0e0f;
        if(tIdx == n){
            step[tIdx] = 0.2e0f;
            rdmn1 = curand_uniform(&state[tIdx]);
            rdmn2 = curand_uniform(&state[tIdx]);

            Xn[tIdx] = X[tIdx] + step[tIdx]*(2.0e0f*rdmn1 - 1.0e0f);
            Yn[tIdx] = Y[tIdx] + step[tIdx]*(2.0e0f*rdmn2 - 1.0e0f);
            DELTA[tIdx] = - (X[tIdx]*X[tIdx] + Y[tIdx]*Y[tIdx]);
            DELTA[tIdx] += Xn[tIdx]*Xn[tIdx] + Yn[tIdx]*Yn[tIdx];

        }
        else{
            x = X[tIdx] - X[n];
            y = Y[tIdx] - Y[n];

            DELTA[tIdx] += -1.0e0f/sqrt(x*x + y*y);
        }           
    }
}

__global__
void DeltaE2(float *X, float *Y,float *Xn, float *Yn,float *DELTA,const int N,const int n){

    int tIdx = blockIdx.x*blockDim.x + threadIdx.x;
    int bIdx = blockIdx.x;

    float x , y;
    float dTotE = 0.0e0f;
    if(tIdx < N){
        if(tIdx != n){
            x = X[tIdx] - Xn[n];
            y = Y[tIdx] - Yn[n];

            DELTA[tIdx] += 1.0e0f/sqrt(x*x + y*y);

        }
        dTotE = DELTA[tIdx];
        dTotE = ReduceSum2(dTotE);
        if(threadIdx.x == 0)DELTA[bIdx] = dTotE;

    }

}
于 2013-03-16T10:56:57.423 に答える