cuda-7.5 未定义标识符“__shfl_down”答案

【问题标题】：identifier "__shfl_down" is undefined for cuda-7.5cuda-7.5 未定义标识符“__shfl_down”
【发布时间】：2019-11-24 14:09:53
【问题描述】：

在 ubuntu 14.04（相当旧的配置）上使用 gcc 4.8.4 在 cuda 7.5 上编译程序时，我收到此错误

error: identifier "__shfl_down" is undefined
      detected during instantiation of "T gmx_shfl_down_sync(unsigned int, T, unsigned int, int) [with T=float]"

哪个指向

template <typename T>
static __forceinline__ __device__
T gmx_shfl_down_sync(const unsigned int activeMask,
                 const T            var,
                 unsigned int       offset,
                 int                width = warp_size)
{
#if GMX_CUDA_VERSION < 9000
    GMX_UNUSED_VALUE(activeMask);
    return __shfl_down(var, offset, width);
#else
    return __shfl_down_sync(activeMask, var, offset, width);
#endif
}

有没有办法解决这个问题？我看到旧 cuda 版本的这个问题，但还没有看到明确的答案。

【问题讨论】：

标签： gcc cuda

【解决方案1】：

Warp shuffle 内在函数仅在计算能力 (cc) 3.0 及更高版本中定义（仅支持）。

在 CUDA 8.0 之后，这些是 nvcc 支持的唯一 GPU，因此即使您针对默认架构 (3.0) 进行编译，它也会正确编译。

但是对于CUDA 8.0 and prior，nvcc 仍然支持 cc 2.x 架构，并且仍然是“默认”架构（如果您没有在 nvcc 编译命令上指定任何架构开关，您会得到什么线）。

因此，在 CUDA 8.0 和更早版本（CUDA 6 中引入了 warp shuffle）上，如果您指定 cc 2.x 架构或不指定架构，您将看到此错误。

由于 cc 2.x 架构不支持 warp shuffle，因此解决方案是在 nvcc 编译命令行中使用 -arch=sm_30 或类似名称明确指定合适的架构。

【讨论】：

我使用的是 SM_20 的 C2075。所以这意味着，我也必须恢复到旧的源代码。
正确。您不能在该 GPU 上使用任何 warp shuffle 代码。