gpu
文章平均质量分 56
jieleiping
be honest...
展开
-
CUDA编程07——向量求和(并行规约,交错配对)
在CUDA编程06——向量求和(并行规约,相邻配对)中介绍了最简单向量求和规约算法。这里补充一个概念:关于warp和half-warp一个warp包含32个threads。warp是调度和执行的基本单位,half-warp是存储器操作的基本单位,这两个非常重要。因此上例中可以看到,由于采用了相邻配对,大多数的累加线程都是在不同的warp内,因此会需要更多的调度开销。因此可以考虑采用交错配对的方式,让那些已经完成了求和任务的线程不在调度。如下图所示:核函数代码// de.原创 2021-12-03 22:55:03 · 622 阅读 · 0 评论 -
CUDA编程06——向量求和(并行规约,相邻配对)
并行规约通常用于处理大输入数据集,将一组输入值规约一个值。数据特点:(1)对于数据集中的元素没有顺序要求。(2)可将数据分为若干小集合,每个线程处理一个集合。操作可以是:求最大值(Max)、求最小值(Min)、求和(Sum)、求乘(Product)。并行规约求和规约求和是常见应用,将输入数据求和得到一个值。如下面简单例子所示:规约求和的最简单思想是:先两两求和,然后再两两直至得到最后结果。核函数代码// device code__global__ vo..原创 2021-12-03 22:05:34 · 1615 阅读 · 0 评论 -
CUDA编程05——矩阵相乘 (shared memory)
在CUDA编程04——矩阵相乘 (去除长度限制)CUDA编程03——矩阵相乘CUDA编程04——矩阵相乘 (去除长度限制)中,另外一个问题是kernel 函数中存在很多global memory的读写操作。这些操作主要是一些重复读取,例如在计算目标矩阵C中的一列元素时,每一个元素的计算都需要读取矩阵A中的一行和B中的一列。行每次都在变化,但是列每次都是一样的。同理计算目标矩阵C中的一行元素时也一样。因此,如果可以共享这些部分,则可以减少对于global memory的访问,提升吞吐量。Shared M.原创 2021-11-27 22:42:56 · 1577 阅读 · 1 评论 -
CUDA编程04——矩阵相乘 (去除长度限制)
在CUDA编程03——矩阵相乘最后,指出了存在长度限制的缺点。本例中我们降尝试解决这个问题。具体做法是:我们将输出矩阵拆分成小块(Tile),把拆分的每个小块放到一个block中,然后通过threadIdx和blockIdx来索引。CPU计算的代码不变,只是现在的矩阵维度从32x32变成了1024x1024。GPU计算代码:__global__ void matrix_mul_device(FLOAT* M, FLOAT* N, FLOAT* P, int width){ in..原创 2021-11-27 17:30:16 · 542 阅读 · 0 评论 -
CUDA编程03——矩阵相乘
矩阵相乘的图示如下:CPU代码显然,最直接的矩阵相乘算法,需要三重循环,代码如下:void matrix_mul_host(FLOAT* M, FLOAT* N, FLOAT* P, int width){ for (int i = 0; i < width; ++i) { for (int j = 0; j < width; ++j) { float sum = 0, a, b; for (int k = 0; k < width; +..原创 2021-11-23 15:49:26 · 1275 阅读 · 0 评论 -
NVIDA CUDA architecture查询
NVIDA CUDA architecture查询原创 2016-11-09 23:14:43 · 11425 阅读 · 1 评论