找回密码
 立即注册

QQ登录

只需一步,快速开始

查看: 280|回复: 2

CUDA矩阵点乘即Matlab中的(.*)运算优化问题

[复制链接]
发表于 2018-6-9 10:47:52 | 显示全部楼层 |阅读模式
ESC4000G3
最近写的一个CUDA程序,运行时间达不到要求,利用NVIDIA Visual Profile分析是矩阵点乘耗费了大量的时间,想要再优化一下但是没有思路。
点乘的意思是矩阵的对应元素相乘,我要实现的点乘是一个是1024*1024*200的三维复数矩阵还有一个是1024*1024的二维复数矩阵,用这个二维矩阵重复去点乘三维矩阵1024*1024的那个面,最后得到的还是1024*10204*200的三维矩阵。
我的点乘代码如下:
  1. //点乘
  2. __global__ void Dot(cuFloatComplex *idata, cuFloatComplex *idata1, size_t pitch1, cuFloatComplex *odata1, unsigned int M, unsigned int N)
  3. {
  4.         int x = blockIdx.x * blockDim.x + threadIdx.x;
  5.         int y = blockIdx.y * blockDim.y + threadIdx.y;
  6.         int z = blockIdx.z * blockDim.z + threadIdx.z;
  7.         if (y < N)
  8.         {
  9.                 cuFloatComplex* row = (cuFloatComplex*)((char*)idata1 + y*pitch1);
  10.                 if ((x < N) && (z < M))
  11.                 {
  12.                         odata1[x + N*y + N*N*z].x = idata[x + N*y + N*N*z].x * row[x].x - idata[x + N*y + N*N*z].y * row[x].y;
  13.                         odata1[x + N*y + N*N*z].y = idata[x + N*y + N*N*z].x * row[x].y + idata[x + N*y + N*N*z].y * row[x].x;
  14.                 }
  15.         }
  16. }
复制代码
idata就是三维矩阵,idata1就是二维矩阵(用cudaMallocPitch和cudaMemcpy2D分配的)M=200,N=1024。用NVIDIA Visual Profile分析是330ms左右,由于程序里要用到12次,时间就花费了近4s。有哪位大神有优化的思路???求赐教!
回复

使用道具 举报

发表于 2018-6-9 16:38:28 | 显示全部楼层
Jetson TX2
Hi, chenjie13,

根据你的文字和代码来看, 这不是一个常规的点乘的过程, 我们常说的点乘是向量的多个分量进行乘法, 并累加得到一个标量.  而根据你的文字和代码看, 这更像是对于你的3D矩阵A的任何一个1024 * 1024 = 1M个元素的面As, 其中的每个元素As(i,j) * C(i,j)的过程, 其中C是一个同样的1024 * 1024的系数矩阵.

如果是基于这种理解, 而不是你的原始描述的话, 根据你的代码:
无论是200层的这个矩阵的任何一个面, 还是系数矩阵的那个面, 大小均为8MB(1024 * 1024 * 2 * sizeof(float)). 这将远远的超过现在的GPU的任何L2 cache的大小. 因此实际上每次乘法, 3D矩阵的每个面和系数的这个平面, 均将反复被从global memory载入. 最终的运行起来将是反复载入8MB * 200 + 8MB * 200 + 8MB * 200 = 约4.8GB每次"你的点乘". 实际上的SM中的SP将基本空载, 而显存将很忙碌.  Profiler将验证我的说法.

考虑到计算能力的不同, 这些写法不一定能充分满载显存.  我建议你分别进入如下步骤:
(1) 使用profiler看一下device memory的utilization程度, 如果不能接近满载. 你应当考虑1个线程计算多组数据(你的这里的2个元素的乘法, 分别得到实部和虚部的过程). 在特定的计算能力的卡上, 这样做是必须的.

(2)如果你的卡目前这个写法能够卡在显存访存的基本峰值. 则你应当考虑优化重复数据. 你有一个巨大的重复系数面(8MB), 每次和200个面的大矩阵中的每个面运算, 你的默认代码将每次都将重新载入一次这个系数面, 这是一种浪费. 你应当考虑通过shared memory之类的手段分批次的缓冲这个面(的一部分), 这样能理论上说最多减少50%的访存. 有大约200%的最多性能提升.

无论是采用了(1)(2)的哪种, 目前这个算法均将卡访存的. (运算量非常少, 基本都在访存上).
一个能实际大约达到300GB/s的卡, 你的这个代码理论说最佳能跑到约8MB * 401 =3.2GB, 3.2GB / 300GB/s = 大约10.x ms (例如1080), 实际能优化到跑到20ms以下就挺好.

Regards,
屠戮人神.
回复 支持 1 反对 0

使用道具 举报

 楼主| 发表于 2018-6-22 09:20:14 | 显示全部楼层
Tesla P100
屠戮人神 发表于 2018-6-9 16:38
Hi, chenjie13,

根据你的文字和代码来看, 这不是一个常规的点乘的过程, 我们常说的点乘是向量的多个分量 ...

谢谢@屠戮人神提供的思路,现在我的程序优化到了34ms,勉强达到了要求
回复 支持 反对

使用道具 举报

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

本版积分规则

关闭

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

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