我可以相信 NVCC 在返回类型中优化 std::pair 吗？答案

【问题标题】：Can I trust NVCC to optimize away std::pair in return types?我可以相信 NVCC 在返回类型中优化 std::pair 吗？
【发布时间】：2020-05-26 02:02:39
【问题描述】：

有时，人们想编写一个（小）CUDA 设备端函数，它返回两个值。在 C 语言中，您可以让该函数采用两个输出参数，例如：

__device__ void pair_maker(float x, float &out1, float& out2);

但在 C++ 中，编写此代码的惯用方式是返回 std::pair（嗯，可能是 std::tuple 或结构，但 C++ 元组很笨重，结构不够通用）：

__device__ std::pair<float, float> pair_maker(float x);

我的问题：我是否可以信任 NVCC（使用 --expt-relaxed-constexpr）来优化指针的构造，并直接分配给我稍后从 .first 和 .second 元素分配给的变量对吗？

【问题讨论】：

标签： cuda std-pair nvcc copy-elision rvo

【解决方案1】：

我没有完整的答案，但根据我有限的经验 - NVCC 似乎可以优化 std::pair 了。插图（也在GodBolt）：

#include <utility>

 __device__ std::pair<float, float> pair_maker(float x) {
    float  sin, cos;
    __sincosf(x, &sin, &cos);
    return {sin, cos};
}

__device__ float foo(float x) {
    auto p = pair_maker(x);
    auto sin = p.first;
    auto cos = p.second;
    return sin + cos;
}

__global__ void bar(float x, float *out) { *out = foo(x); }

__global__ void baz(float x, float *out) {
    float sin, cos;
    __sincosf(x, &sin, &cos);
    *out = sin + cos;
}

内核 bar() 和 baz() 编译成相同的 PTX 代码：

ld.param.f32    %f1, [param_0];
ld.param.u64    %rd1, [param_1];
cvta.to.global.u64      %rd2, %rd1;
sin.approx.f32  %f2, %f1;
cos.approx.f32  %f3, %f1;
add.f32         %f4, %f2, %f3;
st.global.f32   [%rd2], %f4;
ret;

没有额外的副本或与构造相关的操作。

【讨论】：