▶ 第十二章，在其他语言中使用 OpenCL
● JOCL(Java Building for OpenCL)，PyOpenCL
● 一个 PyOpenCL 的例子代码，需要 pyopencl 包
1 importpyopencl as cl2 importnumpy as np3 importnumpy.linalg as la4
5 a = np.random.rand(50000).astype(np.float32)6 b = np.random.rand(50000).astype(np.float32)7
8 ctx =cl.create_some_context()9 queue =cl.CommandQueue(ctx)10
11 mf =cl.mem_flags12 a_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf =a)13 b_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf =b)14 c_buf =cl.Buffer(ctx, mf.WRITE_ONLY, b.nbytes)15
16 code = """
17 __kernel void vecadd(__global const float *a, __global const float *b, __global float *c)18 {19 int gid=get_global_id(0);20 c[gid] = a[gid] + b[gid];21 }22 """
23 prg =cl.Program(ctx, code).build()24
25 prg.vecadd(queue, a.shape, None, a_buf, b_buf, c_buf)26
27 a_plus_b =np.empty_like(a)28 cl.enqueue_copy(queue, a_plus_b, c_buf)29
30 print(sum(a_plus_b - (a + b)))
● OpenCL 对象的生命周期不是由一个 C 作用域来定义，而是由引用计数操作来定义。
▶ 第十三章，OpenCL 的性能剖析和调试
● 查看事件状态
1 void initializeCL(void)2 {3 cl_int status = 0;4 cl_ulong time;5 cl_event ev;6 cl_platform_id platform;7 cl_device_id device;8 cl_context context;9 cl_command_queue commandQueue;10
11 clGetPlatformIDs(1, &platform, NULL);12 clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);13 context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL);14 commandQueue = clCreateCommandQueue(context, device, CL_QUEUE_PROFILING_ENABLE, &status);//创建队列时传入事件分析属性值，之后无法变更
15
16 ...//关于ev 的一些操作
17
18 clGetEventProfilingInfo(ev, CL_PROFILING_COMMAND_QUEUED, sizeof(cl_ulong), &time, NULL);//获取事件进入命令队列的时刻
19 clGetEventProfilingInfo(ev, CL_PROFILING_COMMAND_SUBMIT, sizeof(cl_ulong), &time, NULL);//获取事件提交到设备的时刻
20 clGetEventProfilingInfo(ev, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &time, NULL); //获取事件开始执行的时刻
21 clGetEventProfilingInfo(ev, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &time, NULL); //获取事件执行完成的时刻
22 }
● 启用 AMD printf 扩展，允许咋个核函数中使用 printf 函数
1 #pragma OPENCL EXTENSION cl_amd_printf: enable
▶ 第十四章，某图像分析应用的性能调优
● Kernel Occupancy，定义为 NWA(一个计算单元上能并发执行的 wavefront 数目)与 NWT(该计算单元能启动的 wavefront 数目的最大值)之比。NWA取决于三个因素：每个 work-item 需要的 GPRS 数目，每个 workgroup 需要的本地内存大小(LDS)和 workgroup 的尺寸。
● LDS 对 KO 的影响：WFLDSmax = min{WGCUmax, WGmax} * WFWG。WFWG为每个 workgroup 的 wavefront 大小(与 LDS 大小没有必然联系)；WGCUmax 为硬件规定一个计算单元中 workgroup 的数量上限，WGmax = LDSCU / LDSWG 为用本地内存计算得到的工作组数量上限，LDSCU 为本地内存大小， LDSWG 为一个 workgroup 需要的 LDS 的大小。
● GPR 对 KO 的影响：WFGPRmax = floor(WFGPR / WFWG) * WFWG。WFGPR = Nregmax / Nregused，为每个 work-item 使用的寄存器的最大数量和。这里相当于把 WFGPR 乡下四舍五入到 的整数倍 WFWG。
● 如果一个 kernel 的性能受 GPRS 的限制，而不受 LDS 的限制，则将部分数据转入 LDS 中会部分优化性能，过多的转移数据会进入全局内存中，严重影响执行速度