This is the part 2 of the OpenCL Webinar Series put out by Intel. There were some good information about optimization in general but lot of the information focused on how to optimize OpenCL running on 3rd Gen Intel Core Processor. I jot down some optimization notes that are applicable to all processors.
Avoid Invariant in OpenCL Kernel
Anything that's independent from kernel execution (invariants, constants), move it to the host.
Bad:
__kernel void test1 (__global int* data, int2 size, int base) { int offset = size.y * base + size.x; float offsetF = (float)offset; ... }Good:
__kernel void test1 (__global int* data, int offset, float offset) { ... }Avoid initialization in Kernel
Move one time initialization to the host or in a different kernel.
Bad: __kernel void something (__global int* data) { size_t tid = get_global_id(0) if (0== tid) { //Do Something } barrier(CLK_GLOBAL_MEM_FENCE) }
Use the Built in Functions
i.e. dot, hypot, clamp, etc.
Trade Accuracy vs Speed
If the output is correct, look at the performance and use "mad" or "native_sin"
Get rid of edge conditions
Bad:
__kernel void myKernel(__global int* data, int maxR, int maxC){ int row = get_global_id(0); int col = get_global_id(1); if (row > maxRow){ ... } else if (col > maxCol) { ... } else{ ... } }Use the ND Range (work within a smaller range. Just avoid the conditionals...)
Use the padded buffers...
Get rid of edge conditions in general
Use logical ops instead of comparison/conditions.
Reduce number of registers to increase parallelism
Avoid Byte/Short Load and Stores
Use load and store in a greater chunk.
i.e. uchar -> uint4
Don't use too many barriers in kernel
You are asking everything to wait until the work items are done.
If the work item doesn't run, the code may hang.
Use the vector types (float8, double4, int4, etc) for better performance on the CPU
Preferred work group size for kernels is 64 or 128 work items.
-Workgroup size multiple of 8
-Query
CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE parameter by calling to clGetKernelWorkgroupInfo.
-Match number of workgroups to logical cores.
Buffer Object vs. Image Objects
CPU Device:
- Avoid using Image Object and use Buffer object instead.