CUDA 10 Memory Transaction的一个现象

1. Introduction近日,在写一些microbenchmark分析cuda程序访存问题时,发现了一个有趣的问题。目前尚未找到合理的解释,先记录下来以待后续分析。实验平台为:NVIDIA GTX950,sm5.0,maxwell架构。2. Global MemoryA memory "request" is an instruction which accesses memory, and a "transaction" is the movement of a unit of da more ...

CUDA Sanitizer Samples使用

1. IntroductionCUDA 10.1推出了新的API:The Compute Sanitizer API,提供了更底层更丰富的Instrumentation API。https://docs.nvidia.com/cuda/sanitizer-docs/SanitizerApiGuide/index.html目前相关文档还比较简单,本文记录下官方Samp more ...

Deepin 15.10 安装cuda toolkit 10.1

1. Introductiondeepin可以按照正常cuda toolkit的方式安装cuda 9.0,但是10.0+就出现了问题。查看安装日志也看不出所以然。在cuda论坛阴差阳错发现了一个用来解决其他问题的方法,但是可以用来解决deepin上cuda toolkit的安装。2. 正常安装toolkit和driver的方法2.1 禁用默认闭源驱动# 使用vim或者其他编辑器添加配置文件 more ...

CUDA二进制探索

本文记录了探索NVIDIA CUDA SASS语法对应的二进制位的过程。1. CUDA二进制文件1.1 SASSNVCC编译过程和解读CUDA汇编PTX(二) SASS nvdisasm工具提过CUDA的汇编SASS,使用cuobjdump工具反编译出的SASS格式如下: more ...

GPU寄存器(二)

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