返回对 CUDA 特定向量类型的引用

Question

我想使用参数和 return 值来实现 f1，与下面的代码完全一样。它因错误而失败：

a reference of type "float1 &" (not const-qualified) cannot be initialized with a value of type "float"

然而，几乎相同的函数 f2 使用原生 C++ 类型 float 而不是特定于 CUDA 的包装器 float1 工作得很好。

__device__ float1& f1(float4& v) {
    return v.x; // ERROR
}

__device__ float& f2(float4& v) {
    return v.x; // OK
}

如何更改 f1 的实现来修复此错误？

Answer 1

这个解决方案在评论中讨论过，我也说我觉得这有点脏，但是如果 CUDA 规范保证 float4 和 float1 值的对齐，那么这可能是一个有效的选项；

__device__ float1& f1(float4& v) {
    return *reinterpret_cast<float1*>(&v);
}

__device__ float& f2(float4& v) {
    return v.x;
}

在此解决方案中，您将 v 的地址重新解释为指向 float1 的指针。然后，您可以取消引用结果，使 v 成为 float1&.

在对齐和偏移方面，请注意 reinterpret_cast 和不同的 struct。

在 Compiler Explorer Example 中，您可以看到两个函数产生完全相同的输出。

Returning references to CUDA-specific vector types