PTX 是否真的有 64 位 warp shuffle 指令？答案

【问题标题】：Does PTX actually have a 64-bit warp shuffle instruction?PTX 是否真的有 64 位 warp shuffle 指令？
【发布时间】：2020-02-13 11:11:35
【问题描述】：

我注意到在__shfl_sync() 的文档中，支持 64 位数据类型（long、double）。

这是否意味着硬件/PTX 本身支持 64 位 warp shuffle，还是在编译代码时将它们分解为一对 32 位 shuffle？

【问题讨论】：

标签： cuda

【解决方案1】：

目前有no 64-bit shuffle instruction in PTX。当前所有 CUDA GPU 中的基本寄存器单元都是 32 位的。 64 位量没有对应的 64 位寄存器，而是占用一对 32 位寄存器。机器级别的 warp shuffle 操作在 32 位寄存器上运行。

编译器通过发出 2 个 PTX（或 SASS）指令来处理 CUDA C++ 的 shfl 内在函数的 64 位操作数。这很容易使用 CUDA binary utilities 发现/确认。

例子：

$ cat t45.cu
typedef double mt;
__global__ void k(mt *d){
        mt x = d[threadIdx.x];
        x = __shfl_sync(0xFFFFFFFF, x, threadIdx.x+1);
        d[threadIdx.x] = x;
}

$ nvcc -c t45.cu
$ cuobjdump -ptx t45.o

Fatbin elf code:
================
arch = sm_30
code version = [1,7]
producer = cuda
host = linux
compile_size = 64bit

Fatbin ptx code:
================
arch = sm_30
code version = [6,2]
producer = cuda
host = linux
compile_size = 64bit
compressed








.version 6.2
.target sm_30
.address_size 64



.visible .entry _Z1kPd(
.param .u64 _Z1kPd_param_0
)
{
.reg .pred %p<3>;
.reg .b32 %r<9>;
.reg .f64 %fd<3>;
.reg .b64 %rd<5>;


ld.param.u64 %rd1, [_Z1kPd_param_0];
cvta.to.global.u64 %rd2, %rd1;
mov.u32 %r5, %tid.x;
mul.wide.u32 %rd3, %r5, 8;
add.s64 %rd4, %rd2, %rd3;
ld.global.f64 %fd1, [%rd4];
add.s32 %r6, %r5, 1;

        mov.b64 {%r1,%r2}, %fd1;

        mov.u32 %r7, 31;
mov.u32 %r8, -1;
shfl.sync.idx.b32 %r4|%p1, %r2, %r6, %r7, %r8;
shfl.sync.idx.b32 %r3|%p2, %r1, %r6, %r7, %r8;

        mov.b64 %fd2, {%r3,%r4};

        st.global.f64 [%rd4], %fd2;
ret;
}


$

【讨论】：