ID 683846
Date 12/19/2022
Public

## 7.9. Defining Preprocessor Macros to Specify Kernel Parameters (-D<macro_name>)

The Intel® FPGA SDK for OpenCL™ Offline Compiler supports preprocessor macros that allow you to pass macro definitions and compile code on a conditional basis.
• To pass a preprocessor macro definition to the offline compiler, invoke the aoc -D <macro_name> <kernel_filename>.cl command.
• To override the existing value of a defined preprocessor macro, invoke the aoc -D <macro_name>=<value> <kernel_filename>.cl command.
Consider the following code snippet for the kernel sum:
#ifndef UNROLL_FACTOR
#define UNROLL_FACTOR 1
#endif

__kernel void sum (__global const int * restrict x,
__global int * restrict sum)
{
int accum = 0;

#pragma unroll UNROLL_FACTOR
for(size_t i = 0; i < 4; i++)
{
accum += x[i + get_global_id(0) * 4];
}
sum[get_global_id(0)] = accum;
}

To override the UNROLL_FACTOR of 1 and set it to 4, invoke the aoc -DUNROLL_FACTOR=4 sum.cl command. Invoking this command is equivalent to replacing the line #define UNROLL_FACTOR 1 with #define UNROLL_FACTOR 4 in the sum kernel source code.

• To use preprocessor macros to control how the offline compiler optimizes your kernel without modifying your kernel source code, invoke the aoc -o <hardware_filename>.aocx -D <macro_name>=<value> <kernel_filename>.cl
Where:

-o is the offline compiler option you use to specify the name of the .aocx file that the offline compiler generates.

<hardware_filename> is the name of the .aocx file that the offline compiler generates using the preprocessor macro value you specify.

Tip: To preserve the results from both compilations on your file system, compile your kernels as separate binaries by using the -o flag of the aoc command.
For example, if you want to compile the same kernel multiple times with required work-group sizes of 64 and 128, you can define a WORK_GROUP_SIZE preprocessor macro for the kernel attribute reqd_work_group_size, as shown below:
__attribute__((reqd_work_group_size(WORK_GROUP_SIZE,1,1)))
__kernel void myKernel(...)
for (size_t i = 0; i < 1024; i++)
{
// statements
}

Compile the kernel multiple times by typing the following commands:

aoc –o myKernel_64.aocx –DWORK_GROUP_SIZE=64 myKernel.cl

aoc –o myKernel_128.aocx –DWORK_GROUP_SIZE=128 myKernel.cl