cuda - 用于 Kepler 架构的加载/存储单元 (LD/ST) 和特殊功能单元 (SFU)

标签 cuda nvidia kepler

Kepler architecture whitepaper , NVIDIA 声明有 32特殊功能单元 (SFU) 和 32在 SMX 上加载/存储单元 (LD/ST)。

SFU 用于“快速近似超越运算”。不幸的是,我不明白这是什么意思。另一方面,在 Special CUDA Double Precision trig functions for SFU据说,它们只能在单精度下工作。这在 K20Xm 上仍然正确吗?

LD/ST 单元显然用于存储和装载。通过这些论文之一是否需要任何内存加载/写入?它们是否也用作单一经线?换句话说,是否可以只有一个当前正在写入或读取的经线?

干杯,
和我

最佳答案

The SFU are for "fast approximate transcendental operations"



SFU 计算函数,如 __cosf() , __expf()等等。

On the other hand here is said, that they only work in single precision, is this still correct on a K20Xm?



据近期CUDA C Programming Guide, section G.5.1它们仍然只能在单精度下工作。

这是有道理的,因为如果您需要 double ,则不太可能使用不准确的数学函数。您可以引用this answer有关 double 算术优化的建议。

double 运算的实现细节可以在/usr/local/cuda-5.5/include/math_functions_dbl_ptx3.h中找到。 (或安装 CUDA Toolkit 的任何地方)。
例如。为 sincos它使用 Payne-Hanek 参数缩减,然后是 Taylor 展开(最多 14 阶)。

对于 double 计算,SFU 似乎仅用于 __internal_fast_rcp__internal_fast_rsqrt ,依次用于 acos , log , cosh和其他几个函数(见 math_functions_dbl_ptx3.h)。所以大多数时候它们会停顿,比如如果没有正在进行的内存事务,LD/ST 单元就会停顿。

Is any memoryload/write required to go through one of theses?



是的,每次访问全局内存。

And are they also used as a single warp? In other words can there be only one warp which is currently writing or reading?



单元数仅限制每个周期发出的指令数。 IE。每个时钟周期可以发出32条读指令,可以返回32条结果。

一条指令最多可以读/写 128 个字节,因此如果 warp 中的每个线程读取 4 个字节并且它们被合并,那么整个 warp 将需要单个加载/存储指令。如果访问未合并,则应发出更多指令。

而且,单位是pipelined ,这意味着单个单元可以同时执行多个读取/存储请求。

关于cuda - 用于 Kepler 架构的加载/存储单元 (LD/ST) 和特殊功能单元 (SFU),我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/20472670/

相关文章:

python - Eclipse CDT 中的 pretty-print : unable to look at any variable

cuda - __host__ __device__ 函数可以知道它在哪里执行吗?

cuda - C 中的统一内存和流

concurrency - 同一多处理器上的并发、唯一内核?

c++ - OpenCV2.4错误: No GPU support in unknown function file

c++ - 为什么我的程序在 nVidia NView 下占用 100% CPU?

CUDA 固定内存从设备刷新

ffmpeg - 如何将FFmpeg的AVPacket转入CUVID的CUVIDSOURCEDATAPACKET?或者如何使用 FFmpeg 的 CUVID,任何演示?

python - ubuntu 12.04 上的 CUDA 安装