苹果妖

摘要： Memorykernel性能高低是不能单纯的从warp的执行上来解释的。比如之前博文涉及到的，将block的维度设置为warp大小的一半会导致load efficiency降低，这个问题无法用warp的调度或者并行性来解释。根本原因是获取global memory的方式很差劲。众所周知，memory... 阅读全文

摘要： Linear Regression with One Variablemodel Representation以上篇博文中的房价预测为例，从图中依次来看，m表示训练集的大小，此处即房价样本数量；x表示输入变量或feature（特征），此处即房子面积；y是输出变量或目标变量，此处即房子价格。（x，y）... 阅读全文

摘要： 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... 阅读全文

摘要： 该博文系列是Andrew NG教授的课程笔记,有兴趣的朋友可以在Coursera或者网易公开课上找到该课程.Supervised Learning下图是一个监督学习回归分析的例子。该图旨在预测房价。对于该场景，训练集都是有标签属性的。以上图为例，图中的每个点，都对应了房子的价格，监督学习就是要在这些... 阅读全文

摘要： 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索引 阅读全文