标签: CUDAExample
例程说明:
例程的每一个Block进行归并排序查找处最小值。由于BLocks间并行执行时没有同步机制,实现过程中例程会测量每一个Block所耗费的时间,同时实践过程使用了动态分配share memory的技巧,关于share memory使用参考另一篇博客《CUDA核函数share memory》。
归并排序查抄最小值方法,参考下图:
数组中有2n个数据,每一个Block开辟n个线程(n小于block线程的最大开辟线程数),每一个线程进行交换,挑选出最小值,进行二分,再一次挑选,直到剩下最后一个值就是需要的最小值,同样适用于查找最大值。实现过程中,块内部同步函数不可缺少。
实现代码核函数
// Copy input.
shared[tid] = input[tid];
shared[tid + blockDim.x] = input[tid + blockDim.x];
// Perform reduction to find minimum.
for (int d = blockDim.x; d > 0; d /= 2)
{
__syncthreads();
if (tid < d)
{
float f0 = shared[tid];
float f1 = shared[tid + d];
if (f1 < f0)
{
shared[tid] = f1;
}
}
}
// Write result.
if (tid == 0) output[bid] = shared[0];
__syncthreads();
例程中实现了保存Block执行的时间,并将开始时刻和结束时刻存储在显存中。核函数中使用的计时函数为:clock
if (tid == 0) timer[bid] = clock();存储开始时刻
if (tid == 0) timer[bid+gridDim.x] = clock();存储结束时刻
CPU 对计时数据处理过程:
clock_t timePerBlock;
for (int i = 1; i < NUM_BLOCKS; i++)
{
timePerBlock = timer[NUM_BLOCKS + i] - timer[i];
printf("Block %d clocks = %d\n", i,(int)(timePerBlock));
}
// Compute the difference between the last block end and the first block start.
clock_t minStart = timer[0];
clock_t maxEnd = timer[NUM_BLOCKS];
for (int i = 1; i < NUM_BLOCKS; i++)
{
minStart = timer[i] < minStart ? timer[i] : minStart;
maxEnd = timer[NUM_BLOCKS+i] > maxEnd ? timer[NUM_BLOCKS+i] : maxEnd;
}
printf("\n minStart clocks = %d\n", (int)(minStart));
printf(" maxEnd clocks = %d\n", (int)(maxEnd));
printf("\n Total clocks = %d\n", (int)(maxEnd - minStart)); // CLOCKS_PER_SEC
end