时间戳

首页 > 解决方案 > 我可以相信 NVCC 在返回类型中优化 std::pair 吗？

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

问题描述

有时，想要编写一个（小）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）来优化指针的构造，并直接分配给我稍后从该对的.firstand.second元素分配给的变量吗？

标签： cudastd-pairnvcccopy-elisionrvo

解决方案

我没有完整的答案，但从我有限的经验来看——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;

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

推荐阅读