5.5.5.7. Enforcing the Order of Pipe Calls

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

Download PDF

ID 683846

Date 6/21/2022

Version

Public

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

Visible to Intel only — Ixiasoft

View Details

5.5.5.7. Enforcing the Order of Pipe Calls

To enforce the order of pipe calls, introduce memory fence or barrier functions in your kernel program to control memory accesses. A memory fence function is necessary to create a control flow dependence between the pipe synchronization calls before and after the fence.

When the Intel® FPGA SDK for OpenCL™ Offline Compiler generates a compute unit, it does not create instruction-level parallelism on all instructions that are independent of each other. As a result, pipe read and write operations might not execute independently of each other even if there is no control or data dependence between them. When pipe calls interact with each other, or when pipes write data to external devices, deadlocks might occur.

For example, the code snippet below consists of a producer kernel and a consumer kernel. Pipes c0 and c1 are unbuffered pipes. The schedule of the pipe read operations from c0 and c1 might occur in the reversed order as the pipe write operations to c0 and c1. That is, the producer kernel writes to c0 but the consumer kernel might read from c1 first. This rescheduling of pipe calls might cause a deadlock because the consumer kernel is reading from an empty pipe.

__kernel void producer (__global const uint * restrict src,
                        const uint iterations,
                        write_only pipe uint __attribute__((blocking)) c0,
                        write_only pipe uint __attribute__((blocking)) c1)
{
    for (int i = 0; i < iterations; i++) {
        write_pipe (c0, &src[2*i  ]);
        write_pipe (c1, &src[2*i+1]); 
    }
}

__kernel void consumer (__global uint * restrict dst,
                        const uint iterations,
                        read_only pipe uint __attribute__((blocking)) c0,
                        read_only pipe uint __attribute__((blocking)) c1)
{
    for (int i = 0; i < iterations; i++) {
        read_pipe (c0, &dst[2*i+1]);
        read_pipe( c1, &dst[2*i]); 
    }
}

To prevent deadlocks from occurring by enforcing the order of pipe calls, include memory fence functions (mem_fence) in your kernel.

Inserting the mem_fence call with each kernel's pipe flag forces the sequential ordering of the write and read pipe calls. The code snippet below shows the modified producer and consumer kernels:

__kernel void producer (__global const uint * src,
                        const uint iterations,
                        write_only_pipe uint __attribute__((blocking)) c0,
                        write_only_pipe uint __attribute__((blocking)) c1)
{
    for (int i = 0; i < iterations; i++)
    {
        write_pipe(c0, &src[2*i  ]);
        mem_fence(CLK_CHANNEL_MEM_FENCE);
        write_pipe(c1, &src[2*i+1]);
    }
}

__kernel void consumer (__global uint * dst;
                        const uint iterations,
                        read_only_pipe uint __attribute__((blocking)) c0,
		              read_only_pipe uint __attribute__((blocking)) c1)
{
    for(int i = 0; i < iterations; i++)
    {
        read_pipe(c0, &dst[2*i  ]);
        mem_fence(CLK_CHANNEL_MEM_FENCE);
        read_pipe(c1, &dst[2*i+1]);
    }
}

In this example, mem_fence in the producer kernel ensures that the pipe write operation to c0 occurs before that to c1. Similarly, mem_fence in the consumer kernel ensures that the pipe read operation from c0 occurs before that from c1.

Defining Memory Consistency Across Kernels When Using Pipes

According to the OpenCL™ Specification version 2.0, memory behavior is undefined unless a kernel completes execution. A kernel must finish executing before other kernels can visualize any changes in memory behavior. However, kernels that use pipes can share data through common global memory buffers and synchronized memory accesses. To ensure that data written to a pipe is visible to the read pipe after execution passes a memory fence, define memory consistency across kernels with respect to memory fences.

To create a control flow dependency between the pipe synchronization calls and the memory operations, add the CLK_GLOBAL_MEM_FENCE flag to the mem_fence call.

For example:

__kernel void producer (__global const uint * restrict src,
                        const uint iterations,
                        write_only pipe uint __attribute__((blocking)) c0,
                        write_only pipe uint __attribute__((blocking)) c1)
{
    for (int i = 0; i < iterations; i++)
    {
        write_pipe(c0, &src[2*i]);
        mem_fence(CLK_CHANNEL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
        write_pipe(c1, &src[2*i+1]);
    }
}

In this kernel, the mem_fence function ensures that the write operation to c0 and memory access to src[2*i] occur before the write operation to c1 and memory access to src[2*i+1]. This allows data written to c0 to be visible to the read pipe before data is written to c1.

Select Your Language

Using Intel.com Search

Quick Links

Recent Searches

Advanced Search

Only search in

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

5.5.5.7. Enforcing the Order of Pipe Calls

Defining Memory Consistency Across Kernels When Using Pipes