每个CUDA线程计算矩阵的一行与向量乘积staticvoid-IT168.PPT

矩阵与向量乘法的CUDA优化 风辰 2010年12月11日 2011年1月8日修订 * 目的 对于CUDA程序开发来说，优化往往是整个开发过程的核心，不同算法，不同存储器组织的程序性能往往差几十倍，本文通过一个简单的例子来展示CUDA开发中一些重要的因素对性能的影响。 * 假设读者拥有以下知识 拥有C语言编程的经验,最好拥有并行编程经验 懂得CUDA,最好用CUDA写过代码 * 测试环境 Intel xeon 5405 2.0 GHz Geforce GTX 295(只使用单核) Gcc 4.3.3 CUDA toolkit 3.1 只测试计算时间,不包括数据传输 * 符号说明 matrix:矩阵数据指针,以行为主序或者列为主序存储 v || vec: 向量指针 r: 矩阵和向量乘的结果指针 rowSize: 表示矩阵的行数,也是r的长度 columnSize:表示矩阵的列数,也是v的长度 所有指向显存的指针加前缀d_ * 编译配置 矩阵尺寸8192*8192　单精度 编译选项-O3 –funroll-loops –msse CPU计时函数采用gettimeofday, clock,GPU计时函数采用CUDA event * 串行C版本 算法:遍历矩阵行,每行和向量相乘,最终结果为一向量 void mxv(const int rowSize, const int columnSize, const float *matrix, const float *v, float *r){ for(int i = 0; i rowSize; i++){ float re = 0.0f; for(int j = 0; j columnSize; j++){ re += matrix[i*columnSize+j]*v[j]; } r[i] = re; } } * 运行时间120 ms,不使用-O3运行耗时490 ms 简单SSE版本 算法:利用sse指令计算矩阵每行和向量的乘积 void mxvSSE(const int rowSize, const int columnSize, const float *matrix, const float *v, float *r){ __m128 *mv = (__m128*)v; __m128 *mm = (__m128*)matrix; for(int i = 0; i rowSize; i++){ __m128 re = _mm_set_ps(0.0f, 0.0f, 0.0f, 0.0f); for(int j = 0; j columnSize/4; j++){ re = _mm_add_ps(re, _mm_mul_ps(mm[i*columnSize/4+j], mv[j])); } float __attribute((aligned(16))) a[4]; _mm_store_ps(a, re); r[i] = a[0] + a[1] + a[2] + a[3]; } } 运行时间99ms * SSE + openmp 算法:使用二线程并行计算行循环 void mxvSSEOpenmp(const int rowSize, const int columnSize, float *matrix, float *vec, float *r){ __m128 *mv = (__m128*)v; __m128 *mm = (__m128*)matrix; #pragma omp parallel for num_threads(2) for(int i = 0; i rowSize; i++){ __m128 re = _mm_set_ps(0.0f, 0.0f, 0.0f, 0.0f); for(int j = 0; j columnSize/4; j++){ re = _mm_add_ps(re, _mm_mul_ps(mm[i*columnSize/4+j], mv[j])); } float __attribute((aligned(16))) a[4]; _mm_store_ps(a, re); r[i] = a[0] + a[1] + a[2] + a[3]; } 运行时间50ms *