cuda - CUDA: sum-reduction --- デバイス関数の呼び出しでデータが失われました

Question

要素 begin_index から end_index までの配列の絶対値の合計を取る CUDA 合計削減コードを作成しています (可変数のスレッドを持つ 1 つのブロックを使用しています)。ただし、配列Aをに渡すとreduce_fabs()、以降のすべてのA[4]インデックスがどういうわけかゼロに設定されます。これがコードと関数呼び出しであり、私が試したことの詳細な説明です。

これは sum-reduction カーネルを呼び出すカーネルです。

__device__ void tridiag(float *A,int *dim, float *diag,float *offdiag) {

    A[0] = 1.0f; A[1] = 2.0f; A[2] = 3.0f;
    A[3] = 4.0f; A[4] = 5.0f; A[5] = 6.0f;
    diag[0] = reduce_fabs(A,0,3);
    __syncthreads();
    diag[1] = reduce_fabs(A,0,4);

    return;
}

これは合計削減カーネルです。

__device__ float reduce_fabs(float *v, int begin_index, int end_index) {
    extern __shared__ float sum_array[];
    int tid = threadIdx.x;
    if(tid >= begin_index && tid <= end_index) {
        sum_array[tid-begin_index] = fabs(v[tid]);
        sum_array[tid+end_index-begin_index+1] = 0;
    }
    __syncthreads();
    for(int j=1;j<=(end_index-begin_index);j*=2) {
        if((tie-begin_index)%(2*j) == 0 && tie >= begin_index && tid <= end_index) {
            sum_array[(tie-begin_index)] += sum_array[(tie-begin_index)+j];
        }
        __syncthreads();
    }

    return sum_array[0];
}

コードを見た後、問題のより具体的な説明は、配列Aをに渡すreduce_fabs()と、要素 4 の値A[4] = 0が間違っているということです。A[0]、A[1]、A[2]、およびA[3]は任意の に対して問題ありませんend_indexが、end_index3 より大きい場合はA、リダクション カーネルへの通過時に 3 を超える要素がゼロに設定されます。

これが私がすでに試したことです：

• diag[0] = A[4]最初の割り当てが機能していることを確認しようとしました。そうだった。
• の合計部分を削除しreduce_fabs()、最初の後に停止しましたが__syncthreads()、問題は解決しませんでした。
• 合計に関心のある要素を超える要素のゼロ化を排除しました。つまり//sum_array[tid+end_index-begin_index+1] = 0、コメントしました (この部分の合計もコメントアウトされました)。失敗。
• if(tie >= begin_index && tid <= end_index)すべてのスレッドに割り当てられるように削除sum_arrayしました（合計もコメントアウトされています）が無駄になりました。

これを実行するためのメイン関数は次のとおりです。

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

int main() {
    int n = 10;
    float *A = (float *)malloc(n*sizeof(*A));
    float *diag = (float *)malloc(n*sizeof(*A));
    float *offdiag = (float *)malloc(n*sizeof(*A));

    int *p_n;
    float *p_A, *p_diag, *p_offdiag;

    cudaMalloc((void**) &p_A,n*sizeof(float));
    cudaMalloc((void**) &p_diag,n*sizeof(float));
    cudaMalloc((void**) &p_offdiag,n*sizeof(float));
    cudaMalloc((void**) &p_n,sizeof(int));

    cudaMemcpy(p_n,&n,n*sizeof(int),cudaMemcpyHostToDevice);

    tridiag<<<1,n>>>(p_A,p_n,p_diag,p_offdiag);

    cudaMemcpy(A,p_A,n*sizeof(float),cudaMemcpyDeviceToHost);
    cudaMemcpy(diag,p_diag,n*sizeof(float),cudaMemcpyDeviceToHost);

    printf("A[0] = %f  A[1] = %f  A[2] = %f\n",A[0],A[1],A[2]);
    printf("A[3] = %f  A[4] = %f  A[5] = %f\n",A[3],A[4],A[4]);
    printf("diag[0] = %f  diag[1] = %f\m",diag[0],diag[1]);

    cudaFree(p_A);
    cudaFree(p_diag);
    cudaFree(p_offdiag);
    free(A);
    free(diag);
    free(offdiag);

    return 0;
}