在 __global__ 函数中调用 __host__ __device__ 函数会导致开销

Question

这是来自this thread 的以下问题。 我的__global__ 函数只包含一个 API Geoditic2ECEF(GPS gps)。使用单个 API 执行该全局函数需要 35 毫秒。但是，如果我在__global__ 函数中编写__host__ __device__ Geoditic2ECEF(GPS gps) 的整个代码而不是将其作为API 调用，则__global__ 函数只需2 ms 即可执行。似乎在__global__ 函数中调用__host__ __device__ API 会导致神秘的开销。

这是我使用 API 时的 PTX 输出

ptxas info    : Compiling entry function '_Z16cudaCalcDistanceP7RayInfoPK4GPS3PK6float6PK9ObjStatusPKfSB_SB_fiiiiii' for 'sm_52'
ptxas info    : Function properties for _Z16cudaCalcDistanceP7RayInfoPK4GPS3PK6float6PK9ObjStatusPKfSB_SB_fiiiiii     0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info    : Used 9 registers, 404 bytes cmem[0]

这是我不使用 API 时的 PTX 输出

ptxas info    : Compiling entry function '_Z16cudaCalcDistanceP7RayInfoPK4GPS3PK6float6PK9ObjStatusPKfSB_SB_fiiiiii' for 'sm_52' 
ptxas info    : Function properties for _Z16cudaCalcDistanceP7RayInfoPK4GPS3PK6float6PK9ObjStatusPKfSB_SB_fiiiiii     0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info    : Used 2 registers, 404 bytes cmem[0]

唯一的区别是 API 版本使用了 9 个寄存器，而非 API 版本使用了 2 个寄存器。我可以从这些信息中推断出什么。

在文件 utils.cu 中，我定义了以下结构和 API

struct GPS {
    float latitude;
    float longtitude;
    float height;
};

struct Coordinate
{
    __host__ __device__ Coordinate(float x_ = 0, float y_ = 0, float z_= 0)
    {
        x = x_;
        y = y_;
        z = z_;
    }
    __host__ __device__ float norm()
    {
        return sqrtf(x * x + y * y + z * z);
    }

    float x;
    float y;
    float z;
};

__host__ __device__ Coordinate Geoditic2ECEF(GPS gps)
{
    Coordinate result;

    float a = 6378137;
    float b = 6356752;
    float f = (a - b) / a;
    float e_sq = f * (2 - f);

    float lambda = gps.latitude / 180 * M_PI;
    float phi = gps.longtitude / 180 * M_PI;

    float N = a / sqrtf(1 - e_sq * sinf(lambda) * sinf(lambda));
    result.x = (gps.height + N) * cosf(lambda) * cosf(phi);
    result.y = (gps.height + N) * cosf(lambda) * sinf(phi);
    result.z = (gps.height + (1 - e_sq) * N) * sinf(lambda);
    return result;
}

在main.cu，我有以下功能

__global__ void cudaCalcDistance(GPS* missile_cur,
                                 int num_faces, int num_partialPix)
{
    int partialPixIdx = threadIdx.x + IMUL(blockIdx.x, blockDim.x);
    int faceIdx = threadIdx.y + IMUL(blockIdx.y, blockDim.y);

    if(faceIdx < num_faces && partialPixIdx < num_partialPix)
    {
        Coordinate missile_pos;
        // API version
        missile_pos = Geoditic2ECEF(missile_cur->gps);
        // non_API version
//        float a = 6378137;
//        float b = 6356752;
//        float f = (a - b) / a;
//        float e_sq = f * (2 - f);

//        float lambda = missile_cur->latitude / 180 * M_PI;
//        float phi = missile_cur->longtitude / 180 * M_PI;

//        float N = a / sqrtf(1 - e_sq * sinf(lambda) * sinf(lambda));
//        missile_pos.x = (missile_cur->height + N) * cosf(lambda) * cosf(phi);
//        missile_pos.y = (missile_cur->height + N) * cosf(lambda) * sinf(phi);
//        missile_pos.z = (missile_cur->height + (1 - e_sq) * N) * sinf(lambda);

    }
}

void calcDistance(GPS * data)
{
    int num_partialPix = 10000;
    int num_surfaces = 4000;


    dim3 blockDim(16, 16);
    dim3 gridDim(ceil((float)num_partialPix / threadsPerBlock),
                 ceil((float)num_surfaces / threadsPerBlock));

    cudaCalcDistance<<<gridDim, blockDim>>>(data,                                 
                             m_Rb2c_cur,num_surfaces,num_partialPix);
    gpuErrChk(cudaDeviceSynchronize());
}

int main()
{
    GPS data(11, 120, 32);
    GPS *d_data;
    gpuErrChk(cudaMallocManaged((void**)&d_data, sizeof(GPS)));
    gpuErrChk(cudaMemcpy(d_data, &data, sizeof(GPS), cudaMemcpyHostToDevice));
    calcDistance(d_data);
    gpuErrChk(cudaFree(d_data));
}

Answer 1

您似乎没有问过我可以看到的问题，所以我假设您的问题类似于“这个神秘的开销是什么，我有什么办法可以减轻它？”

当对__device__ 函数的调用与该函数的定义位于不同的编译单元中时，编译器不能inline 该函数（通常）。

这会对性能产生多种影响：

• 调用指令会产生一些开销
• 函数调用具有保留寄存器的 ABI，这会产生可能影响代码性能的寄存器压力
• 编译器可能必须通过堆栈将其他函数参数传输到寄存器之外。这会增加额外的开销。
• 编译器无法（通常）跨函数调用边界进行优化。

所有这些都会在不同程度上对性能产生影响，您可以在提及这些的cuda 标签上找到其他问题。

我知道的最常见的解决方案是：

1. 将函数的定义移动到与调用环境相同的编译单元（如果可能，从编译命令行中删除-rdc=true 或-dc）。
2. 在最近的 CUDA 版本中，使用link-time optimization。