Bruce_0712的博客

作者：谭旭链接：https://www.zhihu.com/question/63219175/answer/206697974来源：知乎著作权归作者所有。商业转载请联系作者获得授权，非商业转载请注明出处。NCCL是Nvidia Collective multi-GPU Communication Library的简称，它是一个实现多GPU的collective communication通信（all-gather, reduce, broadcast）库，Nvidia做了很多优化，以在PCI.

dockr虽然安装好了，但是安装到这一步的docker只能在cpu下使用或者将docker的gpu环境导出到宿主机上使用，无法在docker中使用gpu。nvidia-docker是一个可以使用GPU的docker，nvidia-docker是在docker上做了一层封装，需要先安装好docker。官网：https://github.com/NVIDIA/nvidia-docker1、跟新nvidia的版本到19.03由于我的nvidia-docker2安装失败了，所以参考官方强调的使用nv

http://itindex.net/detail/59715-nvidia-gpu-linuxhttps://people.centos.org/arrfab/shim/results/kernel/20181108233701/3.10.0-957.el7.x86_64/

使用更少线程意味着每个线程拥有更多的寄存器。　　　　每个线程的寄存器数：　　GF100：在100%占用率时有20个，在33%占用率时63个，为3倍。　　GT200：在100%占用率时有16个，在12.5%占用率时约有128个，为8倍　　那么每个线程有更多的寄存器是不是更好呢?　　只有寄存器的速度才能足够达到峰值。考虑这样一个计算: a*b+c：2个f

问题很简单，当我们在编写KERNEL的时候，分配了一个数组，那么这段数组空间是在register中，还是local memory中呢？通过几个测试，我们可以来看一下：首先一些定义：[cpp] view plain copy#define BLOCK_SIZE 32  #define GRID_SIZE 1  #define A

CUDA程序优化CUDA程序优化应该考虑的点：精度：只在关键步骤使用双精度，其他部分仍然使用单精度浮点以获得指令吞吐量和精度的平衡；延迟：需要首先缓冲一部分数据，缓冲的大小应该可以保证每个内核程序处理的一批数据能够让GPU慢负荷工作；计算量：计算量太小的程序使用CUDA很不合算；当需要计算的问题的计算密集度很低的时候，执行计算的时间远远比IO花费的时间短，整个程序的瓶颈出现在PCI

对指针和地址比较熟悉的童鞋可以比较容易的理解这个概念，为了更好的让更多的人理解cudaMalloc的参数，写这篇文章，以飨读者。首先看下此运行时函数的原型：[cpp] view plain copy cudaError_t cudaMalloc (void **devPtr, size_t  size );   此函数返回值是CUDA中定义的一个错误

统一寻址（Unified Memory）：可直接访问CPU内存、GPU显存，不需要手动拷贝数据。CUDA 6在现有的内存池结构上增加了一个统一内存系统，程序员可以直接访问任何内存/显存资源，或者在合法的内存空间内寻址，而不用管涉及到的到底是内存还是显存。CUDA 6的数据拷贝由程序员的手动转移，变成自动执行，因此，它仍然受制于PCI-E的带宽和延迟。NVIDIA的统一内存寻

1 定位bug出现bug的第一想法自然是定位bug。cuda比较奇特的地方在于，有时报错bug在500行，但500行出的代码没有错误，而是在1000行的地方逻辑错了，十分头疼。下面介绍三种我总结的定位bug方法：1.1 二分法一半一半的注释代码，定位bug。比较笨拙和麻烦，但是十分好用。1.2 输出定位法将整体代码分为几个模块，正常的CUDA代码大概可以分为数据初始化，

在CUDA中统计运算时间，大致有三种方法： 使用cutil.h中的函数unsigned int timer=0;//创建计时器cutCreateTimer(&timer);//开始计时cutStartTimer(timer);{  //统计的代码段  …………}//停止计时cutStopTimer(timer);//获得从开始计时到停

在CUDA2.2以下，仅提供cudaMallocHost函数用于分配页锁定内存，与C语言函数malloc分配分页内存相对应。而从CUDA2.2开始，页锁定内存增加三种新的类型用于主机多线程的portable，用于高效写回write-combined以及零拷贝的mapping，用cudaHostAlloc进行分配，其中采用四个可选参数标志用以指定使用何种特性：cudaHostAllocDef

Tesla图形与计算架构：流处理器阵列（scalable streaming processor array，SPA）+存储器系统，由片上互联网络连接；存储器系统：存储器控制器（MMC）、固定功能的光栅操作单元（raster operation processors，ROP），二级纹理缓存；            MMC：负责控制片外的DRAM显存，每个存储器可以提供64bit位宽；

PCI-E控制器，即PCI-EXPRESS LANES控制器，可以支持显示卡。PCI Express接口模式 通常用于显卡网卡等,主板类接口卡.满足条件:主板必须有PCI Express专用插槽。优势与性能介绍:-与PCI和AGP插槽相比，PCI-Express更具有潜在的生产价值。-比PCI总线具有更高的可测量性。能够满足硬盘控制器，千兆网卡以及其他一些对带宽需求较大

cuda没有全局线程同步的函数,__syncthreads()只用于block内线程的同步,调用cudaThreadSynchronize()函数,会使cpu处于等待状态,等待所有的线程都执行完毕.但是是,cudaThreadSynchronize()函数并不能在kernel中使用。

https://devtalk.nvidia.com/default/topic/463316/branch-divergent-branches/

有很多同鞋问怎么使用CUDA和其它的编译器连用呢？混合编程？先吧代码贴出来：文件1 : test1.cu[cpp] view plain copy //文件：test1.cu  #include   #include   #include     #define ROWS 32  #define COLS 16  #define CHECK(res) i

What is GPU Boost?GPU Boost is a new user controllable feature to change theprocessor clock speed on the Tesla K40 GPU. NVIDIA is currentlysupporting 4 selectable Stream Processor clock speeds and t

简单可以理解为：cudaMemcpy是同步的，而cudaMemcpyAsync是异步的。具体理解需要弄清以下概念：1.CUDA Streams在cuda中一个Stream是由主机代码发布的一系列再设备上执行的操作，必须确保顺序执行。不同streams里面的操作可以交叉执行或者并发执行。2.默认stream设备操作包括：数据传输和kernels，在cuda中，所有的设备操作都在str

背景qgzang@ustc:~$ nvidia-smi -h1输出如下信息：NVIDIA System Management Interface – v352.79NVSMI provides monitoring information for Tesla and select Quadro devices. The data is presented in

问：L1 cache只能用来缓存global memory，L2 cache即能用来缓存local也能用来缓存global memory是吧？如果是这样的话，我在进行global memory访问时，假设L1 cache，L2 cache都默认开启，我如何知道访问的数据是被缓存在L1中还是L2中，或者两者都有？答：所有的计算能力（不算1.x)， L1都能缓冲local memory的，而不

Catalog已经达到极限了？影响代码性能的两个主要方面优化代码内存操作看代码内存操作是否有效——DRAM utilizationcoalesce合并从little’s Law中找继续优化的方法SM中的occupancy占用率优化代码计算性能减小线程发散度选择效率更高的数学计算已经达到极限了？经过了对行、各元素并行化处理，OK，似乎目前得到了一个还不

之前我们有介绍shared Memory对于提高性能的好处，在CC3.0以上，支持了shuffle指令，允许thread直接读其他thread的寄存器值，只要两个thread在 同一个warp中，这种比通过shared Memory进行thread间的通讯效果更好，latency更低，同时也不消耗额外的内存资源来执行数据交换。这里介绍warp中的一个概念lane，一个lane就是一个warp中

到目前为止，所有kernel都是在host端调用，GPU的工作完全在CPU的控制下。CUDA Dynamic Parallelism允许GPU kernel在device端创建调用。Dynamic Parallelism使递归更容易实现和理解，由于启动的配置可以由device上的thread在运行时决定，这也减少了host和device之间传递数据和执行控制。我们接下来会分析理解使用Dynamic

CUDA 5.0中引入动态并行化，使得在device端执行的kernel的线程也能跟在host上一样launch kernels，只有支持CC3.5或者以上的设备中才能支持。动态并行化使用CUDA Device Runtime library（cudadevrt），它是一个能在device code中调用的CUDA runtime子集。编译链接为了支持动态并行化，必须使用两步分离编

CUDA的原子操作可以理解为对一个变量进行“读取-修改-写入”这三个操作的一个最小单位的执行过程，这个执行过程不能够再分解为更小的部分，在它执行过程中，不允许其他并行线程对该变量进行读取和写入的操作。基于这个机制，原子操作实现了对在多个线程间共享的变量的互斥保护，确保任何一次对变量的操作的结果的正确性。原子操作确保了在多个并行线程间共享的内存的读写保护，每次只能有一个线程对该变

从硬件角度分析，支持CUDA的NVIDIA 显卡，都是由多个multiprocessors 组成。每个 multiprocessor 里包含了8个stream processors，其组成是四个四个一组，也就是两组4D的处理器。每个 multiprocessor 还具有 很多个（比如8192个）寄存器，一定的（比如16KB） share memory，以及 texture cache 和 cons

目前 CUDA 装置中，每个 multiprocessor 有 16KB 的 shared memory。Shared memory 分成16 个 bank。如果同时每个 thread 是存取不同的 bank，就不会产生任何问题，存取 shared memory 的速度和存取寄存器相同。不过，如果同时有两个（或更多个） threads 存取同一个bank 的数据，就会发生 bank conflic

问： __global__ static void HelloCUDA(  )与__global__  void HelloCUDA(  )有什么区别吗？static关键字可以放在__global__之前吗？如果可以_global__ static与static _global__ 表达的是同一个意思吧？ 答： （1）有区别的。加上了static修饰后将限制符号的作用范围在

之前在写程序的时候，经常用弄混同步函数，现做出总结。_syncthreads（）：线程块内线程同步；保证线程会肿的所有线程都执行到同一位置； 当整个线程块走向同一分支时才可以使用_syncthreads()，否则造成错误；一个warp内的线程不需要同步；即当执行的线程数小于warpsize时，不需要同步函数，调用一次至少需要四个时钟周期，一般需要更多时钟周期，应尽量避免使用。

静态分配加上前缀 shared__shared__ int _ss[1024];1动态分配当我们在编程时，不清楚shared memory 数组开多大，就要用到动态分配。 分为两部分： 1， 声明extern __shared__ int _s[];12, 在调用kernel 时加上数组的大小。xxx_kernelgrid, block, sha

隐藏内存访问延迟，使用相同的说明方式，但针对内存操作。　　所需并行度 = 延迟 * 吞吐量　　　　所以隐藏内存延迟意味着保持100KB的数据读取速率，当然如果kernel是计算限制(compute bound)的，则这个数值可以变小。　　那么，多少线程可以达到100KB呢，有多种方法：　　1. 使用更多的线程　　2. 使用指令级并行(每个线程进行更多

这两天看到Vasily Volkov的ppt，对如何更有效的使用GPU做通用计算提出了增加线程级并行以外的另一种方法---增加指令级并行，受益匪浅。刚好也没找到中文版本，就翻译过来与大家交流下，有错误请各位指正，所有的实验结果和图片均出自原ppt。请参考《Better Performance at Lower Occupancy》后面两个案例。　　以下为译文：　　为提升GPU的效率

CUDA Online英伟达官方培训网站：英文, 中文NVIDIA DEVELOPER ZONE：https://developer.nvidia.com/GPU技术会议（GTC）资料：http://on-demand-gtc.gputechconf.com/gtcnew/on-demand-gtc.phpGTC：http://www.gputechconf.com/GTC中国：http:

问：对于结构体数组 typedef struct{float x; float y; float z;}  float3;  float3 d=data[id], id为线程索引号，则对于一个wrap，为其中的每个线程读取4字节需要几次访存呢？为其中的每一个线程读取12字节需要几次访存呢？答：如果只是为warp中的每个线程读1个该结构体的实例的一个分量（4B）的话。那么如果所有线程在都访问

在之前我们采取了两个主要的措施分别取隐藏和减少latency：1 . 我们一方面通过大量线程并行的方法去不断读取内存（当一个 thread 读取内存，开始等待结果的时候，GPU 就可以立刻切换到下一个 thread，并读取下一个内存位置）来尽可能的隐藏latency。2 . 另一方面我们采取了连续的内存存取模式，尽量减少latency，关于所谓的连续存储我们再详细说明一下：其实更精确的

其实这两天一直不知道什么叫bank conflict冲突，这两天因为要看那个矩阵转置优化的问题，里面有讲到这些问题，但是没办法，为了要看懂那个bank conflict冲突，我不得不去找资料，说句实话我现在不是完全弄明白，但是应该说有点眉目了，现在我就把网上找的整理一下，放在这边，等哪天完全弄明白了我就在修改里面的错误。    Tesla 的每个 SM 拥有 16KB 共享存储器，用于