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

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

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.