这些cuda资料要学会总结

CUDA并行运算，block ，thread计算测试【续-优化算法】

  (2014-03-25 20:01:34)

转载▼

星期一给导师发邮件的时候，总结一周工作。

脑袋里面又想起了前面做测试的6种加法算法。

http://blog.sina.com.cn/s/blog_98740ded0101jkgo.html

考虑这里面的第4种算法相当不错，但是觉得有点不太完美。

第四种方法是使用单block，多线程的方法，kernel运行完毕后，经过cudaDeviceSynchronize();

需要将256个thread得到的加法结果copy到cpu这边来，然后cpu对它们做最终的merge。

感觉不太好，这里面存在的问题，有两点：

1.copy到cpu，数据需要经过pci e的总线，(虽然256个数据不是很多，

但是总归觉得不好)

2.cpu做256个数据的加法，总归没有gpu快，(虽然256个数据不是很多，

但是总归觉得不好)

然后就，灵感来了O(∩_∩)O哈哈哈~。。

汇总操作也放到kernel里面，但是，由某一个线程来执行汇总，岂不更妙！！！

说干就干了，原先的方法四先贴出来。

//方法4/

__global__ static void sum_ThreadOptimization(int *num, int* result)

{

const int tid = threadIdx.x;

int sum = 0;

int i;

for(i = tid; i < DATA_SIZE; i += THREAD_NUM) {

sum += num[i] + num[i] + num[i] + num[i] + num[i] + num[i]+ num[i] + num[i] + num[i] + num[i] + num[i] + num[i];

}

result[tid] = sum;

}

extern "C" void sum_ThreadOptimization_mode4(const int *num, int* result,float *time){

StopWatchInterface *timer = NULL;

int *gpudata;

int *gpuresult;

int *cpuPartialResult;

sdkCreateTimer(&timer);

sdkResetTimer(&timer);

sdkStartTimer(&timer);

//为输入数据分配gpu显存空间。

cudaMalloc((void**) &gpudata, sizeof(int) * DATA_SIZE);

//将内存中的数据拷贝到显存中去。

cudaMemcpy(gpudata, num, sizeof(int) * DATA_SIZE,

cudaMemcpyHostToDevice);

//为输出数据分配显存空间,只需要长度为THREAD_NUM的地址

cudaMalloc((void**) &gpuresult, sizeof(int)*THREAD_NUM);

//执行运算

sum_ThreadOptimization<<<1, THREAD_NUM, 0>>>(gpudata, gpuresult);

//等待计算执行完毕

cudaDeviceSynchronize();

//申请cpu存储空间，存多个线程结果

cpuPartialResult=(int *)malloc(sizeof(int)*THREAD_NUM);

//运算完毕后，将数据回收到cpu，

cudaMemcpy(cpuPartialResult, gpuresult, sizeof(int) * THREAD_NUM,

cudaMemcpyDeviceToHost);

//结果汇总，

for(int i = 0; i < THREAD_NUM; i++) {

*result += cpuPartialResult[i]; 

}

//终止定时器

sdkStopTimer(&timer);

//计算得到定时器的时间

*time =   (sdkGetTimerValue(&timer));

//删除定时器

sdkDeleteTimer(&timer);

//释放gpu的空间

cudaFree(gpudata);

cudaFree(gpuresult);

//释放cpu空间

free(cpuPartialResult);

}

经过前面这样的考虑，思路就清晰了，修改为方法7。主要改动内容为：

1.kernel里面使用share memory(这样所有的线程都能访问了) 存储各个

thread的加法结果。

2.使用 __syncthreads();线程同步语句，保证所有线程执行完毕，数据都准备

好了，再做汇总。

3.只需要将一个最终的结果copy到cpu就可以了。

4.cudaDeviceSynchronize();似乎可以不再使用了，我试过不用这个，也可以得

到最终的结果。

//方法7/

__global__ static void sum_ThreadOptimization_sync(int *num, int* result)

{

//一个block中共享shared memory

__shared__ int Result[THREAD_NUM];

const int tid = threadIdx.x;

int *partResult=Result + threadIdx.x;

int i;

*partResult=0;

for(i = tid; i < DATA_SIZE; i += THREAD_NUM) {

*partResult += num[i] + num[i] + num[i] + num[i] + num[i] + num[i]+ num[i] + num[i] + num[i] + num[i] + num[i] + num[i];

}

__syncthreads();

 

//同步这一个block中的所有线程

//结果汇总，

if (tid==0)

{

//这个地方是硬伤啊，悲剧，不能放在if 外面了，否则这么多thread同时对result赋值，就跪了

*result=0;

for( i = 0; i < THREAD_NUM; i++) {

*result += Result[i]; 

}

}

}

extern "C" void sum_ThreadOptimization_sync_mode7(const int *num, int* result,float *time){

StopWatchInterface *timer = NULL;

int *gpudata;

int *gpuresult;

sdkCreateTimer(&timer);

sdkResetTimer(&timer);

sdkStartTimer(&timer);

//为输入数据分配gpu显存空间。

cudaMalloc((void**) &gpudata, sizeof(int) * DATA_SIZE);

//将内存中的数据拷贝到显存中去。

cudaMemcpy(gpudata, num, sizeof(int) * DATA_SIZE,

cudaMemcpyHostToDevice);

//为输出数据分配显存空间,只需要长度为1的地址

cudaMalloc((void**) &gpuresult, sizeof(int));

//执行运算

sum_ThreadOptimization_sync<<<1, THREAD_NUM, 0>>>(gpudata, gpuresult);

//等待计算执行完毕

// cudaDeviceSynchronize();

 

//运算完毕后，将数据回收到cpu，

cudaMemcpy(result, gpuresult, sizeof(int),

cudaMemcpyDeviceToHost);

//终止定时器

sdkStopTimer(&timer);

//计算得到定时器的时间

*time =   (sdkGetTimerValue(&timer));

//删除定时器

sdkDeleteTimer(&timer);

//释放gpu的空间

cudaFree(gpudata);

cudaFree(gpuresult);

}