大型数组的 C++/CUDA 奇怪行为答案

【问题标题】：C++/CUDA weird behavior with large arrays大型数组的 C++/CUDA 奇怪行为
【发布时间】：2020-12-15 10:32:45
【问题描述】：

我尝试实现某种雅可比算法并测量不同网格大小所花费的时间。

对于具有相同数量的迭代，无论网格有多大，我都不使用某种残差，而是让算法始终运行 4000 次迭代（但数组大小不同）。这完全可以正常工作，直到我超过 510x510 网格（双倍）。 510x510 大约需要 2763498 微秒，然后 520x520 需要 1778 微秒。

我已经尝试从 double 数组更改为 float 数组，以确保它不是某种内存短缺，但我不知道我的问题真正隐藏在哪里。

__global__ void Jacobi(double *a, double *b, double *c, int L){
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;

if(row > 0 && col > 0 && row < L-1 && col < L-1){
    a[row * L + col] = 1.0/4.0 * (b[col+1 + row*L] + b[col - 1 + row*L] + b[col + (row+1)*L] + b[col + (row-1)*L] - c[col + row*L]);
    __syncthreads();
    b[row*L + col] = a[row*L+col];
    }
}

int main(){
int L;
int Iterations;
double *h_phi1;
double *h_phi2;

double *h_f;
FILE * temp = fopen("Timings.out", "w");
for (L=10;L<10000;L+=10){
    long long int size = L*L*sizeof(double);
    h_f = (double*) malloc(size);
    h_phi1 = (double*) malloc(size);
    h_phi2 = (double*) malloc(size);


    for(int i=0;i<L;i++){
        for(int j=0;j<L;j++){
            h_f[j+i*L] = (pow(1.0/float(L),2.0))*exp(-100.0*(pow((float(i)/float(L) - float(1.0/3.0) ),2.0) + pow(( float(j)/float(L) - float(1.0/3.0) ),2.0))) - 
            (pow(1.0/ float(L),2.0))*exp(- 100.0*(pow(( float(i)/ float(L) -(1.0- 1.0/3.0)),2.0) + pow(( float(j)/float(L)-(1.0-float(1.0/3.0))),2.0)));
            h_phi1[j+i*L] = 0;
            h_phi2[j+i*L] = 0;
        }
    }

    //allocate memory on GPU
    double *d_phi1;
    cudaMalloc(&d_phi1, size);
    double *d_phi2;
    cudaMalloc(&d_phi2, size);
    double *d_f;
    cudaMalloc(&d_f, size);
    
    //set CTA
    int threads = 16;
    int blocks = (L+threads-1)/threads;
    double epsCalc;

    //Setup Kernel launch parameters
    dim3 dimBlock(threads, threads);
    dim3 dimGrid(blocks, blocks);

    //Setup timing and Cpy Memory from Host to Device
    Iterations = 0;
    auto t1 = std::chrono::high_resolution_clock::now();
    cudaMemcpy(d_phi2, h_phi2, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_f, h_f, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_phi1, h_phi2, size, cudaMemcpyHostToDevice);

    //Launch Kernel
    for (int j=0;j<4000;j++){
        Iterations += 1;
        Jacobi<<<dimBlock, dimGrid>>>(d_phi2,d_phi1,d_f,L);

    }
    auto t2 = std::chrono::high_resolution_clock::now();
    auto duration = std::chrono::duration_cast<std::chrono::microseconds>( t2 - t1 ).count();
    
    fprintf(temp, "%lf % %d \n", L, duration);
    printf("I reached the end of Jacobi after %d Iterations!\n Time taken was %d in milliseconds", Iterations, duration);
    cudaFree(d_f); cudaFree(d_phi2), cudaFree(d_phi1);
    free(h_f); free(h_phi2); free(h_phi1);
}

return 0;
}

我希望有人能指导我发现我的错误所在。

【问题讨论】：

如果你担心这是否是内存问题，你真的应该测试malloc是否返回0。
你的代码中的时间测量被破坏了。
我不认为，时间测量是错误的。执行速度确实快得多..从 (520,530....) 开始的迭代时间确实再次增加。由于对 cudaLaunchKernel 的 CUDA API 调用中的“无效配置参数”，Cuda-Memcheck 导致“”程序命中 cudaErrorInvalidConfiguration（错误 9）。“”“”在 520。之后它再次运行，没有问题。我不明白它
错误地写入毫秒而不是微秒。对不起@molbdnilo
我也认为您的时序测量被破坏了，此外您可能想了解异步内核启动的性质，因为这也会影响您的时序测量.根据我的测试，您的时序测量至少有两个单独的问题。我知道您可能不同意；随心所欲。

标签： c++ algorithm cuda large-data

【解决方案1】：

在CUDA中，当指定kernel execution configuration arguments时，首先是网格配置，其次是块配置。

因此，鉴于变量dimGrid和dimBlock的定义和用法，这是不正确的：

    Jacobi<<<dimBlock, dimGrid>>>

应该是：

    Jacobi<<<dimGrid, dimBlock>>>

网格和块配置变量具有不同的限制，因此由于混合违反硬件限制而发生的第一次启动失败发生在问题维度 520,520

【讨论】：