06 2015 档案
摘要:CUDA SHARED MEMORYshared memory在之前的博文有些介绍,这部分会专门讲解其内容。在global Memory部分,数据对齐和连续是很重要的话题,当使用L1的时候,对齐问题可以忽略,但是非连续的获取内存依然会降低性能。依赖于算法本质,某些情况下,非连续访问是不可避免的。使用...
阅读全文
摘要:主成分分析(PCA)是用来提升无监督特征学习速度的数据降维算法。看过下文大致可以知道,PCA本质是对角化协方差矩阵,目的是让维度之间的相关性最小(降噪),保留下来的维度能量最大(去冗余),PCA在图像数据的降维上很实用,因为图像数据相邻元素的相关性是很高的。为了方便解释,我们以二维数据降一维为例(实...
阅读全文
摘要:CUDA Libraries简介上图是CUDA 库的位置,本文简要介绍cuSPARSE、cuBLAS、cuFFT和cuRAND,之后会介绍OpenACC。cuSPARSE线性代数库,主要针对稀疏矩阵之类的。cuBLAS是CUDA标准的线代库,不过没有专门针对稀疏矩阵的操作。cuFFT傅里叶变换cuR...
阅读全文
摘要:Memory Access Patterns大部分device一开始从global Memory获取数据,而且,大部分GPU应用表现会被带宽限制。因此最大化应用对global Memory带宽的使用时获取高性能的第一步。也就是说,global Memory的使用就没调节好,其它的优化方案也获取不到什...
阅读全文
摘要:nnCostFunction消耗公式:a1 = [ones(m,1) X];z2 = a1*Theta1';pre = sigmoid(a1*Theta1');a2 = [ones(m,1) pre];z3 = a2*Theta2';a3 = sigmoid(z3);y_vec = zeros(m,...
阅读全文
该文被密码保护。
摘要: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...
阅读全文