我可以信任NVCC来优化返回类型中的std :: pair吗？

Question

有时，人们想编写一个（小）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元素分配给的变量？

Answer 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;

没有额外的副本或与施工有关的操作。