初学cuda程序的一点记录

编cuda程序比较少，尤其是性能分析器也不会看，过段时间还要找工作，把今天看的调的一点点东西记下来，也不一定对，先记录着。程序优化后面再说，就是个naive版代码和分析。

kernel函数分成两个，在主机端调用两次

//按列求和，得到width个结果

__global__ void colAdd(unsigned char* image, int width, int height, int leftUpX, int leftUpY, int x, int y, int* colSum){int tidX = threadIdx.x;int sum = 0;if (tidX < x){for (int i = 0; i < y; i++){sum += imageCon[__mul24((i + leftUpY), width) + leftUpY + tidX];}//printf("imageCon[%d]=%d\n", tidX, imageCon[tidX]);colSum[tidX] = sum;}}

//矢量求和，简单求和，不用多cuda core__global__ void vecAdd(int *vec, int n, int *res){//int n = 256;int sum = 0;for (int i = 0; i < n; i++){sum += vec[i];}*res = sum;}

结果如下

使用的blocks=1， threads=256可以看到活跃的warps太少，活跃的blocks也太少，因此调整blocks数量为（4,64）， （16,16）， （64,4），（256,1），发现（64,4）是效率较高

活跃warps数量增加了，理论活跃blocks也增加了。耗时减少为原来的40%左右。

后分析主机端调用程序需要设置参数，发射指令，也需要耗时，依据参数数量不同耗时也不用，这里分别为3us和6us，而launch时间达到了15us以上，相对于这里cuda kernel函数的执行来说也算是比较影响体验了

从时间轴上看执行要比lauch要晚很多，不知道为什么，这块还有很多不懂，不太会看这些，以后要找专门讲这东西的资料看看。

为减少launch和参数传递延迟，将两个函数合并如下

//按列求和，得到width个结果__global__ void colVecAdd(unsigned char* image, int width, int height, int leftUpX, int leftUpY, int x, int y, int* colSum,int *dev_average){int tidX = threadIdx.x + blockIdx.x * 4;int sum=0;if (tidX < x){for (int i = 0; i < y; i++){ sum+= imageCon[__mul24((i + leftUpY), width) + leftUpY + tidX];}//printf("imageCon[%d]=%d\n", tidX, imageCon[tidX]);colSum[tidX] = sum;}__syncthreads();if (tidX < 1){int sum = 0;for (int i = 0; i < width; i++){sum += colSum[i];}*dev_average = sum;}}

结果变成，launch和传参总耗时减半，cuda  kernel运行时间也减少了

运行时间进一步缩短。

后面可以考虑将该程序使用device端调用，尝试减少耗时。