学校建设网站前的市场分析,博客 系统 wordpress,网站用途,网站集约化建设讲话在OpenCL标准中#xff0c;没有给出查看计算设备一共有多少寄存器#xff0c;至少能分配给每个work-item多少寄存器使用的特征查询。而由于一个段内核代码是否因寄存器紧缺而导致性能严重下降也是一个比较重要的因素#xff0c;因此我这边提供一个比较基本的方法来猜测当前计…在OpenCL标准中没有给出查看计算设备一共有多少寄存器至少能分配给每个work-item多少寄存器使用的特征查询。而由于一个段内核代码是否因寄存器紧缺而导致性能严重下降也是一个比较重要的因素因此我这边提供一个比较基本的方法来猜测当前计算设备至少能为每个work-item分配多少可用的寄存器。
这个方法的思路是先定义四个临时变量然后在一个大规模循环里面做一定规模的计算。然后把时间统计出来。随后再定义八个临时变量仍然在与前者相同次数的循环里做一定规模的计算再把时间统计出来。一般如果寄存器不爆或者由于Cache的缘故性能影响不大的话两者消耗时间一般在2倍左右。如果后者比前者超了2.2倍以上那么我们即可认为寄存器爆了
这个方法对于一般的GPU更有用些。由于CPU往往拥有L1 Data Cache当寄存器不够用的时候编译器会将不太常用的数据放到栈中而栈在此时往往能获得高命中率的Cache访问因此性能不会过受影响。而GPU端当寄存器不够用时编译器往往会采取将不常用数据直接存放到VRAM中而对外部VRAM的访问往往是比较慢的因此如果临时变量太多使得频繁访问外部存储器会使得整体计算性能大幅下降。当然现在不少GPU也有了L1 Cache但是空间也十分有限。因此这里用“猜”这个词
下面先提供四个临时变量的kernel代码
kernel void QueryRegisterCount(__global int *pInOut)
{int index get_global_id(0);int i0 pInOut[(index * 4 0) * 4];int i1 pInOut[(index * 4 1) * 4];int i2 pInOut[(index * 4 2) * 4];int i3 pInOut[(index * 4 3) * 4];for(int i 0; i 100000; i){i1 i0 1;i2 i1 1;i3 i2 1;i0 i3 1;i1 i0 1;i2 i1 1;i3 i2 1;i0 i3 1;i1 i0 2;i2 i1 2;i3 i2 2;i0 i3 2;i1 i0 3;i2 i1 3;i3 i2 3;i0 i3 3;}pInOut[(index * 4 0) * 4] i0;pInOut[(index * 4 1) * 4] i1;pInOut[(index * 4 2) * 4] i2;pInOut[(index * 4 3) * 4] i3;
}再提供八个临时变量的kernel代码
kernel void QueryRegisterCount(__global int *pInOut)
{int index get_global_id(0);int i0 pInOut[(index * 8 0) * 4];int i1 pInOut[(index * 8 1) * 4];int i2 pInOut[(index * 8 2) * 4];int i3 pInOut[(index * 8 3) * 4];int i4 pInOut[(index * 8 4) * 4];int i5 pInOut[(index * 8 5) * 4];int i6 pInOut[(index * 8 6) * 4];int i7 pInOut[(index * 8 7) * 4];for(int i 0; i 100000; i){i1 i0 1;i2 i1 1;i3 i2 1;i4 i3 1;i5 i4 1;i6 i5 1;i7 i6 1;i0 i7 1;i1 i0 1;i2 i1 1;i3 i2 1;i4 i3 1;i5 i4 1;i6 i5 1;i7 i6 1;i0 i7 1;i1 i0 2;i2 i1 2;i3 i2 2;i4 i3 2;i5 i4 2;i6 i5 2;i7 i6 2;i0 i7 2;i1 i0 3;i2 i1 3;i3 i2 3;i4 i3 3;i5 i4 3;i6 i5 3;i7 i6 3;i0 i7 3;}pInOut[(index * 8 0) * 4] i0;pInOut[(index * 8 1) * 4] i1;pInOut[(index * 8 2) * 4] i2;pInOut[(index * 8 3) * 4] i3;pInOut[(index * 8 4) * 4] i4;pInOut[(index * 8 5) * 4] i5;pInOut[(index * 8 6) * 4] i6;pInOut[(index * 8 7) * 4] i7;
}而像16个、32个临时变量的方法依此类推
然后给出主机端代码
/** Prepare for running an OpenCL kernel program to get register count *//*Step 4: Creating command queue associate with the context.*/commandQueue clCreateCommandQueue(context, device, CL_QUEUE_PROFILING_ENABLE, NULL);/*Step 5: Create program object */// Read the kernel code to the bufferkernelPath [[NSBundle mainBundle] pathForResource:reg ofType:ocl];aSource [[NSString stringWithContentsOfFile:kernelPath encoding:NSUTF8StringEncoding error:nil] UTF8String];kernelLength strlen(aSource);program clCreateProgramWithSource(context, 1, aSource, kernelLength, NULL);/*Step 6: Build program. */status clBuildProgram(program, 1, device, NULL, NULL, NULL);/*Step 7: Initial inputs and output for the host and create memory objects for the kernel*/const size_t memSize global_work_size[0] * 1024 * 4 * 4;cl_int *orgBufer (cl_int*)malloc(memSize);memset(orgBufer, 1, memSize);outputMemObj clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR, memSize, orgBufer, NULL);/*Step 8: Create kernel object */kernel clCreateKernel(program, QueryRegisterCount, NULL);/*Step 9: Sets Kernel arguments.*/status | clSetKernelArg(kernel, 0, sizeof(outputMemObj), outputMemObj);/*Step 10: Running the kernel.*/for(int i 0; i 5; i){NSTimeInterval beginTime [[NSProcessInfo processInfo] systemUptime];status | clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);clFinish(commandQueue);NSTimeInterval endTime [[NSProcessInfo processInfo] systemUptime];NSLog(Time spent: %f, endTime - beginTime);}free(orgBufer);if(status ! CL_SUCCESS){NSLog(Program built failed!);return;}clReleaseMemObject(outputMemObj);clReleaseProgram(program);clReleaseKernel(kernel);clReleaseCommandQueue(commandQueue);clReleaseContext(context);以上由于是在macOS下开发的因此直接用Objective-C文件读写更方便些。但是大部分都是C代码很容易读懂。
其中最后一断代码中我们做5次循环统计时间。我们比较的时候往往选出5次执行时间中最小耗费的时间进行比较。
在2013年的MacBook Air中的Intel HD 5000中的测试结果为
四个临时变量耗费0.061020秒八个临时变量耗费0.121868秒十六个临时变量耗费0.243470秒三十二个临时变量耗费0.719506秒
很显然我们可以猜得Intel HD Graphics 5000至少可以为每个work-item分配16个寄存器。
我们如果要用在实际应用场合可以通过动态生成kernel字符串依次执行进行检测直到相邻两段kernel的执行时间超过2.2倍那么我们即可终止。
