不能在 Cuda 内核中使用 __m128i

Posted 2023-04-15

【中文标题】不能在 Cuda 内核中使用 __m128i

【英文标题】：Can't use __m128i in Cuda kernel

【发布时间】：2020-12-09 16:54:30

【问题描述】：

我正在尝试使用 cuda 编译一个使用 __m128i 的简单程序，但是当我在 Linux 上使用 nvcc (nvcc test.cu -o test) 进行编译时，我得到了 "__m128i" is a vector, which is not supported in device code。 这是我要编译的程序

#include <stdio.h>
#include <emmintrin.h>

__global__ void hello()
    printf("%d\n",threadIdx.x);
    __m128i x;


int main()
   hello<<<3,3>>>();

当我输入nvcc --version 时，我得到Cuda compilation tools, release 10.2, V10.2.89

我实际上在更大范围内遇到了这个问题，试图使用 CUDA 实现一些 cpp 代码，而这个 cpp 代码使用__m128i，我所展示的是我所面临问题的简单版本，所以我想知道是否有一种方法可以在 CUDA 内核中使用 __m128i 或其他替代方法。谢谢

【问题讨论】：

@TedLyngmo：那篇文章谈到了 GNU C __uint128_t，它与 __m128i 完全无关，只是大小相同。 SSE 整数向量不是 128 位整数类型；最宽的元素大小是_mm_add_epi64。 （除非您只使用按位布尔运算，否则元素边界无关紧要。）

@TedLyngmo：但__m128i 不是a 128 位整数；这是一个 SIMD 向量。在 GNU C 中，定义为 typedef long long __m128i __attribute__((vector_size(16), may_alias))。拥有受 CUDA 支持的标量 128 位整数类型不会帮助您编译使用 __m128i 和 _mm_shuffle_epi32、_mm_add_epi32 等内在函数的代码（将其视为 4x 32 位整数的向量） ，或_mm_minpos_epu16（16 位无符号元素的水平最小和最小位置），或其他 SSE 硬件操作。您不能将 __m128i 用作单个 128 位整数，因此这不是 OP 想要的。

@PeterCordes 啊......现在我明白你在说什么了。 ：） 对困惑感到抱歉。删除评论。

【参考方案1】：

我想知道是否有办法在 CUDA 内核中使用 __m128i ...

没有。 CUDA 有native 128 bit integer types，它满足与__m128i 相同的对齐属性，但不支持宿主向量类型。

或其他替代方法

如上所述，有 16 字节对齐类型可用于加载和存储数据，但 NVIDIA GPU 中没有原生 128 位 SIMD 内在支持。存在的SIMD instructions 仅限于32位类型。

---

CPU SIMD 使用短向量（如 128 位 __m128i）完成。 GPU SIMD 是跨线程完成的，并且通常不像__m128i CPU SIMD 那样对软件可见，您只需将其编写为标量代码即可。

无法为 GPU 编译使用 __m128i 手动矢量化的代码。如果它有一个标量后备版本，请使用它，例如#undef __SSE2__.

（如果您的数据很窄，如 16 位整数对或 4 个 8 位整数，则 32 位块内的 CUDA SIMD 可让您更多地利用每个 GPU 执行单元中的 32 位宽 ALU。因此，如果您的 SSE 内在代码使用 _mm_add_epi8，您可能仍会受益于 CUDA 中的手动矢量化，其 4x 8 位操作而不是 16x 8 位。）

【讨论】：

我知道你知道，我应该对你的社区 wiki 答案进行编辑，因为它是为了让未来的读者受益，他们不明白 CPU 和 GPU 以不同的方式将 SIMD 暴露给程序员。更新：已编辑。