Visible to Intel only — GUID: GUID-D4B1A2A6-C2EB-433E-8033-D3DDD8511DD5
Legal Information
Getting Help and Support
Introduction
Coding for the Intel® Processor Graphics
Platform-Level Considerations
Application-Level Optimizations
Optimizing OpenCL™ Usage with Intel® Processor Graphics
Check-list for OpenCL™ Optimizations
Performance Debugging
Using Multiple OpenCL™ Devices
Coding for the Intel® CPU OpenCL™ Device
OpenCL™ Kernel Development for Intel® CPU OpenCL™ device
Mapping Memory Objects
Using Buffers and Images Appropriately
Using Floating Point for Calculations
Using Compiler Options for Optimizations
Using Built-In Functions
Loading and Storing Data in Greatest Chunks
Applying Shared Local Memory
Using Specialization in Branching
Considering native_ and half_ Versions of Math Built-Ins
Using the Restrict Qualifier for Kernel Arguments
Avoiding Handling Edge Conditions in Kernels
Using Shared Context for Multiple OpenCL™ Devices
Sharing Resources Efficiently
Synchronization Caveats
Writing to a Shared Resource
Partitioning the Work
Keeping Kernel Sources the Same
Basic Frequency Considerations
Eliminating Device Starvation
Limitations of Shared Context with Respect to Extensions
Why Optimizing Kernel Code Is Important?
Avoid Spurious Operations in Kernel Code
Perform Initialization in a Separate Task
Use Preprocessor for Constants
Use Signed Integer Data Types
Use Row-Wise Data Accesses
Tips for Auto-Vectorization
Local Memory Usage
Avoid Extracting Vector Components
Task-Parallel Programming Model Hints
Visible to Intel only — GUID: GUID-D4B1A2A6-C2EB-433E-8033-D3DDD8511DD5
Using Vector Data Types
To maximize CPU vector unit utilization, try to use vector data types in your kernel code. This technique enables you to map vector data types directly to the hardware vector registers. Thus, the data types used should match the width of the underlying SIMD instructions.
Consider the following recommendations:
- On the 2nd Generation Intel® Core™ Processors and higher with Intel® AVX support, use data types such as float8 or double4, so you bind code to the specific register width of the underlying hardware. This method provides maximum performance on a specific platform. However, performance on other platforms and supported Intel processors might be less than optimal.
- You may use wider data types, such as float16, to transparently cover many SIMD hardware register widths. However, using types wider than the underlying hardware is similar to loop unrolling. This method might improve performance in some cases, but also increases register pressure. Still consider using uchar16 data type to process four pixels simultaneously when operating on eight-bit-per-component pixels.
- When manually “vectorizing” an original kernel that uses scalar data types (like float) to use vector data types (like float8) instead, remember that each work-item processes N elements (for float/float8 example). Make sure you reduce the global size accordingly, so it is dividable by N.
- The int8 data type improves performance for the 4th Generation Intel® Core™ processors and higher.
Using this coding technique, you plan the vector-level parallelism yourself instead of relying on the implicit vectorization module (see the "Benefitting from Implicit Vectorization" section). This approach is useful in the following scenarios:
- You are porting code originally used Intel SSE/AVX/AVX2 instructions.
- You want to benefit from hand-tuned vectorization of your code.
The following example shows a multiplication kernel that targets the 256-bit vector units of the 2nd Generation Intel Core Processors:
__kernel __attribute__((vec_type_hint(float8)))
void edp_mul(__constant float8 *a,
__constant float8 *b,
__global float8 *result)
{
int id = get_global_id(0);
result[id] = a[id]* b[id];
}
In this example, the data passed to the kernel represents buffers of float8. The calculations are performed on eight elements together.
The attribute added before the kernel, signals the compiler, or the implementation that this kernel has an optimized vectorized form, so the implicit vectorization module does not operate on it. Use vec_type_hint to indicate to the compiler that your kernel already processes data using mostly vector types. For more details on this attribute, see the OpenCL™ 1.2 Specification.
See Also
Benefitting from Implicit Vectorization
The OpenCL™ 1.2 Specification at http://www.khronos.org/registry/cl/specs/opencl-1.2.pdf
Did you find the information on this page useful?
Feedback Message
Characters remaining: