随笔分类 -  CUDA

CUDA ---- Stream and Event
摘要：Stream 一般来说，cuda c并行性表现在下面两个层面上： Kernel level Grid level 到目前为止，我们讨论的一直是kernel level的，也就是一个kernel或者一个task由许多thread并行的执行在GPU上。Stream的概念是相对于后者来说的，Grid le 阅读全文
posted @ 2016-09-20 23:38 苹果妖 阅读(29568) 评论(3) 推荐(5) 编辑
CUDA ---- Constant Memory
摘要：CONSTANT MEMORY constant Memory对于device来说只读但是对于host是可读可写。constant Memory和global Memory一样都位于DRAM，并且有一个独立的on-chip cache，比直接从constant Memory读取要快得多。每个SM上c 阅读全文
posted @ 2015-08-06 00:25 苹果妖 阅读(9962) 评论(0) 推荐(0) 编辑
CUDA ---- Shared Memory
摘要：CUDA SHARED MEMORYshared memory在之前的博文有些介绍，这部分会专门讲解其内容。在global Memory部分，数据对齐和连续是很重要的话题，当使用L1的时候，对齐问题可以忽略，但是非连续的获取内存依然会降低性能。依赖于算法本质，某些情况下，非连续访问是不可避免的。使用... 阅读全文
posted @ 2015-06-28 14:35 苹果妖 阅读(34797) 评论(7) 推荐(3) 编辑
CUDA ---- CUDA库简介
摘要：CUDA Libraries简介上图是CUDA 库的位置，本文简要介绍cuSPARSE、cuBLAS、cuFFT和cuRAND，之后会介绍OpenACC。cuSPARSE线性代数库，主要针对稀疏矩阵之类的。cuBLAS是CUDA标准的线代库，不过没有专门针对稀疏矩阵的操作。cuFFT傅里叶变换cuR... 阅读全文
posted @ 2015-06-21 02:47 苹果妖 阅读(9635) 评论(5) 推荐(1) 编辑
CUDA ---- Memory Access
摘要：Memory Access Patterns大部分device一开始从global Memory获取数据，而且，大部分GPU应用表现会被带宽限制。因此最大化应用对global Memory带宽的使用时获取高性能的第一步。也就是说，global Memory的使用就没调节好，其它的优化方案也获取不到什... 阅读全文
posted @ 2015-06-13 15:21 苹果妖 阅读(6903) 评论(3) 推荐(4) 编辑
CUDA ---- Memory Model
摘要：Memorykernel性能高低是不能单纯的从warp的执行上来解释的。比如之前博文涉及到的，将block的维度设置为warp大小的一半会导致load efficiency降低，这个问题无法用warp的调度或者并行性来解释。根本原因是获取global memory的方式很差劲。众所周知，memory... 阅读全文
posted @ 2015-06-09 22:17 苹果妖 阅读(11535) 评论(4) 推荐(3) 编辑
CUDA ---- Dynamic Parallelism
摘要：Dynamic Parallelism到目前为止，所有kernel都是在host端调用，GPU的工作完全在CPU的控制下。CUDA Dynamic Parallelism允许GPU kernel在device端创建调用。Dynamic Parallelism使递归更容易实现和理解，由于启动的配置可以... 阅读全文
posted @ 2015-06-06 19:30 苹果妖 阅读(4904) 评论(0) 推荐(1) 编辑
CUDA ---- Branch Divergence and Unrolling Loop
摘要：Avoiding Branch Divergence有时，控制流依赖于thread索引。同一个warp中，一个条件分支可能导致很差的性能。通过重新组织数据获取模式可以减少或避免warp divergence（该问题的解释请查看warp解析篇）。The Parallel Reduction Probl... 阅读全文
posted @ 2015-06-02 23:43 苹果妖 阅读(5104) 评论(2) 推荐(2) 编辑
CUDA ---- Kernel性能调节
摘要：Exposing Parallelism这部分主要介绍并行分析，涉及掌握nvprof的几个metric参数，具体的这些调节为什么会影响性能会在后续博文解释。代码准备下面是我们的kernel函数sumMatrixOnGPUD：__global__ void sumMatrixOnGPU2D(float... 阅读全文
posted @ 2015-06-01 23:35 苹果妖 阅读(5822) 评论(3) 推荐(1) 编辑
CUDA ---- Warp解析
摘要：Warp逻辑上，所有thread是并行的，但是，从硬件的角度来说，实际上并不是所有的thread能够在同一时刻执行，接下来我们将解释有关warp的一些本质。Warps and Thread Blockswarp是SM的基本执行单元。一个warp包含32个并行thread，这32个thread执行于S... 阅读全文
posted @ 2015-05-31 00:02 苹果妖 阅读(31561) 评论(5) 推荐(6) 编辑
CUDA ---- GPU架构（Fermi、Kepler）
摘要：GPU架构 SM（Streaming Multiprocessors）是GPU架构中非常重要的部分，GPU硬件的并行性就是由SM决定的。 以Fermi架构为例，其包含以下主要组成部分： CUDA cores Shared Memory/L1Cache Register File Load/Store 阅读全文
posted @ 2015-05-30 03:04 苹果妖 阅读(22703) 评论(0) 推荐(3) 编辑
CUDA ---- device管理
摘要：device管理NVIDIA提供了集中凡是来查询和管理GPU device，掌握GPU信息查询很重要，因为这可以帮助你设置kernel的执行配置。本博文将主要介绍下面两方面内容：CUDA runtime API functionNVIDIA系统管理命令行使用runtime API来查询GPU信息你可... 阅读全文
posted @ 2015-05-30 00:25 苹果妖 阅读(10778) 评论(0) 推荐(0) 编辑
CUDA ---- 线程配置
摘要：前言 线程的组织形式对程序的性能影响是至关重要的，本篇博文主要以下面一种情况来介绍线程组织形式： 2D grid 2D block 线程索引 矩阵在memory中是row-major线性存储的： 在kernel里，线程的唯一索引非常有用，为了确定一个线程的索引，我们以2D为例： 线程和block索引 阅读全文
posted @ 2015-05-29 23:28 苹果妖 阅读(7491) 评论(4) 推荐(1) 编辑
CUDA ---- 简介
摘要：CUDA简介CUDA是并行计算的平台和类C编程模型，我们能很容易的实现并行算法，就像写C代码一样。只要配备的NVIDIA GPU，就可以在许多设备上运行你的并行程序，无论是台式机、笔记本抑或平板电脑。熟悉C语言可以帮助你尽快掌握CUDA。CUDA编程CUDA编程允许你的程序执行在异构系统上，即CUP... 阅读全文
posted @ 2015-05-28 22:36 苹果妖 阅读(7642) 评论(2) 推荐(1) 编辑
CUDA ---- Hello World From GPU
摘要：本篇博文仅实现hello world，先看到效果，具体细节将在后续博文解释。准备如果你是第一次使用CUDA，在Linux下可以使用下面的命令来检查CUDA编译器是否安装正确：$ which nvcc一般，该指令输出为：/usr/local/cuda/bin/nvcc另外，你可能还需要检查下你机器上的... 阅读全文
posted @ 2015-05-28 21:33 苹果妖 阅读(11570) 评论(6) 推荐(1) 编辑