Intel® FPGA SDK for OpenCL™ Standard Edition: Programming Guide

ID 683342
Date 4/22/2019
Public
Document Table of Contents

7.7. 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