在 global 函数内调用 host device 函数导致开销

Question

这是来自 this thread 的以下问题。我的 __global__ 函数只包含一个 API Geoditic2ECEF(GPS gps)。使用单个 API 执行该全局函数需要 35 毫秒。但是，如果我在 __global__ 函数中编写 __host__ __device__ Geoditic2ECEF(GPS gps) 的整个代码，而不是将其作为 API 调用，则 __global__ 函数的执行时间仅为 2 毫秒。似乎在 __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__ 函数的调用与该函数的定义位于不同的编译单元中时，编译器无法该函数（通常）。

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

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

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

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

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

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

Calling __host__ __device__ function inside __global__ function causing an overhead

cuda

gpu

在 global 函数内调用 host device 函数导致开销

Calling host device function inside global function causing an overhead