分类: 学术

CUDA二进制探索

本文记录了探索NVIDIA CUDA SASS语法对应的二进制位的过程。 1. CUDA二进制文件 1.1 SASS NVCC编译过程和解读CUDA汇编PTX(二) SASS nvdisasm工具提过CUDA的汇编SASS,使用cuobjdump工具反编译出的SASS格式如下: /* 0x083fc400e3e007f6 */ /*0008*/ MOV R1, c[0x0][0x20]; /* 0x4c98078000870001 */ /*0010*/ S2R R2, SR_TID.X; /*...

GPU寄存器(二)

1. Introduction 本文介绍了NVIDIA GPU寄存器的相关内容。 2. GPU寄存器 2.1 物理寄存器的映射 关于gpu寄存器之前我还整理过:GPU寄存器 一个程序的近机器语言级别的中间语言中适用的寄存器,我们称之为“体系结构寄存器,architected register”,这些寄存器会被处理器映射到物理寄存器(Physical Registers)上。 CPU使用寄存器重命名(register renaming)和寄存器重命名表(Register Rename Table)...

Cache替换策略

1. Introduction 本文主要内容系Reference的整理,介绍了cache访问模式的分类和几种cache替换策略。 2. Cache访问模式分类 $$ \begin{align} &A:( a_1 , a_2 , … , a_{k-1} , a_k , a_k , a_{k-1} , … , a_2 , a_1 )^ N\text{ for any k}\\ &B:( a_1 , a_2 , … , a_k )^ N \text{...

半精度浮点数Half

1. Introduction 本文介绍了半精度浮点数的基本概念以及f32到f16转换的截断法。 混合精度逐渐成为提升深度学习速度的一种有效方法,其本质上,是以运算的精度换速度,当然前提是精度需要在可接受的范围内,或者说应用本身具有容错性(error tolerant)。 在cuda中,half2以及tensorcore的应用,就是对于精度损失容忍性的体现。 在线进制转换工具 2. 半精度浮点数 2.1 位宽 一个float单精度浮点数一般是4bytes(32bit)来表示,由三部分组成:符号位...

PTXPlus笔记

修改ptxplus建议 在gpgpusim源码src/cuda-sim/ptx_loader.cc的函数 char *gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptxfilename, const std::string elffilename,const std::string sassfilename)中,std::ifstream fileStream(fname_ptxplus, std::ios::i...

NVIDIA存储架构速度

GPU不同级别的存储架构有着不同的运行速度,下面是官方的ptx手册中的访问速度表格,比较粗略,以后想起来在哪篇论文里有更详细的再补充。 Space Time Notes Register 0   Shared 0   Constant 0 Amortized cost is low, first access is high Local > 100 clocks   Parameter 0   Immediate 0   Global &g...

使用GPGPU-SIM做实验

1. Introduction 本文说明了使用gpgpusim做实验需要注意的问题以及做出的修改。 强烈不推荐使用gpgpusim做实验。 如果是改SASS,建议maxas或者asfermi(尽管他们可能会有这样那样的问题,而且使用起来难度不小)。使用模拟器,是被reviewer攻击的常用点,除非你有足够强有力的解释,否则这会成为你论文的掣肘。 使用模拟器发的A,一般工作量会很多。 2. 环境搭建 2.1 编译运行环境 docker里跑gpgpusim gpgpusim runtime的dock...

cuda寄存器限制launch_bound和maxrregcount

一个CUDA程序如果使用的寄存器数量过多,会导致在SM上同时驻留的线程和block数量减少,继而导致程序性能不足。 __launch_bounds__和maxrregcount都可以用来限制cuda程序的寄存器数量,但是两者是不同的机制。 __launch_bounds__ __global__ void __launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor) MyKernel(…) { … } maxThrea...

paper阅读01-GPU Scheduling on the NVIDIA TX2: Hidden Details Revealed

以前读完论文并没有写keynotes的习惯,顶多在mendeley上做些标注,但是过段时间再看,还是需要花费一定的时间来理顺思路。所以准备从现在开始对阅读过的论文做一些自己认为重要的备注。 GPU Scheduling on the NVIDIA TX2: Hidden Details Revealed 2017 IEEE Real-Time Systems Symposium (RTSS) (2017) Paris, France Dec 5, 2017 to Dec 8, 2017 ISSN...

gpgpusim runtime的dockerfile

Introduction 之前写过一篇在docker里跑gpgpusim(下称参考文献1),今天花了一点时间,写了个完整的dockerfile,并上传到了docker hub上去。现在简化一下运行的步骤。 gpgpusim-runtime 在docker hub的地址: https://hub.docker.com/r/findhao/gpgpusim_runtime/ 运行 执行参考文献1的第4节 拷贝必须的文件到自己的目录,比如/home/find/e/gpgpusim/。 执行: # 可从...