▶ 参考书中的代码，写了
● 代码，核函数文件包含三中算法
// kernel.cl
__kernel void bitonicSort01(__global uint *data, const uint stage, const uint subStage, const uint direction)// 基本的元素对调整
{
const uint gid = get_global_id();
const uint isAscend = ((gid / ( << stage)) % ) ? - direction : direction; // 判断本工作项的元素对应该排成升序还是降序
const uint distance = << (stage - subStage); // 元素对下标差
const uint lid = (gid / distance) * distance * + gid % distance; // 寻找元素对的左右元素
const uint rid = lid + distance;
const uint lElement = data[lid], rElement = data[rid];
if (lElement > rElement && isAscend || lElement < rElement && !isAscend) // 不符合排序要求，交换两元素
data[lid] = rElement, data[rid] = lElement;
} __kernel void bitonicSort02(__global uint *data, const uint stage, const uint subStage, const uint direction, __local uint *localMem)// 使用局部内存调整
{
const uint gid = get_global_id(), mid = get_local_id(); // 同 bitonicSort01
const uint isAscend = ((gid / ( << stage)) % ) ? - direction : direction;
const uint distance = << (stage - subStage);
const uint lid = (gid / distance) * distance * + gid % distance;
const uint rid = lid + distance; localMem[mid * + ] = data[lid]; // 读取 data 的时候读进局部内存，与局部内存相关的下标用的都是 mid 而不是 gid
localMem[mid * + ] = data[rid];
barrier(CLK_LOCAL_MEM_FENCE); if (localMem[mid * + ] > localMem[mid * + ] && isAscend || localMem[mid * + ] < localMem[mid * + ] && !isAscend)
data[lid] = localMem[mid * + ], data[rid] = localMem[mid * + ];
} #define STRIDE 4 // aux 中四个元素一组，表示一个工作项的元素对索引和值，依照 main.c 中给定的第 5 参数的大小进行相等的调整 __kernel void bitonicSort03(__global uint *data, const uint stage, const uint subStage, const uint direction, __local uint *localMem, __local uint *aux)// 使用两个局部内存，感觉多此一举？
{
const uint gid = get_global_id(), mid = get_local_id(); // 同 bitonicSort02
const uint isAscend = ((gid / ( << stage)) % ) ? - direction : direction;
const uint distance = << (stage - subStage);
const uint lid = (gid / distance) * distance * + gid % distance;
const uint rid = lid + distance; localMem[mid * + ] = data[lid];
localMem[mid * + ] = data[rid];
barrier(CLK_LOCAL_MEM_FENCE); aux[mid * STRIDE + ] = lid; // 开始向aux 中填充
aux[mid * STRIDE + ] = rid;
if (localMem[mid * + ] > localMem[mid * + ] && isAscend || localMem[mid * + ] < localMem[mid * + ] && !isAscend)
aux[mid * STRIDE + ] = localMem[mid * + ], aux[mid * STRIDE + ] = localMem[mid * + ];
else
aux[mid * STRIDE + ] = localMem[mid * + ], aux[mid * STRIDE + ] = localMem[mid * + ];
barrier(CLK_LOCAL_MEM_FENCE); data[aux[mid * STRIDE + ]] = aux[mid * STRIDE + ], data[aux[mid * STRIDE + ]] = aux[mid * STRIDE + ]; // 向原数组中填充数据
/*// 书中的填充方法，一个工作组仅使用一个工作项串行地向原数组中填充，美名曰“无冲突”，实际上花费了 5 倍的时间
if (mid == 0)
{
for (int i = 0; i < get_local_size(0); i++)
data[aux[i * STRIDE + 0]] = aux[i * STRIDE + 1], data[aux[i * STRIDE + 2]] = aux[i * STRIDE + 3];
}
*/
}
// main.c
#include <stdio.h>
#include <stdlib.h>
#include <cl.h> //#define PRINT_RESULT // 输出排序前后的数组元素（数据量较大时不用）
#define ASCENDING 1 // 升序
#define DESCENDING 0 // 降序
#define DATA_SIZE (1<<20) // 数据规模
#define GROUP_SIZE 128 // 工作组大小 const char *sourceText = "D:/Code/OpenCL/OpenCLProjectTemp/OpenCLProjectTemp/kernel.cl";
const unsigned int sortOrder = ASCENDING; // ASCENDING / DESCENDING int readText(const char* kernelPath, char **pcode)// 读取文本文件放入 pcode，返回字符串长度
{
FILE *fp;
int size;
//printf("<readText> File: %s\n", kernelPath);
fopen_s(&fp, kernelPath, "rb");
if (!fp)
{
printf("Open kernel file failed\n");
getchar();
exit(-);
}
if (fseek(fp, , SEEK_END) != )
{
printf("Seek end of file failed\n");
getchar();
exit(-);
}
if ((size = ftell(fp)) < )
{
printf("Get file position failed\n");
getchar();
exit(-);
}
rewind(fp);
if ((*pcode = (char *)malloc(size + )) == NULL)
{
printf("Allocate space failed\n");
getchar();
exit(-);
}
fread(*pcode, , size, fp);
(*pcode)[size] = '\0';
fclose(fp);
return size + ;
} int main()
{
int i;
srand(); cl_int status;
cl_uint nPlatform;
clGetPlatformIDs(, NULL, &nPlatform);
cl_platform_id *listPlatform = (cl_platform_id*)malloc(nPlatform * sizeof(cl_platform_id));
clGetPlatformIDs(nPlatform, listPlatform, NULL);
cl_uint nDevice;
clGetDeviceIDs(listPlatform[], CL_DEVICE_TYPE_ALL, , NULL, &nDevice);
cl_device_id *listDevice = (cl_device_id*)malloc(nDevice * sizeof(cl_device_id));
clGetDeviceIDs(listPlatform[], CL_DEVICE_TYPE_ALL, nDevice, listDevice, NULL);
cl_context context = clCreateContext(NULL, nDevice, listDevice, NULL, NULL, &status);
cl_command_queue queue = clCreateCommandQueue(context, listDevice[], CL_QUEUE_PROFILING_ENABLE, &status); int *hostA = (cl_int*)malloc(DATA_SIZE * sizeof(cl_int));
for (i = ; i < DATA_SIZE; hostA[i++] = rand());
#ifdef PRINT_RESULT // 输出排序前的数组
printf("\n");
for (i = ; i < DATA_SIZE; i++)
{
printf("%5d,", hostA[i]);
if ((i + ) % == )
printf("\n");
}
printf("\n");
#endif
cl_mem deviceA = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, DATA_SIZE * sizeof(cl_int), hostA, &status); char *code;
size_t codeLength = readText(sourceText, &code);
cl_program program = clCreateProgramWithSource(context, , (const char**)&code, &codeLength, &status);
status = clBuildProgram(program, nDevice, listDevice, NULL, NULL, NULL);
if (status)
{
char info[];
clGetProgramBuildInfo(program, listDevice[], CL_PROGRAM_BUILD_LOG, , info, NULL);
printf("\n%s\n", info);
}
//cl_kernel kernel = clCreateKernel(program, "bitonicSort01", &status); // 选择使用的核函数
//cl_kernel kernel = clCreateKernel(program, "bitonicSort02", &status);
cl_kernel kernel = clCreateKernel(program, "bitonicSort03", &status); size_t globalSize = DATA_SIZE / , groupSize = GROUP_SIZE;
cl_uint stageCount, stage, subStage;
for (i = DATA_SIZE, stageCount = ; i > ; stageCount++, i >>= ); // 需要的总轮数
clSetKernelArg(kernel, , sizeof(cl_mem), (void*)&deviceA);
clSetKernelArg(kernel, , sizeof(cl_uint), (void*)&sortOrder);
clSetKernelArg(kernel, , GROUP_SIZE * * sizeof(cl_uint), NULL); // bitonicSort02 和 bitonicSort03 需要的局部内存参数
clSetKernelArg(kernel, , GROUP_SIZE * * sizeof(cl_uint), NULL); // bitonicSort03 需要的局部内存参数 cl_event exeEvent; // 计时用的事件
cl_ulong executionStart, executionEnd, timeCount1, timeCount2; // 计时器， 精确到 ns
for (stage = timeCount1 = timeCount2 = ; stage < stageCount; stage++) // 当前轮数
{
clSetKernelArg(kernel, , sizeof(cl_uint), (void*)&stage);
for (subStage = ; subStage < stage + ; subStage++) // 当前轮中归并的步骤数
{
clSetKernelArg(kernel, , sizeof(cl_uint), (void*)&subStage);
//clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &globalSize, &groupSize, 0, NULL, NULL); // 不计时的核函数调用
clEnqueueNDRangeKernel(queue, kernel, , NULL, &globalSize, &groupSize, , NULL, &exeEvent);
clWaitForEvents(, &exeEvent);
clGetEventProfilingInfo(exeEvent, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &executionStart, NULL);
clGetEventProfilingInfo(exeEvent, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &executionEnd, NULL);
timeCount1 += (executionEnd - executionStart) / , timeCount2 += (executionEnd - executionStart) % ;// ms 的整数和小数部分分开记录
}
}
printf("Kernel time: %lu.%lu ms\n", timeCount1 + timeCount2 / , timeCount2 % ); // 内核运行总时间 clEnqueueReadBuffer(queue, deviceA, CL_TRUE, , DATA_SIZE * sizeof(cl_int), hostA, , NULL, NULL);
#ifdef PRINT_RESULT // 输出排序后的结果
printf("\n");
for (i = ; i < DATA_SIZE; i++)
{
printf("%5d,", hostA[i]);
if ((i + ) % == )
printf("\n");
}
printf("\n");
#else
for (i = ; i < DATA_SIZE - ; i++) // 不输出结果，只做检查，有错误的时候输出第一个错误的位置
{
if (sortOrder != (hostA[i + ] >= hostA[i]))
{
printf("Error at i == %d, hostA[i] == %d, hostA[i+1] == %d", i, hostA[i], hostA[i + ]);
break;
}
}
#endif free(listPlatform);
free(listDevice);
clReleaseContext(context);
clReleaseCommandQueue(queue);
clReleaseProgram(program);
clReleaseKernel(kernel);
clReleaseEvent(exeEvent);
free(hostA);
clReleaseMemObject(deviceA);
getchar();
return ;
}
● 输出结果，统一采用 (1<<20) 的数据规模，尝试不同的工作组大小。使用局部内存并没有明显提升，尤其是使用两个局部内存的方法，严重拖后腿。
bitonicSort01
// groupSize = 64
kernels time: 7.941984 ms
// groupSize = 128
kernels time: 7.947360 ms
// groupSize = 256
kernels time: 7.935616 ms
// groupSize = 512
kernels time: 7.919776 ms
// groupSize = 1024
kernels time: 7.970080 ms bitonicSort02
// groupSize = 64
kernels time: 7.980160 ms
// groupSize = 128
kernels time: 7.957984 ms
// groupSize = 256
kernels time: 7.964032 ms
// groupSize = 512
kernels time: 7.959072 ms
// groupSize = 1024
kernels time: 8.517120 ms bitonicSort03
// groupSize = 64
kernels time: 9.901120 ms
// groupSize = 128
kernels time: 9.936384 ms
// groupSize = 256
kernels time: 9.950976 ms
// groupSize = 512
kernels time: 10.134176 ms
// groupSize = 1024
kernels time: 10.533516 ms bitonicSort03（书中的串行写入）
// groupSize = 64
kernels time: 51.590080 ms
// groupSize = 128
kernels time: 31.607904 ms
// groupSize = 256
kernels time: 35.344244 ms
// groupSize = 512
kernels time: 50.841184 ms
// groupSize = 1024
kernels time: 93.284544 ms
● 总结
■ CPU 版双调排序使用递归，代码比较简洁，也可以使用本篇中的方法家拿其转化为循环迭代。GPU暂时不使用递归，主要精力在于不停地向命令队列中入队不同层次的任务，同一层次内的任务可以并行执行
■ 这本书就是垃圾，就 T M D 是 垃 圾！①书里代码引用不全，比如核函数里上来就使用变量 threadId，声明定义都没有；②代码混乱，比如在 bitonicSort01 和 bitonicSort02 里 threadId 的含义是不同的，尤其书中 bitonicSort02 里同时用该变量表示 get_global_id(0) 和 get_local_id(0)，一想就知道代码执行结果肯定是错的；③ 机翻，好不容易写到倒数第二章了，翻译都懒得弄了，这里抄一段原文看谁看得懂：“我们基本上执行的操作是引入名为 sharedMem 的变量，并且使用加载这些值的简单策略：每个线程将在共享内存数据存储器中存储两个值（邻近值），在随后的代码部分中读出此线程，并且用于引用全局内存的所有读取操作目前在本地 / 共享内存中执行”。④ 残缺不堪，感觉 bitonicSort03 的优化应该是有一定道理的，但是在这里完全没有看出该优化的目的和优势所在，书中也没有说清楚。