《OpenCL异构并行编程实战》补充笔记散点,第十二至十四章

Posted cuancuancuanhao

tags:

篇首语:本文由小常识网(cha138.com)小编为大家整理,主要介绍了《OpenCL异构并行编程实战》补充笔记散点,第十二至十四章相关的知识,希望对你有一定的参考价值。

? 第十二章,在其他语言中使用 OpenCL

● JOCL(Java Building for OpenCL),PyOpenCL

● 一个 PyOpenCL 的例子代码,需要 pyopencl 包

 1 import pyopencl as cl
 2 import numpy as np
 3 import numpy.linalg as la
 4 
 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_flags
12 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 中会部分优化性能,过多的转移数据会进入全局内存中,严重影响执行速度

以上是关于《OpenCL异构并行编程实战》补充笔记散点,第十二至十四章的主要内容,如果未能解决你的问题,请参考以下文章

OpenCL异构并行计算

OpenCL

C++实战之OpenCL 并行优化编程从零学起系列文章

OpenCL现在都支持哪些并行设备

C++ 实战之OpenCL环境搭建

从零开始学习OpenCL开发架构