算法加速优化瓶颈如何解决

  • 3 replies
  • 353 views
算法加速优化瓶颈如何解决
« 于: 五月 11, 2019, 08:20:28 pm »
如题所示,刚接触cuda几个月,算法加速遇到瓶颈了,不知如何进行下去, :'(系统是win10,软件是vs2017+cuda10.0,GPU是P5000;核函数执行时dim3 Grid1(4, 4);dim3 Block1(32, 32)); 该算法位于for循环中(有数据依赖),需要运行100多次;
下面说下我算法的具体情况:
第一点说的是此算法的具体过程以及原始代码,
第二点是是基于原始代码做的一点改进;
第三点说的是visual profiler中的核函数测试的时间;
第四点是我的问题。
1.算法其实是光学中由zenike系数生成波面的算法,很简单,见下图(第一次发帖,不知道怎么加图片,附件中上传了图),输入是个1*9向量,输出是1024*1024的矩阵。1024矩阵中每个元素(i,j)的值由以下过程得到:  a.将0~1023取值的i和j的值变换到-1~1的p,q;   b.由p和q的值计算出一个1*9的向量;   c.将b中计算出的向量与输入向量相乘,即得到输出矩阵的(i,j)的值。核函数代码如下,d_dz为输入向量,Z_Num是向量维数,此处是9,d_wave是输出矩阵,S是矩阵大小,此处是1024:
__global__ void WAVE(float *d_wave, int S, float *d_dz, int Z_Num)
{
   int x = threadIdx.x + blockIdx.x*blockDim.x;//表示x列
   int y = threadIdx.y + blockIdx.y*blockDim.y;//表示y行
   int nx = gridDim.x*blockDim.x;
   int ny = gridDim.y*blockDim.y;
   float p, q, r, c2, s2, c3, s3;
   float *Z = (float*)malloc(Z_Num * sizeof(float));

   for (int i = x; i < S; i += nx)
   {
      p = (float)2.0 * (i + 1) / S - 1;
      for (int j = y; j < S; j += ny)
      {
         q = (float)2.0 * (j + 1) / S - 1;
         r = p * p + q * q;
         if (r < 1)
         {            
            c2 = p * p - q * q;
            s2 = 2 * p*q;
            c3 = p * c2 - q * s2;
            s3 = q * c2 + p * s2;
            Z[0] = 2.0*p;
            Z[1] = 2.0*q;
            Z[2] = sqrt((float)3)*(2 * r - 1);
            Z[3] = sqrt((float)6)*s2;
            Z[4] = sqrt((float)6)*c2;
            Z[5] = sqrt((float)8 )*q*(3 * r - 2);
            Z[6] = sqrt((float)8 )*p*(3 * r - 2);
            Z[7] = sqrt((float)8 )*s3;
            Z[8] = sqrt((float)8 )*c3;
            float temp = 0;
            for (int k = 0; k < Z_Num; k++)
            {
               temp += Z[k] * d_dz[k];
            }
            d_wave[i*S + j] = temp;
         }
      }
   }
   delete[]Z;
}
2.上面的核函数运行很慢,所以做了点小改进,即在循环之前就将计算出1中b步骤中的向量,得到的是个1024*1024*9的一个矩阵,矩阵我是采用一维存储的;然后循环中直接将这个读取这个1024*1024*9的矩阵计算波面矩阵d_wave;另外由于一个核函数中d_dz需要多次读取,可将其设为共享存储,具体代码如下:
这是循环前计算1024*1024*9矩阵d_ZArray的核函数:
__global__ void SWAVE(float *d_ZArray, int S, int Z_Num)
{
   int x = threadIdx.x + blockIdx.x*blockDim.x;//表示x列
   int y = threadIdx.y + blockIdx.y*blockDim.y;//表示y行
   int nx = gridDim.x*blockDim.x;
   int ny = gridDim.y*blockDim.y;

   float p, q, r, c2, s2, c3, s3;
   for (int i = x; i < S; i += nx)
   {
      p = (float)2.0 * (i + 1) / S - 1;
      for (int j = y; j < S; j += ny)
      {
         q = (float)2.0 * (j + 1) / S - 1;
         r = p * p + q * q;
         if (r < 1)
         {         
            c2 = p * p - q * q;
            s2 = 2 * p*q;
            c3 = p * c2 - q * s2;
            s3 = q * c2 + p * s2;
            d_ZArray[(i*S + j)*Z_Num + 0] = 2.0*p;
            d_ZArray[(i*S + j)*Z_Num + 1] = 2.0*q;
            d_ZArray[(i*S + j)*Z_Num  + 2] = sqrt((float)3)*(2 * r - 1);
            d_ZArray[(i*S + j)*Z_Num  + 3] = sqrt((float)6)*s2;
            d_ZArray[(i*S + j)*Z_Num  + 4] = sqrt((float)6)*c2;
            d_ZArray[(i*S + j)*Z_Num  + 5] = sqrt((float)8 )*q*(3 * r - 2);
            d_ZArray[(i*S + j)*Z_Num  + 6] = sqrt((float)8 )*p*(3 * r - 2);
            d_ZArray[(i*S + j)*Z_Num  + 7] = sqrt((float)8 )*s3;
            d_ZArray[(i*S + j)*Z_Num  + 8] = sqrt((float)8 )*c3;
         }
      }
   }
}

这是循环中调用的核函数:
__global__ void WAVE0(float *d_wave, float *d_ZArray, float *d_dz, int S, int Z_Num)
{
   int x = threadIdx.x + blockIdx.x*blockDim.x;//表示x列
   int y = threadIdx.y + blockIdx.y*blockDim.y;//表示y行
   int nx = gridDim.x*blockDim.x;
   int ny = gridDim.y*blockDim.y;

   extern __shared__ float dz[];
   int id = threadIdx.x;
   if(id<Z_Num)
   {
      dz[id]=d_dz[id];
   }
   __syncthreads();

   for (int i = x; i < S; i += nx)
   {
      for (int j = y; j < S; j += ny)
      {
         float temp = 0;
         for (int k = 0; k < Z_Num; k++)
         {
            temp += d_ZArray[(i*S + j)*Z_Num + k] * d_dz[k];
         }
         d_wave[i*S + j] = temp;
      }
   }
}
3. visual profiler中release版程序测试结果:1中原始核函数WAVE单次需要1ms左右;2中核函数WAVE0单次执行需要640us,占据总时间的40%左右,依旧十分耗时
4.求助的问题是:a.对于1中的算法还能进行哪种更好的优化 b.对于改进后的2中算法还能从哪方面继续突破加速瓶颈 c.cuda中矩阵存储方式对性能影响有多大,比方说这里的ZArray矩阵我是采用一维存储的,如果用三维存储(三维存储有点麻烦,所以就没尝试了),预计,大概能有多少提升?麻烦大家帮忙回答时尽量详细点,感激不尽!

Re: 算法加速优化瓶颈如何解决
« 回复 #1 于: 五月 13, 2019, 02:17:43 pm »
周末不在,所以才看到。我说一下对你的代码几点考虑:

你的代码等于是在做一个特殊的滤波,滤波的参数只是需要根据S参数(编译时刻未知的浮点值),和Z_Num才能确定(我猜测Z_Num可能等于8,但是编译器不一定能知道)。则基于这种理解,显然的有如下考虑:

没有必要将临时的滤波参数保存到device side的heap上(你的kernel内部的分配和释放),这种分配和释放非常缓慢。也没有必要进行无辜的读取和写入,因为该参数的唯一用途是进行某种点乘,直接用掉了就可以了。这样可以抵消一次写入和一次读取。而一旦直接用掉了,则也节省了对应的空间的分配和释放(无论是从device端还是从host端进行的)。

你可能会说,这个参数可能很大,我需要手工的分配空间。那么此时,我依然建议你就地直接使用,因为如果在实际卡上的实际编译的时候,编译器认为的确很大(而不是你说的,或者你估计的,因为这种估计往往不准),编译器会自动考虑使用local memory. 而local memory自带跨线程合并效果(参考本论坛sisiy妹子的阅读CUDA手册100天系列中的,关于local memory的描述),你自行分配的,确是在线程间strided的访问。更何况,编译器很可能会直接消除对中间结果的保存写入+重新读取的完全使用。

你的第二个版本中的对shared memory的使用是很好的,这里没有什么说的,但是你将kernel拆分成两步了,一旦拆分成两步后,这种中间值的计算-保存-写入就无可避免了,所以应当考虑避免这种情况。而保留你对shared memory d的使用。

那么综合起来看,应当直接写成:统一的1个kernel,里面就地计算一个参数值,就地点乘累加一次(用你的第二个半分中的shared使用部分)。这样是最好的。

其次,当S是编译时刻未知的时候,编译器可能会生成昂贵的除法操作,这种操作可能在你修改了上面(采用上面几段的建议后)无法被掩盖(这个不一定,要看profiler最终分析结果)。则如果S实际上是你知道的常量(而编译器不知道),我建议写成const your_type S = ....; 这样编译器可以直接将相应的除法操作,尽量的在编译时刻就转换成乘法操作。此时可能会有正面效果(这是次要建议)。你也可以用#define来让编译器知道这是一个常量。很多滤波操作都会吃亏在这里。

大致这样。欢迎反馈。

EDIT: 修改了几个错别字。
« 最后编辑时间: 五月 13, 2019, 02:19:54 pm 作者 屠戮人神 »

Re: 算法加速优化瓶颈如何解决
« 回复 #2 于: 五月 15, 2019, 04:46:25 pm »
周末不在,所以才看到。我说一下对你的代码几点考虑:

你的代码等于是在做一个特殊的滤波,滤波的参数只是需要根据S参数(编译时刻未知的浮点值),和Z_Num才能确定(我猜测Z_Num可能等于8,但是编译器不一定能知道)。则基于这种理解,显然的有如下考虑:

没有必要将临时的滤波参数保存到device side的heap上(你的kernel内部的分配和释放),这种分配和释放非常缓慢。也没有必要进行无辜的读取和写入,因为该参数的唯一用途是进行某种点乘,直接用掉了就可以了。这样可以抵消一次写入和一次读取。而一旦直接用掉了,则也节省了对应的空间的分配和释放(无论是从device端还是从host端进行的)。

你可能会说,这个参数可能很大,我需要手工的分配空间。那么此时,我依然建议你就地直接使用,因为如果在实际卡上的实际编译的时候,编译器认为的确很大(而不是你说的,或者你估计的,因为这种估计往往不准),编译器会自动考虑使用local memory. 而local memory自带跨线程合并效果(参考本论坛sisiy妹子的阅读CUDA手册100天系列中的,关于local memory的描述),你自行分配的,确是在线程间strided的访问。更何况,编译器很可能会直接消除对中间结果的保存写入+重新读取的完全使用。

你的第二个版本中的对shared memory的使用是很好的,这里没有什么说的,但是你将kernel拆分成两步了,一旦拆分成两步后,这种中间值的计算-保存-写入就无可避免了,所以应当考虑避免这种情况。而保留你对shared memory d的使用。

那么综合起来看,应当直接写成:统一的1个kernel,里面就地计算一个参数值,就地点乘累加一次(用你的第二个半分中的shared使用部分)。这样是最好的。

其次,当S是编译时刻未知的时候,编译器可能会生成昂贵的除法操作,这种操作可能在你修改了上面(采用上面几段的建议后)无法被掩盖(这个不一定,要看profiler最终分析结果)。则如果S实际上是你知道的常量(而编译器不知道),我建议写成const your_type S = ....; 这样编译器可以直接将相应的除法操作,尽量的在编译时刻就转换成乘法操作。此时可能会有正面效果(这是次要建议)。你也可以用#define来让编译器知道这是一个常量。很多滤波操作都会吃亏在这里。

大致这样。欢迎反馈。

EDIT: 修改了几个错别字。
首先谢谢你的回答,
1.你说的S和Z_Num问题可以直接将二者定义为全局变量来解决,即#define S 1024;#define Z_Num 9;
2.不知道我们理解的拆分kernel函数的前后的具体情况有没有异议,详细说一下:
拆分前
for(int i=0;i<200;i++)
{
    WAVE<<<>>>();
}
拆分后
SWAVE<<<>>>();
for(int i=0;i<200;i++)
{
    WAVE<<<>>>();
}
也就是说拆分后SWAVE只需要在循环前执行一次;
你给的建议是不要拆分,仅仅是将WAVE核函数中引入共享内存的使用,其余部分都不变动,是这样吗?
如果是这样的话,应该是更慢,帖子中给了拆分前后耗时情况;visual profiler中release版程序测试结果:1中原始核函数WAVE单次需要1ms左右(这里没有使用共享内存,即便使用,加速效果也是很有限的);2中核函数WAVE0单次执行需要640us,所以拆分后应该是更快的,也就是说读取矩阵比计算矩阵中每个值要省时

纠正原帖子中手误,最后一个核函数中temp += d_ZArray[(i*S + j)*Z_Num + k] * d_dz[k];
应该改为temp += d_ZArray[(i*S + j)*Z_Num + k] * dz[k];

Re: 算法加速优化瓶颈如何解决
« 回复 #3 于: 五月 16, 2019, 10:17:55 am »
又一个笔误,第二个for循环中应该是 WAVE0<<<>>>();
我现在的想法是:
1.就目前实验数据来看还是拆分后更快,如果有新的方法,能让速度更快,那是最好了
2.对于拆分的方法还想继续优化,目前的想法就是能不能继续优化ZArray矩阵的读取速度,毕竟矩阵这么大,读取应该很耗时,目前还在看关于cuda存储方面的知识,希望你能给个意见,比方说这样做行不行得同?如果行,该怎么做,预计效果提升有多少?