Могу ли я доверять NV CC для оптимизации std :: pair в возвращаемых типах?

Question

Иногда хочется написать (маленькую) функцию на стороне устройства CUDA, которая возвращает два значения. В C эта функция будет принимать два выходных параметра, например:

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

, но в C ++ идиоматический способ c написать это - вернуть std::pair (хорошо, может быть std::tuple или структура, но кортежи C ++ неуклюжи и структура не достаточно обобщенная c достаточно:

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

Мой вопрос: могу ли я доверять NV CC (с --expt-relaxed-constexpr) чтобы оптимизировать конструкцию указателя и просто назначить непосредственно переменным, которые я позже назначу из .first и .second элементов пары?

Answer 1

У меня нет полного ответа, но из моего ограниченного опыта - кажется, что NV CC может оптимизировать 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;

Без дополнительных копий или строительные работы.