#
Because the compiler can do lots of optimizations for performance, the order of memory accesses in the machine code might not be the same as the order of memory accesses implemented in the source code. If the memory is only accessed by a single thread in the program, the memory access order changed by the compiler usually does not matter. Otherwise it is a compiler bug. However, if the memory is accessed by multiple threads or even multiple processes, the memory access order has to be carefully managed. In many cases, we don’t want the compiler to optimize the memory access order for those which are critical for the desired program behavior. This is a common problem in parallel programming not only for multi-threaded CPU programming but also for CUDA programming on GPU.
In CUDA programming, we can use CUDA memory fence functions to enforce a sequentially-consistent ordering on memory accesses. In this blog post, I would like to discuss the memory fence functions and their motivations in CUDA programming.
#
In computing, a memory fence or memory barrier, is a type of fence instruction that causes a processor or compiler to enforce an ordering constraint on memory operations issued before and after the fence instruction. This typically means that operations issued prior to the fence are guaranteed to be performed before operations issued after the fence.
The following example shows why memory fence is necessary in parallel programming. Suppose we have two threads, one thread is modifying a variable x, and the other thread is printing the modified x.
|
1234567891011121314151617
|
#include <iostream>int x{0};int f{0};void print_modified_x(){ while (f == 0) {} std::cout << "x = " << x << std::endl;}void modify_x(){ x = 42; // Indicate that x has been modified f = 1;}
|
The desired behavior of the program is that the print_modified_x function will hang until the modify_x function modifies x, and then it should print x = 42. However, the program might not behave as expected because the compiler might optimize the memory access order. In the print_modified_x function, the order of the while loop and the x print statement might be changed by the compiler. In the modify_x function, the order of the x assignment and the f assignment might also be changed by the compiler. For example, if f is assigned before x in the modify_x function and x is printed from the print_modified_x function before x is assigned, the program will print x = 0 instead of x = 42.
The desired behavior of the program can be ensured if the while loop is performed before the x print statement in the print_modified_x function, and the x assignment is performed before the f assignment in the modify_x function. This sequentially-consistent ordering on memory accesses can be enforced via memory fence.
#
There are three CUDA memory fence functions: __threadfence_block, __threadfence, and __threadfence_system. The __threadfence_block function enforces a sequentially-consistent ordering on memory accesses issued by the calling thread and is visible to all threads in the thread block. The __threadfence function enforces a sequentially-consistent ordering on memory accesses issued by the calling thread and is visible to all threads in the device. The __threadfence_system function enforces a sequentially-consistent ordering on memory accesses issued by the calling thread and is visible to all threads in the system, including all threads in the device, host threads, and all threads in peer devices.
In the following example, thread 1 executes write_xy, while thread 2 executes read_xy.
|
1234567891011121314
|
__device__ int X{1};__device__ int Y{2};__device__ void write_xy(){ X = 10; Y = 20;}__device__ void read_xy(){ int B{Y}; int A{X};}
|
Because there is no memory fence that preserves sequentially-consistent ordering on memory accesses, the following four outcomes of the variables A and B are possible:
A = 1andB = 2A = 10andB = 2A = 1andB = 20A = 10andB = 20
If we insert a memory fence function __threadfence_system in the write_xy between the assignments to X and Y, the assignment of X will be performed before the assignment of Y. Similarly, if we insert a memory fence function __threadfence_system in the read_xy before the assignments to A and B, the assignment of B will be performed before the assignment of A.
|
12345678910111213141516
|
__device__ int X{1};__device__ int Y{2};__device__ void write_xy(){ X = 10; __threadfence_system(); Y = 20;}__device__ void read_xy(){ int B{Y}; __threadfence_system(); int A{X};}
|
With the memory fence functions, the following three outcomes of the variables A and B are possible:
A = 1andB = 2A = 10andB = 2A = 10andB = 20
The A = 1 and B = 20 outcome is no longer possible because the memory fence function in write_xy ensures that when Y = 20 is performed, X = 10 has already been performed and are visible to all threads in the system.
With the memory fence functions, all the six possible execution orders of the write_xy and read_xy are listed as follows.
X = 10,Y = 20,B = Y,A = XX = 10,B = Y,Y = 20,A = XX = 10,B = Y,A = X,Y = 20B = Y,X = 10,Y = 20,A = XB = Y,X = 10,A = X,Y = 20B = Y,A = X,X = 10,Y = 20
Notice that X = 10 is always performed before Y = 20 and B = Y is always performed before A = X.
Now the question is, if we can ensure the sequentially-consistent ordering on memory accesses in one thread, why do we need to ensure the sequentially-consistent ordering on memory accesses in all threads in the thread block, all threads in the device, and all threads in the system and have three distinct memory fence functions, __threadfence_block, __threadfence, and __threadfence_system?
This is because usually we have memory hierarchy in the system, and the memory fence functions are used to ensure the sequentially-consistent ordering on memory accesses in different levels of the memory hierarchy.
In our case, if __threadfence is used as the memory fence function in write_xy and read_xy but the thread executing write_xy and the thread executing read_xy are from two different devices, because X and Y might be cached in the device memory on both devices, the __threadfence function only ensures that all the threads in the device where write_xy is executed see the X and Y assignments in the desired order, and the thread executing read_xy reads X and Y in the desired order. However, it is possible that X and Y assignments are performed in the desired order in the cache and when the assignments are flushed to the global memory, the desired order is not preserved. Therefore, using __threadfence, when write_xy is performed on one device and read_xy is performed on another device, the A = 1 and B = 20 outcome is still possible.
|
12345678910111213141516
|
__device__ int X{1};__device__ int Y{2};__device__ void write_xy(){ X = 10; __threadfence(); Y = 20;}__device__ void read_xy(){ int B{Y}; __threadfence(); int A{X};}
|
In order words, the memory fence function __threadfence only ensure the sequentially-consistent ordering on memory accesses for the memory that is observable to the same device but it does not ensure the sequentially-consistent ordering on memory accesses for the memory that is observable to all the threads in the system.
Similarly, the memory fence function __threadfence_block only ensures the sequentially-consistent ordering on memory accesses for the memory that is observable to the same thread block, and it does not ensure the sequentially-consistent ordering on memory accesses for the memory that is observable to the all the threads in the same device or all the threads in the system.
Apparently, the memory fence function __threadfence_system is the most powerful memory fence function that ensures the sequentially-consistent ordering on memory accesses for the memory that is observable to all the threads in the system, but it is also the most expensive memory fence function than __threadfence and __threadfence_block. The memory fence function __threadfence is also more expensive than __threadfence_block.
If we can know what level of sequentially-consistent ordering on memory accesses is required for the memory that is observable to the threads, we can choose the appropriate memory fence function to use and preserve the performance of the program as much as possible. For example, in our case, if we know that the write_xy and read_xy are always executed on the same device but from two different thread blocks, we can just use the memory fence function __threadfence, and if we know that the write_xy and read_xy are always executed on the same thread block, we can just use the memory fence function __threadfence_block.
#
Memory fence ensures the memory accesses before the fence are guaranteed to be performed before the memory accesses after the fence. Synchronization ensures the operations before the synchronization are performed before the operations after the synchronization for all the threads involved in the synchronization.
#
CUDA Memory Fence