随笔分类 - c/c++
摘要：Stream 一般来说，cuda c并行性表现在下面两个层面上： Kernel level Grid level 到目前为止，我们讨论的一直是kernel level的，也就是一个kernel或者一个task由许多thread并行的执行在GPU上。Stream的概念是相对于后者来说的，Grid le
阅读全文
摘要：CONSTANT MEMORY constant Memory对于device来说只读但是对于host是可读可写。constant Memory和global Memory一样都位于DRAM，并且有一个独立的on-chip cache，比直接从constant Memory读取要快得多。每个SM上c
阅读全文
摘要：findClosestCentroids.mm = size(X,1);for i=1:m [value index] = min(sum((repmat(X(i,:),K,1)-centroids).^2,2)); idx(i) = index;endcomputeCentroids....
阅读全文
摘要：CUDA SHARED MEMORYshared memory在之前的博文有些介绍，这部分会专门讲解其内容。在global Memory部分，数据对齐和连续是很重要的话题，当使用L1的时候，对齐问题可以忽略，但是非连续的获取内存依然会降低性能。依赖于算法本质，某些情况下，非连续访问是不可避免的。使用...
阅读全文
摘要：CUDA Libraries简介上图是CUDA 库的位置，本文简要介绍cuSPARSE、cuBLAS、cuFFT和cuRAND，之后会介绍OpenACC。cuSPARSE线性代数库，主要针对稀疏矩阵之类的。cuBLAS是CUDA标准的线代库，不过没有专门针对稀疏矩阵的操作。cuFFT傅里叶变换cuR...
阅读全文
摘要：Memory Access Patterns大部分device一开始从global Memory获取数据，而且，大部分GPU应用表现会被带宽限制。因此最大化应用对global Memory带宽的使用时获取高性能的第一步。也就是说，global Memory的使用就没调节好，其它的优化方案也获取不到什...
阅读全文
摘要：Memorykernel性能高低是不能单纯的从warp的执行上来解释的。比如之前博文涉及到的，将block的维度设置为warp大小的一半会导致load efficiency降低，这个问题无法用warp的调度或者并行性来解释。根本原因是获取global memory的方式很差劲。众所周知，memory...
阅读全文
摘要：Dynamic Parallelism到目前为止，所有kernel都是在host端调用，GPU的工作完全在CPU的控制下。CUDA Dynamic Parallelism允许GPU kernel在device端创建调用。Dynamic Parallelism使递归更容易实现和理解，由于启动的配置可以...
阅读全文
摘要：Avoiding Branch Divergence有时，控制流依赖于thread索引。同一个warp中，一个条件分支可能导致很差的性能。通过重新组织数据获取模式可以减少或避免warp divergence（该问题的解释请查看warp解析篇）。The Parallel Reduction Probl...
阅读全文
摘要：Exposing Parallelism这部分主要介绍并行分析，涉及掌握nvprof的几个metric参数，具体的这些调节为什么会影响性能会在后续博文解释。代码准备下面是我们的kernel函数sumMatrixOnGPUD：__global__ void sumMatrixOnGPU2D(float...
阅读全文
摘要：Warp逻辑上，所有thread是并行的，但是，从硬件的角度来说，实际上并不是所有的thread能够在同一时刻执行，接下来我们将解释有关warp的一些本质。Warps and Thread Blockswarp是SM的基本执行单元。一个warp包含32个并行thread，这32个thread执行于S...
阅读全文
摘要：GPU架构 SM（Streaming Multiprocessors）是GPU架构中非常重要的部分，GPU硬件的并行性就是由SM决定的。 以Fermi架构为例，其包含以下主要组成部分： CUDA cores Shared Memory/L1Cache Register File Load/Store
阅读全文
摘要：device管理NVIDIA提供了集中凡是来查询和管理GPU device，掌握GPU信息查询很重要，因为这可以帮助你设置kernel的执行配置。本博文将主要介绍下面两方面内容：CUDA runtime API functionNVIDIA系统管理命令行使用runtime API来查询GPU信息你可...
阅读全文
摘要：前言 线程的组织形式对程序的性能影响是至关重要的，本篇博文主要以下面一种情况来介绍线程组织形式： 2D grid 2D block 线程索引 矩阵在memory中是row-major线性存储的： 在kernel里，线程的唯一索引非常有用，为了确定一个线程的索引，我们以2D为例： 线程和block索引
阅读全文
摘要：CUDA简介CUDA是并行计算的平台和类C编程模型，我们能很容易的实现并行算法，就像写C代码一样。只要配备的NVIDIA GPU，就可以在许多设备上运行你的并行程序，无论是台式机、笔记本抑或平板电脑。熟悉C语言可以帮助你尽快掌握CUDA。CUDA编程CUDA编程允许你的程序执行在异构系统上，即CUP...
阅读全文
摘要：本篇博文仅实现hello world，先看到效果，具体细节将在后续博文解释。准备如果你是第一次使用CUDA，在Linux下可以使用下面的命令来检查CUDA编译器是否安装正确：$ which nvcc一般，该指令输出为：/usr/local/cuda/bin/nvcc另外，你可能还需要检查下你机器上的...
阅读全文
摘要：同一编译单元内部重名符号在编译期便被阻止，而不同编译单元之间重名符号要到链接期才会被发现。编译单元：#include 完头文件的内容之后（即将头文件内容粘贴到cpp中之后）的cpp文件就是编译单元，简单说便是经过预处理的cpp文件。条件编译是为了防止同一个.c文件中包含多个相同的.h文件。C++ p...
阅读全文
摘要：iconv命令实现linux下字符集编码的转换windows下的文件复制到linux下时常会乱码，因为windows下文件编码为GBK，linux下默认文件编码为UTF-8，故需要libiconv库转码。1.iconv命令用法如下：iconv [选项...] [文件...]输入/输出格式规范：-f,...
阅读全文
摘要：Linux下静态库生成和使用一.静态库概念1.库是预编译的目标文件(objectfiles)的集合，它们可以被链接进程序。静态库以后缀为”.a”的特殊的存档（archivefile）存储。2.标准系统库可在目录/usr/lib与/lib中找到。比如，在类Unix系统中C语言的数序库一般存储为文件/u...
阅读全文
摘要：c++声明与定义 声明是将一个名称引入程序。定义提供了一个实体在程序中的唯一描述。声明和定义有时是同时存在的。 如 inta; externintb=1; 只有当extern中不存在初始化才是声明。其他情况既是定义也是声明。 但是在下列情况下，声明仅仅是声明：1：仅仅提供函数原型。如voi...
阅读全文