GPU寄存器(二)

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

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{ k > cache size}\\&C:(( a_ more ...

半精度浮点数Half

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



使用GPGPU-SIM做实验

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

cuda寄存器限制launch_bound和maxrregcount

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



解读CUDA汇编PTX(二) SASS nvdisasm工具

Introduction解读CUDA汇编PTX--目录NVIDIA CUDA的NVCC编译过程之前已经介绍过了,编译ptx后,会生成cubin文件。cubin文件是包含了CUDA执行代码节的ELF格式文件。类似于我们常见运行文件。而官方提供了两个工具来反编译cubin文件到sass文件(类似常见的汇编), more ...