OpenCL Attributes

Description

The ALWAYS_INLINE attribute indicates that a function must be inlined. This attribute is a standard feature of GCC, and a standard feature of theSDxtool compilers.

This attribute enables a compiler optimization to have a function inlined into the calling function. The inlined function is dissolved and no longer appears as a separate level of hierarchy in the RTL.

In some cases, inlining a function allows operations within the function to be shared and optimized more effectively with surrounding operations in the calling function. However, an inlined function can no longer be shared with other functions, so the logic might be duplicated between the inlined function and a separate instance of the function which can be more broadly shared. While this can improve performance, this will also increase the area required for implementing the RTL.

ForOpenCLkernels, theSDxcompiler uses its own rules to inline or not inline a function. To directly control inlining functions, use the ALWAYS_INLINE or NOINLINE attributes.

By default, inlining is only performed on the next level of function hierarchy, not sub-functions.

Syntax

__attribute__((always_inline))

Examples

This example adds the ALWAYS_INLINE attribute to functionfoo:

__attribute__((always_inline)) void foo ( a, b, c, d ) { ... }

This example prevents the inlining of the functionfoo:

__attribute__((noinline)) void foo ( a, b, c, d ) { ... }

See Also

• https://gcc.gnu.org
• SDAccel Environment Profiling and Optimization Guide(UG1207)

Description

Loop unrolling is the first optimization technique available in theSDAcceldevelopment environment. The purpose of the loop unroll optimization is to expose concurrency to the compiler. This newly exposed concurrency reduces latency and improves performance, but also consumes more FPGA fabric resources.

The OPENCL_UNROLL_HINT attribute is part of theOpenCLLanguage Specification, and specifies that loops (for,while,do) can be unrolled by theOpenCLcompiler.

The OPENCL_UNROLL_HINT attribute qualifier must appear immediately before the loop to be affected. You can use this attribute to specify full unrolling of the loop, partial unrolling by a specified amount, or to disable unrolling of the loop.

Syntax

Place the attribute in theOpenCLsource before the loop definition:

__attribute__((opencl_unroll_hint()))

Where:

• is an optional loop unrolling factor and must be a positive integer, or compile time constant expression. An unroll factor of 1 disables unrolling.
  TIP:If is not specified, the compiler automatically determines the unrolling factor for the loop.

Example 1

The following example unrolls theforloop by a factor of 2. This results in two parallel loop iterations instead of four sequential iterations for the compute unit to complete the operation.

__attribute__((opencl_unroll_hint(2))) for(int i = 0; i < LENGTH; i++) { bufc[i] = bufa[i] * bufb[i]; }

Conceptually the compiler transforms the loop above to the following code:

for(int i = 0; i < LENGTH; i+=2) { bufc[i] = bufa[i] * bufb[i]; bufc[i+1] = bufa[i+1] * bufb[i+1]; }

See Also

• SDAccel Environment Profiling and Optimization Guide(UG1207)
• https://www.khronos.org/
• The OpenCL C Specification

Description

WhenOpenCLAPI kernels are submitted for execution on anOpenCLdevice, they execute within an index space, called an ND range, which can have 1, 2, or 3 dimensions. This is called the global size in theOpenCLAPI. The work-group size defines the amount of the ND range that can be processed by a single invocation of a kernel compute unit. The work-group size is also called the local size in theOpenCLAPI. TheOpenCLcompiler can determine the work-group size based on the properties of the kernel and selected device. After the work-group size (local size) is determined, the ND range (global size) is divided automatically into work-groups, and the work-groups are scheduled for execution on the device.

Although theOpenCLcompiler can define the work-group size, the specification of the REQD_WORK_GROUP_SIZE attribute on the kernel to define the work-group size is highly recommended for FPGA implementations of the kernel. The attribute is recommended for performance optimization during the generation of the custom logic for a kernel. SeeOpenCL Execution Modelin theSDAccel Environment Profiling and Optimization Guide(UG1207)for more information.

OpenCLkernel functions are executed exactly one time for each point in the ND range index space. This unit of work for each point in the ND range is called a work-item. Work-items are organized into work-groups, which are the unit of work scheduled onto compute units. The optional REQD_WORK_GROUP_SIZE attribute defines the work-group size of a compute unit that must be used as thelocal_work_sizeargument toclEnqueueNDRangeKernel. This allows the compiler to optimize the generated code appropriately for this kernel.

Syntax

Place this attribute before the kernel definition, or before the primary function specified for the kernel:

__attribute__((reqd_work_group_size(, , )))

Where:

• , , : Specifies the ND range of the kernel. This represents each dimension of a three dimensional matrix specifying the size of the work-group for the kernel.

Examples

The followingOpenCLC kernel code shows a vector addition design where two arrays of data are summed into a third array. The required size of the work-group is 16x1x1. This kernel will execute 16 times to produce a valid result.

#include  // For VHLS OpenCL C kernels, the full work group is synthesized __attribute__ ((reqd_work_group_size(16, 1, 1))) __kernel void vadd(__global int* a, __global int* b, __global int* c) { int idx = get_global_id(0); c[idx] = a[idx] + b[idx]; }

See Also

• SDAccel Environment Profiling and Optimization Guide(UG1207)
• https://www.khronos.org/
• The OpenCL C Specification

Description

The optional__attribute__((vec_type_hint()))is part of theOpenCLLanguage Specification, and hints to theOpenCLcompiler representing the computational width of the kernel, providing a basis for calculating processor bandwidth usage when the compiler is looking to auto-vectorize the code.

By default, the kernel is assumed to have the__attribute__((vec_type_hint(int)))qualifier. This lets you specify a different vectorization type.

Implicit in autovectorization is the assumption that any libraries called from the kernel must be re-compilable at runtime to handle cases where the compiler decides to merge or separate workitems. This means that these libraries can never be hard-coded binaries or that hard-coded binaries must be accompanied either by source or some re-targetable intermediate representation. This might be a code security question for some.

Syntax

__attribute__((vec_type_hint()))

Where:

• : is one of the built-in vector types listed in the following table, or the constituent scalar element types.
  Note:When not specified, the kernel is assumed to have an INT type.

Examples

The following example autovectorizes assuming double-wide integer as the basic computation width:

#include  // For VHLS OpenCL C kernels, the full work group is synthesized __attribute__((vec_type_hint(double))) __attribute__ ((reqd_work_group_size(16, 1, 1))) __kernel void ...

See Also

• SDAccel Environment Profiling and Optimization Guide(UG1207)
• https://www.khronos.org/
• The OpenCL C Specification

Description

The work-group size in theOpenCLAPI standard defines the size of the ND range space that can be handled by a single invocation of a kernel compute unit. WhenOpenCLkernels are submitted for execution on anOpenCLdevice, they execute within an index space, called an ND range, which can have 1, 2, or 3 dimensions.See "OpenCL Execution Model" inSDAccel Environment Profiling and Optimization Guide(UG1207)for more information.

OpenCLkernel functions are executed exactly one time for each point in the ND range index space. This unit of work for each point in the ND range is called a work-item. Unlikeforloops in C, where loop iterations are executed sequentially and in-order, anOpenCLruntime and device is free to execute work-items in parallel and in any order.

Work-items are organized into work-groups, which are the unit of work scheduled onto compute units. The optional WORK_GROUP_SIZE_HINT attribute is part of theOpenCLLanguage Specification, and is a hint to the compiler that indicates the work-group size value most likely to be specified by thelocal_work_sizeargument toclEnqueueNDRangeKernel. This allows the compiler to optimize the generated code according to the expected value.

Syntax

Place this attribute before the kernel definition, or before the primary function specified for the kernel:

__attribute__((work_group_size_hint(, , )))

Where:

• , , : Specifies the ND range of the kernel. This represents each dimension of a three dimensional matrix specifying the size of the work-group for the kernel.

Examples

The following example is a hint to the compiler that the kernel will most likely be executed with a work-group size of 1:

__attribute__((work_group_size_hint(1, 1, 1))) __kernel void ...

See Also

• SDAccel Environment Profiling and Optimization Guide(UG1207)
• https://www.khronos.org/
• The OpenCL C Specification

Description

An advantage of using the FPGA over other compute devices forOpenCLprograms is the ability for the application programmer to customize the memory architecture all throughout the system and into the compute unit. By default, theSDAccelcompiler generates a memory architecture within the compute unit that maximizes local and private memory bandwidth based on static code analysis of the kernel code. Further optimization of these memories is possible based on attributes in the kernel source code, which can be used to specify physical layouts and implementations of local and private memories. The attribute in theSDAccelcompiler to control the physical layout of memories in a compute unit isarray_partition.

For one-dimensional arrays, the XCL_ARRAY_PARTITION attribute implements an array declared within kernel code as multiple physical memories instead of a single physical memory. The selection of which partitioning scheme to use depends on the specific application and its performance goals. The array partitioning schemes available in theSDAcceltool compiler arecyclic,block, andcomplete.

Syntax

Place the attribute with the definition of the array variable:

__attribute__((xcl_array_partition(, , )))

Where:

• : Specifies one of the following partition types:
  • cyclic: Cyclic partitioning is the implementation of an array as a set of smaller physical memories that can be accessed simultaneously by the logic in the compute unit. The array is partitioned cyclically by putting one element into each memory before coming back to the first memory to repeat the cycle until the array is fully partitioned.
  • block: Block partitioning is the physical implementation of an array as a set of smaller memories that can be accessed simultaneously by the logic inside of the compute unit. In this case, each memory block is filled with elements from the array before moving on to the next memory.
  • complete: Complete partitioning decomposes the array into individual elements. For a one-dimensional array, this corresponds to resolving a memory into individual registers. The defaultiscomplete.
• : For cyclic type partitioning, thespecifies how many physical memories to partition the original array into in the kernel code. For block type partitioning, thespecifies the number of elements from the original array to store in each physical memory.
  IMPORTANT:For completetype partitioning, the > is not specified.
• : Specifies which array dimension to partition. Specified as an integer from 1 to .SDAccelenvironment supports arrays of N dimensions and can partition the array on any single dimension.

Example 1

For example, consider the following array declaration:

int buffer[16];

The integer array, named buffer, stores 16 values that are 32-bits wide each. Cyclic partitioning can be applied to this array with the following declaration:

int buffer[16] __attribute__((xcl_array_partition(cyclic,4,1)));

In this example, the cyclicattribute tells theSDAccelcompiler to distribute the contents of the array among four physical memories. This attribute increases the immediate memory bandwidth for operations accessing the array buffer by a factor of four.

All arrays inside of a compute unit in the context of theSDAccelenvironment are capable of sustaining a maximum of two concurrent accesses. By dividing the original array in the code into four physical memories, the resulting compute unit can sustain a maximum of eight concurrent accesses to the array buffer.

Example 2

Using the same integer array as found in Example 1, block partitioning can be applied to the array with the following declaration:

int buffer[16] __attribute__((xcl_array_partition(block,4,1)));

Because the size of the block is four, theSDAccelcompiler will generate four physical memories, sequentially filling each memory with data from the array.

Example 3

Using the same integer array as found in Example 1, complete partitioning can be applied to the array with the following declaration:

int buffer[16] __attribute__((xcl_array_partition(complete, 1)));

In this example, the array is completely partitioned into distributed RAM, or 16 independent registers in the programmable logic of the kernel. Because complete is the default, the same effect can also be accomplished with the following declaration:

int buffer[16] __attribute__((xcl_array_partition));

While this creates an implementation with the highest possible memory bandwidth, it is not suited to all applications. The way in which data is accessed by the kernel code through either constant or data dependent indexes affects the amount of supporting logic that theSDAccelcompiler has to build around each register to ensure functional equivalence with the usage in the original code. As a general best practice guideline for theSDxenvironment, the complete partitioning attribute is best suited for arrays in which at least one dimension of the array is accessed through the use of constant indexes.

See Also

• xcl_array_reshape
• pragma HLS array_partition
• SDAccel Environment Profiling and Optimization Guide(UG1207)
• Vivado Design Suite User Guide: High-Level Synthesis(UG902)

Description

This attribute combines array partitioning with vertical array mapping.

The XCL_ARRAY_RESHAPE attribute combines the effect of XCL_ARRAY_PARTITION, breaking an array into smaller arrays, and concatenating elements of arrays by increasing bit-widths. This reduces the number of block RAM consumed while providing parallel access to the data. This attribute creates a new array with fewer elements but with greater bit-width, allowing more data to be accessed in a single clock cycle.

Given the following code:

void foo (...) { int array1[N] __attribute__((xcl_array_reshape(block, 2, 1))); int array2[N] __attribute__((xcl_array_reshape(cycle, 2, 1))); int array3[N] __attribute__((xcl_array_reshape(complete, 1))); ... }

The ARRAY_RESHAPE attribute transforms the arrays into the form shown in the following figure:

Figure:ARRAY_RESHAPE

Syntax

Place the attribute with the definition of the array variable:

__attribute__((xcl_array_reshape(,, )))

Where:

• : Specifies one of the following partition types:
  • cyclic: Cyclic partitioning is the implementation of an array as a set of smaller physical memories that can be accessed simultaneously by the logic in the compute unit. The array is partitioned cyclically by putting one element into each memory before coming back to the first memory to repeat the cycle until the array is fully partitioned.
  • block: Block partitioning is the physical implementation of an array as a set of smaller memories that can be accessed simultaneously by the logic inside of the compute unit. In this case, each memory block is filled with elements from the array before moving on to the next memory.
  • complete: Complete partitioning decomposes the array into individual elements. For a one-dimensional array, this corresponds to resolving a memory into individual registers. The defaultiscomplete.
• : For cyclic type partitioning, thespecifies how many physical memories to partition the original array into in the kernel code. For Block type partitioning, thespecifies the number of elements from the original array to store in each physical memory.
  IMPORTANT:For completetype partitioning, the should not be specified.
• : Specifies which array dimension to partition. Specified as an integer from 1 to .SDAccelenvironment supports arrays of dimensions and can partition the array on any single dimension.

Example 1

Reshapes (partition and maps) an 8-bit array with 17 elements, AB[17], into a new 32-bit array with five elements using block mapping.

int AB[17] __attribute__((xcl_array_reshape(block,4,1)));

Example 2

Reshapes the two-dimensional array AB[6][4] into a new array of dimension [6][2], in which dimension 2 has twice the bit-width:

int AB[6][4] __attribute__((xcl_array_reshape(block,2,2)));

Example 3

Reshapes the three-dimensional 8-bit array, AB[4][2][2] in functionfoo, into a new single element array (a register), 128 bits wide (4*2*2*8):

int AB[4][2][2] __attribute__((xcl_array_reshape(complete,0)));

See Also

• xcl_array_partition
• pragma HLS array_reshape
• SDAccel Environment Profiling and Optimization Guide(UG1207)
• Vivado Design Suite User Guide: High-Level Synthesis(UG902)

Description

The XCL_LATENCY attribute specifies a minimum, or maximum latency value, or both, for the completion of functions, loops, and regions. Latency is defined as the number of clock cycles required to produce an output. Function or region latency is the number of clock cycles required for the code to compute all output values, and return. Loop latency is the number of cycles to execute all iterations of the loop. See "Performance Metrics Example" ofVivado Design Suite User Guide: High-Level Synthesis(UG902).

TheVivadoHigh-Level Synthesis (HLS) tool always tries to minimize latency in the design. When the XCL_LATENCY attribute is specified, the tool behavior is as follows:

• When latency is greater than the minimum, or less than the maximum: The constraint is satisfied. No further optimizations are performed.
• When latency is less than the minimum: If the HLS tool can achieve less than the minimum specified latency, it extends the latency to the specified value, potentially increasing sharing.
• When latency is greater than the maximum: If the HLS tool cannot schedule within the maximum limit, it increases effort to achieve the specified constraint. If it still fails to meet the maximum latency, it issues a warning, and produces a design with the smallest achievable latency in excess of the maximum.

Syntax

__attribute__((xcl_latency(min, max)))

Where:

• : Specifies the minimum latency for the function, loop, or region of code.
• : Specifies the maximum latency for the function, loop, or region of code.

Example 1

Theforloop in thetestfunction is specified to have a minimum latency of 4 and a maximum latency of 8:

__kernel void test(__global float *A, __global float *B, __global float *C, int id) { for (unsigned int i = 0; i < id; i++) __attribute__((xcl_latency(4, 12))) { C[id] = A[id] * B[id]; } }

See Also

• pragma HLS latency
• SDAccel Environment Profiling and Optimization Guide(UG1207)
• Vivado Design Suite User Guide: High-Level Synthesis(UG902)

Description

The XCL_LATENCY attribute specifies a minimum, or maximum latency value, or both, for the completion of functions, loops, and regions. Latency is defined as the number of clock cycles required to produce an output. Function or region latency is the number of clock cycles required for the code to compute all output values, and return. Loop latency is the number of cycles to execute all iterations of the loop. See "Performance Metrics Example" ofVivado Design Suite User Guide: High-Level Synthesis(UG902).

TheVivadoHigh-Level Synthesis (HLS) tool always tries to minimize latency in the design. When the XCL_LATENCY attribute is specified, the tool behavior is as follows:

• When latency is greater than the minimum, or less than the maximum: The constraint is satisfied. No further optimizations are performed.
• When latency is less than the minimum: If the HLS tool can achieve less than the minimum specified latency, it extends the latency to the specified value, potentially increasing sharing.
• When latency is greater than the maximum: If the HLS tool cannot schedule within the maximum limit, it increases effort to achieve the specified constraint. If it still fails to meet the maximum latency, it issues a warning, and produces a design with the smallest achievable latency in excess of the maximum.

Syntax

__attribute__((xcl_latency(min, max)))

Where:

• : Specifies the minimum latency for the function, loop, or region of code.
• : Specifies the maximum latency for the function, loop, or region of code.

Example 1

Theforloop in thetestfunction is specified to have a minimum latency of 4 and a maximum latency of 8:

__kernel void test(__global float *A, __global float *B, __global float *C, int id) { for (unsigned int i = 0; i < id; i++) __attribute__((xcl_latency(4, 12))) { C[id] = A[id] * B[id]; } }

See Also

• pragma HLS latency
• SDAccel Environment Profiling and Optimization Guide(UG1207)
• Vivado Design Suite User Guide: High-Level Synthesis(UG902)