找回密码
 立即注册

QQ登录

只需一步,快速开始

查看: 153|回复: 4

关于复杂结构体的疑问

[复制链接]
发表于 2018-4-9 13:41:07 | 显示全部楼层 |阅读模式
ESC4000G3
显卡:GTX1080   cuda 8.0 想使用GPU来加速人脸特征的比较, 特征结构体是一个544字节的结构体,计算过程会使用到其中三个成员,有一个成员是特征值 513 Byte的数组
核函数运算之前已经申请了1500W 该结构体数组的显卡内存等待运算, 核函数的参数的一个结构体,这个1500W的数组,以及一个存放运算结果的1500W int类型数组。特征对比的过程大概是A1*B1 + ....A513*B513  再加上几个乘法和除法。目前的运算速度跟32核的CPU相比,相差无几(600+ms)。
每条线程都是读取不同的结构体。

启动参数是 <<<20, 1024>>>  换成<<<20*32,32>>>可以略微提速

我想问问,这种频繁读取全局内存是否会造成无法提速,以及该如何设计加速较好,谢谢大家。
回复

使用道具 举报

发表于 2018-4-9 19:06:34 | 显示全部楼层
Jetson TX2
Hi, 楼主,

收到你的帖子。目前没有你的代码,根据文字描述来分析,你:
有一堆513B (544B)的数据,数据类型应当为uint8_t或者int8_t, 和另外的一个用来做系数的513B单一数组。
然后你试图做的是,对每个这513个元素的uint8_t[]数组,乘以这个固定的单一系数数组,并累加为普通整数类型,也就是点乘。

在这个假设的分析的基础上(如果这个分析不对,请提供具体代码而不是文字描述),你分别有:

(1)标准的循环了1500W次(或者对应的OpenMP或者手工线程展开),核心部分为这个运算的,作为参考的CPU实现。
(2)直接展开为2W个CUDA线程的GPU实现,每个CUDA线程则做刚才的这个核心运算,同时余外需要再循环多次,补齐一共1500W这个计算。而无其他数据的结构上安排变化。

===那么在上面的对你的原帖的分析的基础上, 和得出的假设的可能情况, 则====
(A)这是一种最基本的直接将CPU实现展开为1个CUDA线程,毫无任何数据和代码变动的最简单改写。
(B)该改写在显存的访问上,进行了大跨步访问(每个CUDA线程可能分别访问:
int id = 0....1500W; //或者等效外层循环
sum = 0;
for (int t = 0; t < 513; t++) sum += A[id].SomeMember[t] * Coeff[t],
result[id] = sum;
类似这种风格的代码。

这样做是很不好的,跨步太大了。结构体本身倒是不复杂的。这样在CUDA上叫不能合并的global memory访问,在你的数据规模下,可能会导致严重的数据cache和TLB miss,严重的降低效率。根据你的数据量来看,等效你才13GB/s的显存速率, 是1080本身的峰值的4%不太到,问题可见一斑。

建议的解决方案:
(1)尝试将这1500W的这个struct实例中的这个513B成员打散,改写成按照CUDA线程ID方式的连续排布。也就是常说的struct of array或者array of structs的问题(SoA vs AoS).
这个建议也是类似问题提问在NV的英文官方论坛上,会得到的标准答案。
但是我不推荐你这样做。因为这需要改动数据的结构,同时可能会对前期的数据出来带来较大的代价。

(2)忘记方案1, 同时将每个线程处理原本的CPU的那个核心过程,改写成1个block(例如64个或者128个线程)共同处理,同时使用规约累加完成并行这个点乘过程。这样你可以额外得到64x - 128x的并行度,和几乎完全合并的global memory访问,有效降低TLB和数据缓存的劳累/抖动。

方案2是可以几乎在你0已有项目的改动下,简单的更换一下在CUDA上的数据映射关系(从1个线程处理1个点乘,到1个block处理一个),就几乎可以得到巨大的性能提升的最短方案。建议采纳。

不妨试验一下,应当会看到惊喜。

Regards,
屠戮人神

PS: 你作为CUDA新人, 建议至少看完CUDA C Programming Guide后再进行CUDA代码书写尝试。
PS: 本回复只对你的中文描述负责。不对你可能有的隐含的没有发出的代码负责。如果推测出来的信息和你的实际实现不同,请提供实际代码。
回复 支持 反对

使用道具 举报

 楼主| 发表于 2018-4-11 10:06:51 | 显示全部楼层
Tesla P100
屠戮人神 发表于 2018-4-9 19:06
Hi, 楼主,

收到你的帖子。目前没有你的代码,根据文字描述来分析,你:

谢谢您的建议。

本次使用了您的建议改动了代码,时间从650ms降低到160ms<<<3000,128>>>左右,而之前的方案改动 block 和 thread的配置后,速度最快也有到达120ms<<<32, 256>>>(之前点积过程手动展开)
  1. for(size_t i = bid; i < pstTask->nConfMaxFeature; i += pstTask->nConfBlocks)
  2.         {
  3.                
  4.                 pstBase = pstBaseFeature + i;

  5.         /*        if(pstBase->nFeatureSize <= 0)
  6.                 {
  7.                         continue;
  8.                 }*/

  9.                 //for (nFeatureIdx = 0; nFeatureIdx < pstTask->nValidFeatureNum; nFeatureIdx++)
  10.                 //{
  11.                 //        a_b_value[nFeatureIdx] = 0;
  12.                 //}

  13.                 for(size_t j = tid; j < 512; j += pstTask->nConfMaxThreads)
  14.                 {
  15.                         buffer[j] = pstBase->featureVal[j] * pstTask->pFeature[0][j];
  16.                 }
  17.                 __syncthreads();
  18.                 for (size_t s = 512 >> 1; s > 32; s >>= 1)
  19.                 {
  20.                         for(size_t k = tid; k < s; k += pstTask->nConfMaxThreads)
  21.                         {
  22.                                 buffer[k] += buffer[k + s];
  23.                         }
  24.                         __syncthreads();
  25.                 }

  26.                 if (tid < 32)
  27.                         warpReduce(buffer, tid);

  28.                 if(tid == 0)
  29.                 {
  30.                         pSimilarity[i] = buffer[0] / (pstBase->fNorm * pstTask->fNorm[0])) + 1.0) / 2) * 10000;
  31.                 }
  32.                 __syncthreads();
  33. }
复制代码


上面就是核函数代码/

我再解释下我的想法,程序需要从1500W的人脸中比较并拿出分数最高的1000个,这个因为没什么好的方案在核函数中做排序筛选,故使用了同样的1500W结果数组来存。


(这里我也不太清楚这些配置为什么会快了好多倍,而且无论增大或者减小都会降低速度,用NSight性能检测看到大部分issue stall都是在memory dependency 88%, No eligible占用 99%, 这是我原先方案通过改配置后发现的)
我想请教一下这个配置如何设置?查阅很多资料总是不得要领,掌握的知识还不够系统,总是对不上号。



回复 支持 反对

使用道具 举报

发表于 2018-4-11 10:27:17 | 显示全部楼层
考虑到楼主对我的说法含有顾虑。我将直接为你提供改好的代码。请等待。
回复 支持 反对

使用道具 举报

发表于 2018-4-11 12:21:02 | 显示全部楼层
4cl 发表于 2018-4-11 10:06
谢谢您的建议。

本次使用了您的建议改动了代码,时间从650ms降低到160ms左右,而之前的方案改动 block ...

既然楼主着急, 不妨看一下这个实现。这个实现也是我说的,通过block完成一组乘法。然后规约求和保存到global memory。根据问题的重点所在,去掉了无关紧要的部分。该kernel在1070上,一半规模(750W), 跑在24ms左右。这等效于你在1080上,1500W规模,小于48ms(甚至更多,因为1080访存快不少。SM也多)。这样从600ms到40多ms,至少证明我的建议不是对你的撒谎。代码如下:


  1. struct st58787
  2. {
  3.         uint8_t data[512]; //user data
  4.         uint8_t _unknown[32];
  5. };


  6. #define TABLES 7500000
  7. __global__ void dot_product(uint32_t * __restrict__ result, const struct st58787 * __restrict__ p, const uint8_t * __restrict__ coeff)
  8. {
  9.         __shared__ uint32_t partial[128];
  10.         partial[threadIdx.x] = 0;
  11.         __syncthreads();

  12.         for (int id = blockIdx.x; id < TABLES; id += gridDim.x)
  13.         {
  14.                 uint32_t t = 0;
  15.                 t += p[id].data[threadIdx.x] * coeff[threadIdx.x];
  16.                 t += p[id].data[threadIdx.x + 128] * coeff[threadIdx.x + 128];
  17.                 t += p[id].data[threadIdx.x + 256] * coeff[threadIdx.x + 256];
  18.                 t += p[id].data[threadIdx.x + 384] * coeff[threadIdx.x + 384];

  19.                 atomicAdd(&partial[threadIdx.x % 32], t);
  20.                 __syncthreads();

  21.                 if (threadIdx.x < 32)
  22.                 {
  23.                         t = partial[threadIdx.x];
  24.                         atomicAdd(&result[id], t);
  25.                 }
  26.         }
  27. }
复制代码

(代码仅供参考,请勿在测试前应用于生产环境。本论坛仅出于示范,不对此代码造成的任何经济和生命安全损失负责。)

几个需要说明的是:
(1)亲的代码中,不需要特意设定一个task描述符之类的东西。这没有必要。直接传递参数给kernel即可。类似的,blockIdx.x之类的东西直接用即可。不需要从host传递过来。这些都是华而不实的东西。
(2)此代码依然可以优化,虽然已经让你的代码从600ms到了40多ms,大约提速了15倍了。显存从4%的理论峰值到了大约60%了。但依然存在优化的空间。

以下工作请问亲:
(1)是否可以规避掉block内部的通信过程(例如在shared memory上的规约)从而可能的提高性能?
(2)每个线程只计算1个uint8_t,这样在读取数据和计算上,是否高效?如果不是,能否每个线程能计算,例如uint32_t, 从而能利用GTX1080自带的深度学习加速能力(__dp4a,对4个uint8_t进行乘法运算并直接累加求和)?
(3)如果(2)能做到,是否需要有结构体(你的结构体)在内存(显存)对齐上的要求?
(4)如果考虑了这点,能否继续直接扩大到每个warp中的每个线程读取16B(16组8-bit值);以及,这样是否能完全消除shared memory使用和之上的规约,从而能减少延迟,或者提高吞吐率?
(5)通过当前实现内的循环( for (int id = blockIdx.x; id < TABLES; id += gridDim.x)),哪些数据得到了重用?编译器通过何种措施,代码中的那些描述提示了它,从而让它能发现这些重用的数据?

好了。本想一步到位的,但那样对论坛和其他阅读者无用。
于是亲不要太着急,我也不直接弄完它。你我各退一步,欢迎继续。

预祝亲能从4%到60%后,继续提升(例如80%理论峰值?)

Regards,
屠戮人神


回复 支持 反对

使用道具 举报

您需要登录后才可以回帖 登录 | 立即注册

本版积分规则

关闭

站长推荐上一条 /1 下一条

快速回复 返回顶部 返回列表