当前位置:首页 > 话题广场 > 攻略专题 > 最新攻略

【0号元素】CUDA加速——基于规约思想的数组元素求和

数组元素的总和顾名思义就是数组中所有元素的总和。例如,阵列X。

x的所有元素如下:

如果按串行顺序上升,请按照step1到stepn的步骤进行操作,如下图所示:S0、S1、S2、Sn-1,结果Sn-1是所有元素的总和。

01

把规约和思想合二为一

如果按照上述串行顺序计算的步骤1 ~步骤1 ~步骤1 ~步骤1 ~步骤1 ~步骤1 ~步骤1 ~步骤1 ~步骤1 ~步骤1 ~步骤1 ~步骤1 ~步骤1 ~步骤3多个聚集步骤并行执行,而不是按先后顺序执行,该怎么办?

阶段1至阶段中后续阶段的计算结果取决于上一阶段的结果,因此必须等到上一阶段的计算结果出来后,才能开始后续阶段的计算。如果直接并行执行第1步至第n步,则无法获得正确的结果,因为计算将在上一步的结果出来之前开始。

为解上述问题,通常采用规约思想来并行计算。下面我们举一个简单的例子来说明规约思想,假设数组X有8个元素,现要使用规约思想求其元素和:

步骤如下:

  • 将x0~x7两两分组:(x0,x1)、(x2,x3)、(x4,x5)、(x6,x7)。然后使用4个线程并行计算每组的元素和,得到4个求和结果:

  • 再将上一步骤得到的4个求和结果两两分组:(S0,S1)、(S2,S3)。然后使用2个线程并行计算每组的元素和,得到2个求和结果:

  • 使用1个线程计算上一步骤得到的S4,S5之和,也即最终求和结果:

以上3个步骤可用下图表示,每个步骤都将数据两两分组,然后并行计算每组的元素和,最后得到一个结果,这就是规约的过程:


02

CUDA实现数组元素的规约求和

CUDA是为并行计算而生的,使用CUDA可以很容易实现上述的数组规约求和算法。不过有一点需要注意,就是必须确保每个步骤的所有线程是同步的,也即所有线程计算完成之后再进入下一步骤的计算,否则会导致结果错误。

比如假设上图的线程a、b、c已完成计算,但是线程d未完成计算,如果不等线程d完成计算就直接进入下一步骤计算S4和S5,由于线程d未完成计算,得到的S3是错误的值,这导致得到的S5也是错误的值,从而导致最终计算的S6也错了——这将导致错误的连锁反应。

在cuda中,可以调用__syncthreads函数方便地同步同一个线程块中的所有线程,因此我们可以使用同一个线程块中的多个线程做规约运算。那么问题来了,如果数据量很大,一个线程块不能完成所有数据的规约运算该怎么办呢?答案是分块处理,将数据平均分成多个部分,每部分都分配给一个线程块做规约运算。因此每个线程块最后得到一个规约结果,最后再将多个规约结果求和,即得到最后结果。如下图所示:

也许这里有人会问,最后得到的多个规约结果还是得按照串行顺序求和呀,这样做有点脱裤子放屁了。我想说是这样,只不过最后得规约结果相对原始数据个数,已经少了很多很多了,因此其串行求和的耗时基本可以忽略。

CUDA核函数代码如下:

#define N (1536 * 20480) //数据总长度 //Para为输入数组,长度为N //blocksum_cuda存储所有线程块的规约结果 __global__ void cal_sum_ker0(float *Para, float *blocksum_cuda) { //计算线程ID号 为线程块的ID号 每个线程块中包含的线程总个数 为每个线程块中的线程ID号 int tid = blockIdx.x * blockDim.x + ; if(tid < N) { for (int index = 1; index < blockDim.x; index = (index*2)) { if ( % (index*2) == 0) { Para[tid] += Para[tid + index]; //规约求和 } __syncthreads(); //同步线程块中的所有线程 } if( == 0) //整个数组相加完成后,将共享内存数组0号元素的值赋给全局内存数组0号元素,最后返回CPU端 blocksum_cuda[blockIdx.x] = Para[tid]; } }

上面代码的for循环,咋一看很难理解,没有关系,我们举个简单例子来说明就好了。假设:blockDim.x=8,也即每个线程块有8个线程;N=8,也即输入数组的长度为8。那么对于第0个线程块(其它线程块也类似),其包含线程id为0~7,计算过程如下,你是否已经发现,其计算过程就是上述讲的规约过程呀~

下面我们写代码来测试上方实现的规约算法是否正确:

首先,定义一个微秒级计时的类,用于计时:

class Timer_Us2 { private: LARGE_INTEGER cpuFreq; LARGE_INTEGER startTime; LARGE_INTEGER endTime; public: double rumTime; void get_frequence(void) { QueryPerformanceFrequency(&cpuFreq); //获取时钟频率 } void start_timer(void) { QueryPerformanceCounter(&startTime); //开始计时 } void stop_timer(char *str) { QueryPerformanceCounter(&endTime); //结束计时 rumTime = - ) * 1000.0f) / c); cout << str << rumTime << " ms" << endl; } Timer_Us2() //构造函数 { QueryPerformanceFrequency(&cpuFreq); } };

接着是测试函数:

void Cal_Sum_Test() { Timer_Us2 timer; //申请长度为N的float型动态内存 float *test_d = (float *)malloc(N * sizeof(float)); for (long long i = 0; i < N; i++) { test_d[i] = 0.5; //将所有元素赋值为0.5 } double Parasum = 0.0; (); //在CPU端按顺序计算数组元素和 for (long long i = 0; i < N; i++) { ParaSum += test_d[i]; //CPU端数组累加 } ("CPU time:"); cout << " CPU result = " << ParaSum << endl; //显示CPU端结果 //设置每个线程块有1024个线程 dim3 sumblock(1024); //设置总共有多少个线程块 dim3 sumgrid(((N%) ? (N/ + 1) : (N/))); float *test_d_cuda; float *blocksum_cuda; float *blocksum_host = (float *)malloc(sizeof(float) * ); //申请GPU端全局内存 cudaMalloc((void **)&test_d_cuda, sizeof(float) * N); cudaMalloc((void **)&blocksum_cuda, sizeof(float) * ); (); //将数据从CPU端拷贝到GPU端 cudaMemcpy(test_d_cuda, test_d, sizeof(float) * N, cudaMemcpyHostToDevice); //调用核函数进行规约求和 cal_sum_ker0 << < sumgrid, sumblock >> > (test_d_cuda, blocksum_cuda); //将每个线程块的规约求和结果拷贝到CPU端 cudaMemcpy(blocksum_host, blocksum_cuda, sizeof(float) * , cudaMemcpyDeviceToHost); //在CPU端对所有线程块的规约求和结果做串行求和 double sum = 0.0; for(int i = 0; i < ; i++) { sum += blocksum_host[i]; } ("GPU time:"); cout << " GPU result = " << sum << endl; //显示GPU端结果 //释放内存 cudaFree(test_d_cuda); cudaFree(blocksum_cuda); free(blocksum_host); free(test_d);

运行结果如下,可以看到对1536*20480长度的数组求元素和,CPU和GPU的计算结果是一致的,不过GPU CUDA计算耗时反而比CPU更多了,一方面是因为GPU计算多了host与device端内存拷贝的耗时,另一方面是因为我们实现的CUDA规约算法没有做到优化的极致,还有不小的优化空间。那么接下来让我们继续尝试优化吧~

1.《【0号元素】CUDA加速——基于规约思想的数组元素求和》援引自互联网,旨在传递更多网络信息知识,仅代表作者本人观点,与本网站无关,侵删请联系页脚下方联系方式。

2.《【0号元素】CUDA加速——基于规约思想的数组元素求和》仅供读者参考,本网站未对该内容进行证实,对其原创性、真实性、完整性、及时性不作任何保证。

3.文章转载时请保留本站内容来源地址,https://www.lu-xu.com/gl/2070672.html

上一篇

【2002八神出招表】拳皇中伤害爆表的五大必杀技,草薙京和八神庵当年根本排不上号

下一篇

2016英雄联盟全明星赛看这里!#2016NBA全明星# 东西部首发名单公布,科比当选票王

【0号元素】专题硅元素简介

【0号元素】专题硅元素简介

0号元素相关介绍,硅是化学元素,化学符号是Si,是以前的硅。原子序数14,相对原子质量28.0855,有无定形硅和晶体硅两种同素异形体,属于元素周期表的第三周期,IVA族的金属元素。 硅也是极为常见的一种元素,然而它极...

关于0号元素我想说氦元素简介

关于0号元素我想说氦元素简介

0号元素相关介绍,氦(Helium),最不活跃的元素,元素符号He是稀有气体的一种。 氦在通常情况下为无色、无味的气体,是唯一不能在标准大气压下固化的物质。氦的应用主要是作为保护气体、气冷式核反应堆的工作流体和超低温冷...

0号元素,干货看这篇!氢元素简介

0号元素,干货看这篇!氢元素简介

0号元素相关介绍,氢是元素符号H在元素周期表中排在第一位的化学元素。 氢通常的单质形态是氢气,无色无味无臭,是一种极易燃烧的由双原子分子组成的气体,氢气是最轻的气体。医学上用氢气来治疗疾病。 [1] 氢气的爆炸极限为4...

0号元素看这里!磷元素简介

  • 0号元素看这里!磷元素简介
  • 0号元素看这里!磷元素简介
  • 0号元素看这里!磷元素简介
【0号元素】专题碳元素简介

【0号元素】专题碳元素简介

0号元素相关介绍,碳(Carbon)是非金属元素,化学符号是C,在室温下具有稳定性,无反应,对人体毒性极低,可安全食用石墨或活性炭,是元素周期表的第二周期IVA族。 碳是一种很常见的元素,它以多种形式广泛存在于大气和地...

【0号元素】专题氟元素简介

0号元素相关介绍,Fluorine是非金属化学元素,化学符号为F,原子序数为9。 氟是卤族元素之一,属周期系ⅦA族,在元素周期表中位于第二周期。氟元素的单质是F2,它是一种淡黄色有剧毒的气体。氟气的腐蚀性很强,化学性质...

【0号元素】磷元素简介

  • 【0号元素】磷元素简介
  • 【0号元素】磷元素简介
  • 【0号元素】磷元素简介

0号元素专题之人类梦寐以求的金属氢,还要再等90年?

0号元素相关介绍,金属氢是未来高密度、高能储存材料,是人类梦寐以求的能源物质。 90多年来,人们一直在试图制造出以金属形态存在的氢气,并为此付出不懈努力,但稳定的金属氢样品始终没能得到。 从理论上来看,在超高压下得到金...