1

我对以下程序（下面的代码）感到困惑。当内核中的两行定义specsin和时，它工作正常并给出正确的结果speccos（注意第二项，即sin(t)）：

specsin+=sin(pi*t/my_tau)*sin(t)*sin(my_omega*(t+my_a0*my_a0/4.0/pi*(2.0*pi*t-my_tau*sin(2*pi*t/my_tau))));
speccos+=sin(pi*t/my_tau)*sin(t)*cos(my_omega*(t+my_a0*my_a0/4.0/pi*(2.0*pi*t-my_tau*sin(2*pi*t/my_tau))));

一旦我将第二sin(t)项更改为sin(t+0.0*my_a0*my_a0)，这不应该改变结果，我得到全零而不是正确答案。

难道是我的内核内存用完了？

#include <stdio.h>
__global__ void Calculate_Spectrum(float * d_Detector_Data, int numCols, int numRows,
                                const float omega_min, float dOmega,
                               const float a0_min, float da0,
                               const float tau_min, float dtau, float dt)
{

    int Global_x = blockIdx.x * blockDim.x + threadIdx.x;
    int Global_y = blockIdx.y * blockDim.y + threadIdx.y;
    int Position1D = Global_y * numCols + Global_x;

    float my_omega=omega_min + Global_x * dOmega;
    float my_a0=a0_min + Global_y*da0;
    float my_tau=tau_min;
    int total_time_steps=int(my_tau/dt);

    float specsin=0.0;
    float speccos=0.0;
    float t=0.0;
    float pi=3.14159265359;

    for(int n=0; n<total_time_steps; n++)
    {
        t=n*dt;
        specsin+=sin(pi*t/my_tau)*sin(t+0.0*my_a0*my_a0)*sin(my_omega*(t+my_a0*my_a0/4.0/pi*(2.0*pi*t-my_tau*sin(2*pi*t/my_tau))));
        speccos+=sin(pi*t/my_tau)*sin(t+0.0*my_a0*my_a0)*cos(my_omega*(t+my_a0*my_a0/4.0/pi*(2.0*pi*t-my_tau*sin(2*pi*t/my_tau))));
    }

    d_Detector_Data[Position1D]=(specsin*specsin+speccos*speccos)*dt*dt*my_a0*my_a0*my_omega*my_omega/4.0/pi/pi;
}


int main(int argc, char ** argv)
{
    const int omega_bins = 1024;
    const int a0_bins = 512;
    const int tau_bins = 1;

    const float omega_min = 0.5;
    const float omega_max = 1.1;
    const float a0_min = 0.05;
    const float a0_max = 1.0;
    const float tau_min = 1200;
    const float tau_max = 600;

    const int steps_per_period=20;  // for integrating
    float dt=1.0/steps_per_period;

    int TotalSize = omega_bins * a0_bins * tau_bins;

    float dOmega=(omega_max-omega_min)/(omega_bins-1);
    float da0=(a0_max-a0_min)/(a0_bins-1);
    float dtau=0.;

    float * d_Detector_Data;
    int * d_Global_x;
    int * d_Global_y;


float h_Detector_Data[TotalSize];

    // allocate GPU memory
    cudaMalloc((void **) &d_Detector_Data, TotalSize*sizeof(float));

    Calculate_Spectrum<<<dim3(1,a0_bins,1), dim3(omega_bins,1,1)>>>(d_Detector_Data, omega_bins, a0_bins, omega_min, dOmega, a0_min, da0, tau_min, dtau, dt);


cudaMemcpy(h_Detector_Data, d_Detector_Data, TotalSize*sizeof(float), cudaMemcpyDeviceToHost);

    FILE * SaveFile;
    char TempStr[255];

    sprintf(TempStr, "result.dat");
    SaveFile = fopen(TempStr, "w");

    int counter=0;

    for(int j=0; j<a0_bins;j++)
    {
        for(int i=0; i<omega_bins; i++)
        {
    fprintf(SaveFile,"%e\t", h_Detector_Data[counter]);
    counter++; 

        }
    fprintf(SaveFile, "\n");
    }
    fclose(SaveFile);

// free GPU memory
return 0;

}
4

1 回答 1

3

我相信这是由于注册限制。

为了启动内核，每个线程的寄存器总数不得超过最大限制（即Maximum number of 32-bit registers per thread编译器应保证的），每个线程的寄存器乘以请求的线程数不得超过最大限制（the Number of 32-bit registers per multiprocessor）。

在您得到不正确结果的情况下，我相信您的内核由于这个原因没有启动（总共请求的寄存器太多）。您没有进行任何cuda 错误检查，但如果您这样做了，我相信您可以确认这一点。

您可以使用任何将总数减少到限制以下的方法来解决此问题。显然，减少每个块的线程数是一种直接的方法。指定编译器的开关等其他事情-G也会影响代码生成，因此可能会影响每个线程的寄存器。解决此问题的另一种方法是指示编译器将其寄存器的使用限制为每个线程的某个最大数量。这在nvcc 手册中有记录，用法是这样的：

nvcc -maxrregcount=xx  ... (rest of compile command)

xx每个线程限制使用的寄存器数量在哪里。如果您将其限制为假设每个线程 20 个，那么即使每个块有 1024 个线程，我仍然只会使用大约 20K 寄存器，这将适合任何支持每块 1024 个线程的设备（cc 2.0 及更高版本）。

您还可以通过要求编译器生成寄存器使用统计信息：

nvcc -Xptxas -v ... (rest of compile command)

这将导致编译器生成有关资源使用情况的各种统计信息，包括每个线程使用/预期的寄存器数量。

请注意，资源使用（包括寄存器使用）可能会影响占用率，这对整体应用程序性能有影响。因此，限制寄存器的使用不仅可以让内核运行，还可以允许多个线程块驻留在一个 SM 上，这通常表明您的占用率有所提高，这可能会提高您的应用程序的性能。

显然0.0vs.的使用0.0f对编译器行为有一些微妙的影响，这在代码生成中出现了。我还推测您可能处于可接受的边界，因此每个线程使用的寄存器的微小变化可能会影响将要运行的内容。您可以使用我上面提到的编译器的资源使用统计打印输出和/或可能通过检查生成的 PTX 代码（中间代码，类似于汇编代码）来进一步调查这一点：

nvcc -ptx ....

如果您选择检查 PTX，您需要参考PTX 手册。

于 2013-07-03T04:10:29.513 回答