cuda - 我可以相信 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;

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