断点 break if, b if break [break-args] if (condition) break main if argc more ...

之前记录的一些cuda的用法中也有关于cuda汇编的一些介绍。本文主要记录了带源代码行信息的cuda汇编文件的获取。主要内容参考CUDA Binary Utilities 程序编译时需要添加的参数 在makefile或者cmakelist文件中，添加如下内容到nvccflag或者手动添加到nvcc编译的参 more ...

在Pacsal及更新的GPU中，managed memory在调用cudaMallocManaged()分配以后， 不一定在device memory上实际malloc。或者说，page和page table直到被GPU或CPU访问以后才被创建。page可以在任意时间迁移到任意memory，driver会采用启发算法来维护数据局部性和防止过多的page faults产生。 Reference https://developer.nvidia.com/blog/unified-memory-cuda-beginners/ more ...

之前的那篇转载浮点数的存储转载内容比较粗糙。今天又花了点时间试了个具体的例子。部分内容摘自参考。 以32bit浮点数 0.123456789的存储为例，通过在线进制转换获得其16进制表示为 3DFCD6E9 more ...

本文主要内容系转载。 标准 在 IEEE-754 规范[39]中，浮点数由三部分组成：符号位、指数部分和尾数部分（标准化表示方式）。单精度浮点数一般是用 4 字节（32bit）来表示。 不同标准的单精度浮点数、半精度浮点数表示方式[39] 数据类型 符号位 指数部分 尾数部分 more ...

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

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

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

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 ...