两个 SASS 分分析案例

本文主要是介绍两个 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.outall: $(OUT)%.out: %.cunvcc $< -o $@%.cubin: %.cunvcc -cubin $< -o $@.PHONY: clean
clean:-rm -rf $(OUT)

2， shfl_sync_up的 机器 sass 汇编代码

这篇关于两个 SASS 分分析案例的文章就介绍到这儿，希望我们推荐的文章对编程师们有所帮助！

---