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

ID 683846
Date 12/13/2021

A newer version of this document is available. Customers should click here to go to the newest version.

Document Table of Contents

5.8. Declaring __constant Address Space Qualifiers

There are several limitations and workarounds you must consider when you include __constant address space qualifiers in your kernel.

Function Scope __constant Variables

The Intel® FPGA SDK for OpenCL™ Offline Compiler does not support function scope __constant variables. Replace function scope __constant variables with file scope constant variables. You can also replace function scope __constant variables with __constant buffers that the host passes to the kernel.

File Scope __constant Variables

If the host always passes the same constant data to your kernel, consider declaring that data as a constant preinitialized file scope array within the kernel file. Declaration of a constant preinitialized file scope array creates a ROM directly in the hardware to store the data. This ROM is available to all work-items in the NDRange.

For example:

__constant int my_array[8] = {0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7};

__kernel void my_kernel (__global int * my_buffer)
   size_t gid = get_global_id(0);
   my_buffer[gid] += my_array[gid % 8];

In this case, the offline compiler sets the values for my_array in a ROM because the file scope constant data does not change between kernel invocations.

Pointers to __constant Parameters from the Host

You can replace file scope constant data with a pointer to a __constant parameter in your kernel code if the data is not fixed across kernel invocations. You must then modify your host application in the following manner:

  1. Create cl_mem memory objects associated with the pointers in global memory.
  2. Load constant data into cl_mem objects with clEnqueueWriteBuffer prior to kernel execution.
  3. Pass the cl_mem objects to the kernel as arguments with the clSetKernelArg function.

For simplicity, if a constant variable is of a complex type, use a typedef argument, as shown in the table below:

Table 8.  Replacing File Scope __constant Variable with Pointer to __constant Parameter
If your source code is structured as follows: Rewrite your code to resemble the following syntax:
__constant int Payoff[2][2] = {{ 1, 3}, {5, 3}};
__kernel void original(__global int * A)
   *A = Payoff[1][2];
   // and so on
__kernel void modified(__global int * A,
__constant Payoff_type * PayoffPtr )
   *A = (PayoffPtr)[1][2];
   // and so on
Attention: Use the same type definition in both your host application and your kernel.