Как я могу обеспечить порядок между записью и чтением в глобальную память?

У меня есть ядро ​​CUDA следующей формы:

Void launch_kernel(..Arguments...) 
{  
    int i = threadIdx.x
    //Load required data 
    int temp1 = A[i];
    int temp2 = A[i+1];
    int temp3= A[i+2]; 
    // compute step 
    int output1 =  temp1 + temp2 + temp3;
    int output2 = temp1 + temp3;  
    // Store the result
    B[i] = output1; 
    C[i] = output2;
}  

Как обсуждалось в руководстве CUDA, модель согласованности для глобальной памяти графического процессора не является последовательной. В результате может показаться, что операции с памятью выполняются в порядке, отличном от порядка исходной программы. Чтобы принудительно упорядочить память, CUDA предлагает функции __threadfence (). Однако, согласно руководству, такая функция обеспечивает относительный порядок операций чтения и относительный порядок операций записи. Цитата из руководства:

Все записи в общую и глобальную память, сделанные вызывающим потоком перед вызовом __threadfence_block (), наблюдаются всеми потоками в блоке вызывающего потока, как происходящие до всех записей в общую память и глобальную память, сделанных вызывающим потоком после вызова. to __threadfence_block ();

Таким образом, очевидно, что __threadfence () недостаточно для обеспечения упорядочения операций чтения и записи.

Как принудительно упорядочить операции чтения и записи в глобальную память. В качестве альтернативы, как мне убедиться, что все операции чтения гарантированно будут завершены перед выполнением раздела вычислений и сохранения вышеуказанного ядра.