Есть ли у PTX 64-битная инструкция тасования деформации?

Question

Я заметил, что в документах для __shfl_sync() и его родственников поддерживаются 64-битные типы данных (long, double).

Означает ли это, что аппаратное обеспечение / PTX изначально поддерживает 64-битные тасовки деформации, или они разбиты на пару 32-битных тасов при компиляции кода?

Answer 1

В настоящее время в PTX нет 64-битной инструкции тасования. Базовый регистр c во всех современных графических процессорах CUDA является 32-разрядным. 64-битные величины не имеют соответствующих 64-битных регистров, но вместо этого занимают пару 32-битных регистров. Операция деформации основы на машинном уровне работает с 32-разрядными регистрами.

Компилятор обрабатывает 64-разрядные операнды для встроенных функций shfl для CUDA C ++, испуская 2 команды PTX (или SASS). Это легко обнаружить / подтвердить, используя двоичные утилиты 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;
}


$