调用__host__ __device__函数在__global__函数中导致开销
调用__host__ __device__函数在__global__函数中导致开销
提问于 2022-01-25 08:47:46
这是来自这条线的以下问题。我的__global__函数只包含一个Geoditic2ECEF(全球定位系统)。用一个API执行全局函数需要35 to。但是，如果我在__host__ __device__ Geoditic2ECEF(GPS gps)函数中编写整个__global__代码，而不是将其作为API调用，则__global__函数只需2 ms即可执行。在__host__ __device__函数中调用__global__ 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));
}回答 1
Stack Overflow用户
回答已采纳
发布于 2022-01-25 17:25:52
你似乎没有问一个我能看到的问题，所以我假设你的问题是“这个神秘的开销是什么，我有什么办法来减轻它？”
当对__device__函数的调用位于与该函数的定义不同的编译单元中时，编译器不能内联该函数(通常)。
这可能会对性能产生各种影响：
- 调用指令会造成一些开销。
- 函数调用具有保留寄存器的ABI，这会造成寄存器压力，从而影响代码性能。
- 编译器可能必须通过堆栈在寄存器之外传输附加的函数参数。这增加了额外的开销。
- 编译器不能(通常)在函数调用边界上进行优化。
所有这些都会在不同程度上产生性能影响，您可以在cuda标记上找到提到这些问题的其他问题。
我所知道的最常见的解决办法是：
- 将函数的定义移到与调用环境相同的编译单元(如果可能，从编译命令行中删除
-rdc=true或-dc)。 - 在最近的CUDA版本中，使用链路时间优化。
页面原文内容由Stack Overflow提供。腾讯云小微IT领域专用引擎提供翻译支持
原文链接：
https://stackoverflow.com/questions/70845729
复制相关文章
相似问题
腾讯云开发者
