Logo Questions Linux Laravel Mysql Ubuntu Git Menu
 

When to use volatile with shared CUDA Memory

Under what circumstances should you use the volatile keyword with a CUDA kernel's shared memory? I understand that volatile tells the compiler never to cache any values, but my question is about the behavior with a shared array:

__shared__ float products[THREADS_PER_ACTION];

// some computation
products[threadIdx.x] = localSum;

// wait for everyone to finish their computation
__syncthreads();

// then a (basic, ugly) reduction:
if (threadIdx.x == 0) {
    float globalSum = 0.0f;
    for (i = 0; i < THREADS_PER_ACTION; i++)
        globalSum += products[i];
}

Do I need products to be volatile in this case? Each array entry is only accessed by a single thread, except at the end, where everything is read by thread 0. Is it possible that the compiler could cache the entire array, and so I need it to be volatile, or will it only cache elements?

Thanks!

like image 393
Taj Morton Avatar asked Mar 11 '13 04:03

Taj Morton


People also ask

Does shared memory need to be volatile?

Yes it is critical. Like you said volatile prevents code breaking optimization on shared memory [C++98 7.1.

When should volatile be used?

A variable should be declared volatile whenever its value could change unexpectedly. In practice, only three types of variables could change: Memory-mapped peripheral registers. Global variables modified by an interrupt service routine.

Can CUDA use shared memory?

Shared memory is a powerful feature for writing well optimized CUDA code. Access to shared memory is much faster than global memory access because it is located on chip. Because shared memory is shared by threads in a thread block, it provides a mechanism for threads to cooperate.

When should I use volatile in C#?

The volatile keyword in C# is used to inform the JIT compiler that the value of the variable should never be cached because it might be changed by the operating system, the hardware, or a concurrently executing thread.


1 Answers

If you don't declare a shared array as volatile, then the compiler is free to optimize locations in shared memory by locating them in registers (whose scope is specific to a single thread), for any thread, at it's choosing. This is true whether you access that particular shared element from only one thread or not. Therefore, if you use shared memory as a communication vehicle between threads of a block, it's best to declare it volatile. However, this sort of communication pattern often also requires execution barriers to enforce ordering of reads/writes, so continue reading about barriers below.

Obviously if each thread only accessed its own elements of shared memory, and never those associated with another thread, then this does not matter, and the compiler optimization will not break anything.

In your case, where you have a section of code where each thread is accessing it's own elements of shared memory, and the only inter-thread access occurs at a well understood location, you could use a memory fence function to force the compiler to evict any values that are temporarily stored in registers, back out to the shared array. So you might think that __threadfence_block() might be useful, but in your case, __syncthreads() already has memory-fencing functionality built in. So your __syncthreads() call is sufficient to force thread synchronization as well as to force any register-cached values in shared memory to be evicted back to shared memory.

By the way, if that reduction at the end of your code is of performance concern, you could consider using a parallel reduction method to speed it up.

like image 137
Robert Crovella Avatar answered Oct 26 '22 23:10

Robert Crovella