Visible to Intel only — GUID: vah1521223862995
Ixiasoft
Visible to Intel only — GUID: vah1521223862995
Ixiasoft
5.8. Declaring __constant Address Space Qualifiers
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:
- Create cl_mem memory objects associated with the pointers in global memory.
- Load constant data into cl_mem objects with clEnqueueWriteBuffer prior to kernel execution.
- 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:
If your source code is structured as follows: | Rewrite your code to resemble the following syntax: |
---|---|
|
|