时间戳

首页 > 解决方案 > PTX 是否真的有 64 位 warp shuffle 指令？

cuda - PTX 是否真的有 64 位 warp shuffle 指令？

问题描述

我注意到在 docs for __shfl_sync()and it's relatives 中，支持 64 位数据类型（long、double）。

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

标签： cuda

解决方案

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

shfl编译器通过发出 2 个 PTX（或 SASS）指令来处理 CUDA C++ 内部函数的64 位操作数。使用 CUDA二进制实用程序很容易发现/确认这一点。

例子：

$ 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;
}


$

推荐阅读