在__global__函数内部调用__host____device__函数导致开销

Question

这是来自该线程的以下问题。我的__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标签上找到其他提到这些问题的问题。

据我所知，最常见的解决方案是：

1. 将函数的定义移动到与调用环境相同的编译单元(如果可能，从编译命令行中删除-rdc=true或-dc(
2. 在最近的CUDA版本中，使用链接时间优化