vec_type_hint
Description
The optional __attribute__((vec_type_hint(<type>)))
is part of the OpenCL Language Specification, and is a hint to the OpenCL compiler
representing the computational width of the kernel, providing a basis for calculating
processor bandwidth utilization when the compiler is looking to autovectorize the code.
By default, the kernel is assumed to have the
__attribute__((vec_type_hint(int)))
qualifier. This lets you specify a
different vectorization type.
Implicit in autovectorization is the assumption that any libraries called from the kernel must be re-compilable at run time to handle cases where the compiler decides to merge or separate workitems. This probably means that such libraries can never be hard coded binaries or that hard coded binaries must be accompanied either by source or some re-targetable intermediate representation. This may be a code security question for some.
Syntax
__attribute__((vec_type_hint(<type>)))
- <type>: is one of the built-in vector types listed in the
following table, or the constituent scalar element types. Note: When not specified, the kernel is assumed to have an INT type.
Type |
Description |
---|---|
charn |
A vector of n 8-bit signed two’s complement integer values. |
ucharn |
A vector of n 8-bit unsigned integer values. |
shortn |
A vector of n 16-bit signed two’s complement integer values. |
ushortn |
A vector of n 16-bit unsigned integer values. |
intn |
A vector of n 32-bit signed two’s complement integer values. |
uintn |
A vector of n 32-bit unsigned integer values. |
longn |
A vector of n 64-bit signed two’s complement integer values. |
ulongn |
A vector of n 64-bit unsigned integer values. |
floatn |
A vector of n 32-bit floating-point values. |
doublen |
A vector of n 64-bit floating-point values. |
Examples
The following example autovectorizes assuming double-wide integer as the basic computation width:
#include <clc.h>
// For VHLS OpenCL C kernels, the full work group is synthesized
__attribute__((vec_type_hint(double)))
__attribute__ ((reqd_work_group_size(16, 1, 1)))
__kernel void
...
See Also
- SDAccel Environment Optimization Guide (UG1207)
- https://www.khronos.org/
- The OpenCL C Specification