随笔分类 - Heterogeneous Parallel Programming
胡文美的CUDA公开课,翻译,总结
摘要:MPI 模型如图MPI的各个运算节点是分布式的.每一个节点可以视为是一个“Thread”,但这里的不同之处在于这些节点没有所谓的共享内存,或者说Global Memory。所以,在后面也会看到,一般会有一个节点专门处理数据传输和分配的问题。MPI和CUDA的另一个不同之处在于MPI只有一级结构,即所...
阅读全文
摘要:C++ AMP一些更高级的概念:1. device内存的分配和拷贝.void vecAdd(float* A, float* B, float* C, int n){ array AA(n), BA(n); array CA(n); copy(A,AA); copy(B,BA); para...
阅读全文
摘要:C++ AMP是专为设计支持C++的异构并行模型.全程是:AcceleratorMassive Parallelism下面是一个Vector C++ AMP的代码,通过这段代码来解释C++ AMP的语法:#include using namespace concurrency;void vecAdd...
阅读全文
摘要:OpenACC:openacc 可以用于fortran, c 和 c++程序,可以运行在CPU或者GPU设备.openacc的代码就是在原有的C语言基础上进行修改,通过添加:compiler directives 编译器指令(pragmas): #pragma 来标示.cuda 中有__syncth...
阅读全文
摘要:什么是HistogrammingHistogramming是一种从大的数据集中提取典型特征和模式的方式.在统计学中,直方图(英语:Histogram)是一种对数据分布情况的图形表示,是一种二维统计图表,它的两个坐标分别是统计样本和该样本对应的某个属性的度量。图像直方图(英语:Image Histog...
阅读全文
摘要:stream是什么nivdia给出的解释是:A sequence of operations that execute in issue-order on the GPU. 可以理解成在GPU上执行的操作序列.比如下面的这些动作.cudaMemcpy()kernel launchdevice syn...
阅读全文
摘要:CPU和GPU内存交互在CUDA编程中,内存拷贝是非常费时的一个动作.从上图我们可以看出:1. CPU和GPU之间的总线bus是PCIe,是双向传输的.2. CPU和GPU之间的数据拷贝使用DMA机制来实现,非常容易理解,为了更快的传输速度.虚拟内存(virtual memory)我们都知道,虽然在...
阅读全文
摘要:和许多多线程并行问题一样,CUDA也存在互斥访问的问题,即当一个线程改变变量X,而另外一个线程在读取变量X的值,执行原子操作类似于有一个自旋锁,只有等X的变量在改变完成之后,才能执行读操作,这样可以保证每一次读取的都是最新的值.在kernel 程序中,做统计累加,都需要使用原子操作:atomicAd...
阅读全文
摘要:1. Prefix Sum前缀求和由一个二元操作符和一个输入向量组成,虽然名字叫求和,但操作符不一定是加法。先解释一下,以加法为例:第一行是输入,第二行是对应的输出。可以看到,Output[1] = Input[0] + Input[1],而Output[length - 1]就是整个输入向量元素之...
阅读全文
摘要:Reduction并行分析:每个线程是基于一个树状的访问模型,从上至下,上一层读取数据相加得到下一层的数据.不停的迭代,直到访问完所有的数据.利用这么多的线程块(thread block)我们需要做的事情如下:1. 处理非常大的数组2. 让GPU的每个处理器保持忙碌3. 每个thread block...
阅读全文
摘要:首先添加上Heterogeneous Parallel Programming class 中 lab: Reduction的代码:myReduction.c// MP Reduction// Given a list (lst) of length n// Output its sum = lst...
阅读全文
摘要:1. ReductionReduction是一种广泛使用的计算模型,特别是在并行计算领域。简单地来说,Reduction就是一系列的划分(Partition)和汇总(Summarize)操作的集合:对输入数据分块,对每一个分块汇总,然后再将汇总后的数据视为新的输入数据,重复分块和汇总,直到得到最终结...
阅读全文
摘要:分析tile并行算法的优化情况:一维卷积的复用情况分析比如8个元素的一维卷积tile优化.M的大小是5,计算8个元素的卷积需要载入 8+5-1 =12, 如果不使用tile,每个元素都需要载入 8*5 =40, 所以全局内存访问带宽减少 40/12 =3.3. 正常我们算40 = 8*5的方式来看,...
阅读全文
摘要:一维卷积Convolution卷积也是很常用的一种计算模式。卷积计算方法如下:对输出数据中的每一个元素,它的值是输入数据中相同位置上的元素与该元素周边元素的值的加权和。卷积中有一个被称为卷积核(Kernel)或卷积码(Mask)的数据段,指定了周边元素的权值。为了避免混淆,以后都称为卷积码。计算如下...
阅读全文
摘要:全局存储带宽(DRAM)全局内存是动态随机访问的方式访问内存.我们希望访问DRAM的时候非常快,实际情况是DRAM中出来的数据非常非常慢,这就好比,理想状态是泄洪,水倾巢而出,气势宏伟,实际取水却像是用吸管在喝饮料,速度非常慢.通常来看,我们会通过优化算法减少DRAM的访问次数.由上图可以看出,用户...
阅读全文
摘要:CPU 矩阵乘法能相乘的两个矩阵,必须满足一个矩阵的行数和第二个矩阵的列数相同.A(N*P) * B(P*M) = C(N*M). 其中P是行数,N是列数, 从宽高的角度来说,即 A的宽度和B的高度是相同的.C矩阵 = ha * wb.其中C(i,j) = A矩阵中的i行和B矩阵中的j列进行点乘得到...
阅读全文
摘要:在CUDA基本概念介绍有简单介绍CUDA memory。这里详细介绍:每一个线程拥有自己的私有存储器,每一个线程块拥有一块共享存储器(Shared memory);最后,grid中所有的线程都可以访问同一块全局存储器(global memory)。除此之外,还有两种可以被所有线程访问的只读存储器:常...
阅读全文
摘要:在HOST端我们会分配block的dimension, grid的dimension。但是对应到实际的硬件是如何执行这些硬件的呢?如下图:lanuch kernel 执行一个grid。一个Grid有8个block,可以有两个硬件执行单元,一个执行一个block,需要执行4次,或者像右边有4个执行单元...
阅读全文
摘要:1.Programming_Massively_Parallel_Processors.pdf2.CUDA_C_Programming_Guide.pdf3.CUDA范例精解通用GPU编程.pdf4.GPU高性能编程CUDA实战中文.pdf5.CUDA深入浅出.pdf
阅读全文