两个 SASS 分分析案例

1. shfl_sync的机器 sass 汇编代码

1.1 实验目标

对比

int ret = __shfl_sync(0xFFFFFFFF, value, 5, 16);

int ret = __shfl_sync(0xFFFFFFFF, value, 5, 32);

不同的 sass 汇编代码

1.2 实验代码

源代码 shfl 16：

shft_sync_test_16.cu

#include <iostream>
#include <stdio.h>

__global__ static void shfffl_test(int *A)
{
    int tid = threadIdx.x;
    int value = tid;

    int ret = __shfl_sync(0xFFFFFFFF, value, 5, 16);

    A[tid] = ret;
}

int main()
{
    int *A = nullptr;
    int *A_h = nullptr;

    cudaMalloc((void**)&A, 32*sizeof(int));
    A_h = (int*)malloc(32*sizeof(int));
    for(int i=0; i<32; i++)
	A_h[i] = i*3;
    cudaMemcpy(A, A_h, 32*sizeof(int), cudaMemcpyHostToDevice);

    shfffl_test<<<1, 32>>>(A);

    cudaDeviceSynchronize();

    cudaMemcpy(A_h, A, 32*sizeof(int), cudaMemcpyDeviceToHost);

    for(int i=0; i<32; i++)
	std::cout<<A_h[i]<<"  "<<std::endl;

    return 0;
}

编译运行：

$ nvcc shft_sync_test_16.cu -o ./shfl_sync_test_16.out

$ ./shfl_sync_test_16.out

源代码 shfl 32：

shfl_sync_test_32.cu

#include <iostream>
#include <stdio.h>

__global__ static void shfffl_test(int *A)
{
    int tid = threadIdx.x;
    int value = tid;

    int ret = __shfl_sync(0xFFFFFFFF, value, 5, 32);

    A[tid] = ret;
}

int main()
{
    int *A = nullptr;
    int *A_h = nullptr;

    cudaMalloc((void**)&A, 32*sizeof(int));
    A_h = (int*)malloc(32*sizeof(int));
    for(int i=0; i<32; i++)
	A_h[i] = i*3;
    cudaMemcpy(A, A_h, 32*sizeof(int), cudaMemcpyHostToDevice);

    shfffl_test<<<1, 32>>>(A);

    cudaDeviceSynchronize();

    cudaMemcpy(A_h, A, 32*sizeof(int), cudaMemcpyDeviceToHost);

    for(int i=0; i<32; i++)
	std::cout<<A_h[i]<<"  ";

    std::cout<<std::endl;

    return 0;
}

编译运行：

$ nvcc shft_sync_test_32.cu -o ./shfl_sync_test_32.out

$ ./shfl_sync_test_32.out

分别执行 cuobjdump -sass xxx.cubin，可以查看器中的机器汇编 sass：

1.3 实验结论

先说结论：

猜测，16个线程一组,warp 分为两组做shfl_sync，与warp 内 32 个线程一大组做 shfl_sync，都是使用一条指令完成；

而不是两条。（为什么可能是两条呢？）如果2*16个线程之间，每16一组，硬件无法跨组传递数据，那么，这里的两种情况，其代码会不一样。

int ret = __shfl_sync(0xFFFFFFFF, value, 5, 16);

int ret = __shfl_sync(0xFFFFFFFF, value, 5, 32);

汇编的不同之处仅仅是最后的参数立即数不同。

32个 warp 线程 shfl_sync 时，立即数参数为 0x1f;

16 个 warp 线程 shfl_sync 时，立即数参数为 0x101f；

但是汇编机器码稍有差别：

/* 0xef17407c30570400 */

/* 0xef17007c30570400 */

其中的数字5，是指warp lane-id == 5 的线程做 shfl 广播；

这里例子是做广播，相对简单。

接下来测试一下 shfl_sync_down 的响应情况；

上例的Makefile：

OUT := 	shfl_sync_test_08.cubin \
	shfl_sync_test_08.out 	\
	shfl_sync_test_16.cubin \
	shfl_sync_test_16.out 	\
	shfl_sync_test_32.cubin \
	shfl_sync_test_32.out

all: $(OUT)

%.out: %.cu
	nvcc $< -o $@


%.cubin: %.cu
	nvcc -cubin $< -o $@

.PHONY: clean
clean:
	-rm -rf $(OUT)