圣诞快乐！请问这个CUDA的程序，效率还能怎么再提高呢？

加速计算 CUDA
1

这个程序运行在一个I7 2600 CPU上，用时20ms。我的GPU是Geforce GTX 570，用时5ms。
所以我觉得还能再次提高速度

代码在2楼，谢谢各位大侠！！

2 
#include <atltime.h>
#include "cuda_runtime.h"
#include <stdio.h>
#include <stdlib.h>
#include <iostream>
#include <iomanip>
#include <time.h>


//行*列
#define ROW 3000
#define COL 2000
#define DATA_SIZE ROW*COL
#define THREAD_NUM 256

int* data;
int* result;

void GenerateNumbers(int *number, int size)
{
   time_t t;
   srand((unsigned) time(&t));
   for(int i = 0; i < size; i++) {
   number[i] = rand() % 10;
   }
}

//适应2D数组和pitch，把行的数据读入共享内存来处理
__global__ static void CheckInGPU5(size_t _row, size_t _col, int* _data, int* _result, size_t _duration, size_t _pitch_col)
{
   const int tid = threadIdx.x;
   const int row = blockIdx.x;

   if(row<_row)
   {
   //拷贝入共享内存
   extern __shared__ float shared1[];
   for(int i = tid; i < _col; i += blockDim.x) {
   shared1[i] = _data[row * _pitch_col + i];
   }
   __syncthreads();

   int _r=row;
   int _close1=0;
   int _close2=0;//当前日期的close，大的日期
   int _index1=0;
   int _index2=0;//大的日期对应的游标
   int _resultIndex=0;
   int _count2=0;
   bool _bRowStart=false;
   bool _bHaveSetValue=false;

   for(int _c=tid; _c<_col; _c+=blockDim.x)
   {
   _index1=_c;
   _index2=_c;
   _resultIndex=_r*_pitch_col+_c;
   _bHaveSetValue=false;
   _bRowStart=false;
   _count2=0;

   //跳过还没符合条件的数据
   if(_c<_duration)
   {
   _result[_resultIndex]=0;
   _bHaveSetValue=true;
   continue;
   }

   //检查取到的当日Close是否不为0，为0则返回-1
   _close2=shared1[_index2];
   if(_close2==0)
   {
   _result[_resultIndex]=-1;
   _bHaveSetValue=true;
   continue;
   }

   //循环_duration长度，检测
   for(int _d=0;_d<_duration;_d++)
   {
   if(_index1>0)
   {
   _close2=shared1[_index2];//大的日期

   _index1=_index2-1;
   _close1=shared1[_index1];
   while(_close1==0)//若为0，则往下移动
   {
   if(_index1>0)
   {
   _index1--;
   _close1=shared1[_index1];
   if(_close1>0)
   break;
   }
   else
   {
   _bRowStart=true;
   break;
   }
   }

   //已经到了行的开头
   if(_bRowStart==true && _count2==0)
   {
   _result[_resultIndex]=0;
   _bHaveSetValue=true;
   break;
   }

   if(_bRowStart==true)
   {
   _result[_resultIndex]=1;
   _bHaveSetValue=true;
   break;
   }

   _count2++;
   _close1=shared1[_index1];

   //判断是否异常的逻辑过程
   if(_close2<_close1)
   {
   _result[_resultIndex]=-2;
   _bHaveSetValue=true;
   break;
   }
   _index2=_index1;
   }
   }

   if(_bHaveSetValue==false)
   _result[_resultIndex]=1;

   continue;
   }
   }
}

//2D数组的版本
int main()
{
   //生成数据
   data = (int*)malloc(sizeof(int) * ROW * COL);
   result = (int*)malloc(sizeof(int) * ROW * COL);

   GenerateNumbers(data, DATA_SIZE);

   LARGE_INTEGER litmp;
   LONGLONG Qpart1,Qpart2;
   double dfMinus,dfMillionSecond,dfFreq;
   QueryPerformanceFrequency(&litmp);
   dfFreq = (double)litmp.QuadPart;

   //准备gpu的数据
   int *gpudata, *resultDev;
   int* resultGPU = (int*)malloc(sizeof(int) * ROW * COL);

   //gpu计算
   ///////////////pitch/////////////
   //使用pitch的方式分配内存
   size_t pitch_data, pitch_result;
   cudaMallocPitch((void**) &gpudata, &pitch_data, sizeof(int) * COL, ROW);
   cudaMemset2D(gpudata, pitch_data, 9, COL*sizeof(int), ROW);
   cudaMallocPitch((void**) &resultDev, &pitch_result, sizeof(int) * COL, ROW);
   cudaMemset2D(resultDev, pitch_result, 9, COL*sizeof(int), ROW);
   cudaMemcpy2D(gpudata, pitch_data, data, sizeof(int) * COL, sizeof(int) * COL, ROW, cudaMemcpyHostToDevice);

   //最快-稳定
   QueryPerformanceCounter(&litmp);
   Qpart1 = litmp.QuadPart; //开始计时
   CheckInGPU5<<<ROW, THREAD_NUM, sizeof(int) * COL>>>(ROW,COL,gpudata,resultDev,3,pitch_data/sizeof(int));

   //pitch方式
   cudaMemcpy2D(resultGPU, sizeof(int) * COL, resultDev, pitch_result, sizeof(int) * COL, ROW, cudaMemcpyDeviceToHost);
   //////////////pitch///////////////

   QueryPerformanceCounter(&litmp);
   Qpart2 = litmp.QuadPart; //终止计时
   cudaFree(gpudata);
   cudaFree(resultDev);
   printf("GPU Result(用时:%.2fms):\n",((double)(Qpart2 - Qpart1)/dfFreq)*1000);

   //清理
   delete data;
   data=NULL;
   delete result;
   result=NULL;
   delete resultGPU;
   resultGPU=NULL;

   system("pause");
   return 0;
}
3

你这个程序的话，我感觉你用的shared次数频繁，但是39行你读入数据shared1[i] = _data[row * _pitch_col + i];没有排除bank冲突，所以如果发生bank冲突的话由于你shared次数频繁可能对性能影响很大。

4

那请问版主，该如何改进？？

5

第一步可以用visual profiler找找程序还有多少上升空间：测程序性能对硬件峰值的百分比。具体方法建议看看我们的best practice guide, 或听听GTC上的fundamental optimization & analysis-driven optimization这两个报告。

6

不要意思，当初没仔细看你程序，39行，你的shared mem可以的，不发生bank冲突的。不过，你还是先按照pengwang111的建议，采用visual profile检测一下你的程序性能瓶颈，然后再有目的的优化会更好一些。

7

谢谢斑竹！我在国外的cuda论坛也问了，大家都说要用visual profiler检测，现在正在学习中，请问斑竹，哪里有比较好的visual profiler的入门教程？

8

cudatoolkit里面就有官方的文档,你找一下,应该在doc文件夹下面