This chapter is from the book
This chapter is from the book
This chapter is from the book
Preprocessor Directives and Macros
The preprocessing directives defined by the C99 specification are supported. These include
# non-directive #if #ifdef #ifndef #elif #else #endif #include #define #undef #line #error #pragma
The defined operator is also included.
The following example demonstrates the use of #if, #elif, #else, and #endif preprocessor macros. In this example, we use the preprocessor macros to determine which arithmetic operation to apply in the kernel. The kernel source is described here:
#define OP_ADD 1
#define OP_SUBTRACT 2
#define OP_MULTIPLY 3
#define OP_DIVIDE 4
kernel void
foo(global float *dst, global float *srcA, global float *srcB)
{
size_t id = get_global_id(0);
#if OP_TYPE == OP_ADD
dst[id] = srcA[id] + srcB[id];
#elif OP_TYPE == OP_SUBTRACT
dst[id] = srcA[id] – srcB[id];
#elif OP_TYPE == OP_MULTIPLY
dst[id] = srcA[id] * srcB[id];
#elif OP_TYPE == OP_DIVIDE
dst[id] = srcA[id] / srcB[id];
#else
dst[id] = NAN;
#endif
}
To build the program executable with the appropriate value for OP_TYPE, the application calls clBuildProgram as follows:
// build program so that kernel foo does an add operation
err = clBuildProgram(program, 0, NULL,
"-DOP_TYPE=1", NULL, NULL);
Pragma Directives
The #pragma directive is described as
#pragma pp-tokensopt new-line
A #pragma directive where the preprocessing token OPENCL (used instead of STDC) does not immediately follow pragma in the directive (prior to any macro replacement) causes the implementation to behave in an implementation-defined manner. The behavior might cause translation to fail or cause the translator or the resulting program to behave in a nonconforming manner. Any such pragma that is not recognized by the implementation is ignored. If the preprocessing token OPENCL does immediately follow pragma in the directive (prior to any macro replacement), then no macro replacement is performed on the directive.
The following standard pragma directives are available.
Floating-Point Pragma
The FP_CONTRACT floating-point pragma can be used to allow (if the state is on) or disallow (if the state is off) the implementation to contract expressions. The FP_CONTRACT pragma definition is
#pragma OPENCL FP_CONTRACT on-off-switch
on-off-switch: one of ON OFF DEFAULT
A detailed description of #pragma OPENCL FP_CONTRACT is found in Chapter 5 in the section "Floating-Point Pragmas."
Compiler Directives for Optional Extensions
The #pragma OPENCL EXTENSION directive controls the behavior of the OpenCL compiler with respect to language extensions. The #pragma OPENCL EXTENSION directive is defined as follows, where extension_ name is the name of the extension:
#pragma OPENCL EXTENSION extension_name: behavior
#pragma OPENCL EXTENSION all : behavior
behavior: enable or disable
The extension_name will have names of the form cl_khr_<name> for an extension (such as cl_khr_fp64) approved by the OpenCL working group and will have names of the form cl_<vendor_name>_<name> for vendor extensions. The token all means that the behavior applies to all extensions supported by the compiler. The behavior can be set to one of the values given in Table 4.9.
Table 4.9. Optional Extension Behavior Description
|
Behavior |
Description |
|
enable |
Enable the extension extension_name. Report an error on the #pragma OpenCL EXTENSION if the extension_name is not supported, or if all is specified. |
|
disable |
Behave (including issuing errors and warnings) as if the extension extension_name is not part of the language definition. If all is specified, then behavior must revert back to that of the nonextended core version of the language being compiled to. Warn on the #pragma OPENCL EXTENSION if the extension extension_name is not supported. |
The #pragma OPENCL EXTENSION directive is a simple, low-level mechanism to set the behavior for each language extension. It does not define policies such as which combinations are appropriate; these are defined elsewhere. The order of directives matters in setting the behavior for each extension. Directives that occur later override those seen earlier. The all variant sets the behavior for all extensions, overriding all previously issued extension directives, but only if the behavior is set to disable.
An extension needs to be enabled before any language feature (such as preprocessor macros, data types, or built-in functions) of this extension is used in the OpenCL program source. The following example shows how to enable the double-precision floating-point extension:
#pragma OPENCL EXTENSION cl_khr_fp64 : enable double x = 2.0;
If this extension is not supported, then a compilation error will be reported for double x = 2.0. If this extension is supported, this enables the use of double-precision floating-point extensions in the program source following this directive.
Similarly, the cl_khr_3d_image_writes extension adds new built-in functions that support writing to a 3D image:
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
kernel void my_func(write_only image3d_t img, ...)
{
float4 coord, clr;
...
write_imagef(img, coord, clr);
}
The built-in functions such as write_imagef with image3d_t in the preceding example can be called only if this extension is enabled; otherwise a compilation error will occur.
The initial state of the compiler is as if the following directive were issued, telling the compiler that all error and warning reporting must be done according to this specification, ignoring any extensions:
#pragma OPENCL EXTENSION all : disable
Every extension that affects the OpenCL language semantics or syntax or adds built-in functions to the language must also create a preprocessor #define that matches the extension name string. This #define would be available in the language if and only if the extension is supported on a given implementation. For example, an extension that adds the extension string cl_khr_fp64 should also add a preprocessor #define called cl_khr_fp64. A kernel can now use this preprocessor #define to do something like this:
#ifdef cl_khr_fp64
// do something using this extension
#else
// do something else or #error
#endif
Macros
The following predefined macro names are available:
- __FILE__ is the presumed name of the current source file (a character string literal).
- __LINE__ is the presumed line number (within the current source file) of the current source line (an integer constant).
- CL_VERSION_1_0 substitutes the integer 100, reflecting the OpenCL 1.0 version.
- CL_VERSION_1_1 substitutes the integer 110, reflecting the OpenCL 1.1 version.
- __OPENCL_VERSION__ substitutes an integer reflecting the version number of the OpenCL supported by the OpenCL device. This reflects both the language version supported and the device capabilities as given in Table 4.3 of the OpenCL 1.1 specification. The version of OpenCL described in this book will have __OPENCL_VERSION__ substitute the integer 110.
- __ENDIAN_LITTLE__ is used to determine if the OpenCL device is a little endian architecture or a big endian architecture (an integer constant of 1 if the device is little endian and is undefined otherwise).
-
__kernel_exec(X, typen) (and kernel_exec(X, typen)) is defined as
__kernel __attribute__((work_group_size_hint(X, 1, 1))) __attribute__((vec_type_hint(typen))).
- __IMAGE_SUPPORT__ is used to determine if the OpenCL device supports images. This is an integer constant of 1 if images are supported and is undefined otherwise.
- __FAST_RELAXED_MATH__ is used to determine if the –cl-fast-relaxed-math optimization option is specified in build options given to clBuildProgram. This is an integer constant of 1 if the – cl-fast-relaxed-math build option is specified and is undefined otherwise.
The macro names defined by the C99 specification but not currently supported by OpenCL are reserved for future use.