Caching and volatile variable with intel integrated GPU

[Dutta__Sankha]

Hello Everyone 

I am working with intel integrated GPU kabylake i7-7700K on ubuntu 16.10. I have a question regarding the volatile usage with the global variable type. Suppose I have a kernel definition as follows:

__kernel void foo(volatile __global int *buffer)

I am interested to know that when a variable in the global space declared as volatile then what happens with respect to the caching at LLC and L3 level. As mentioned in the gen9 architecture manual that L3 and LLC are used for caching all the global memory access. But I was wondering that if a variable is declared as volatile then does it get cached only at LLC and not L3 or it doesn't get cached at any level or something else. If someone could provide me some information regarding how different level of caching works with respect to volatile variables then it would greatly appreciated. Thank you.

[Ben_A_Intel]

Short answer: "volatile" affects how the compiler generates code, but does not affect caching, at least for your Kabylake GPU.

If you're trying to access memory concurrently from multiple work items (or even the host and multiple work items) then I'd recommend using the OpenCL 2.0 atomics instead, which are very similar to C11 atomics.  There's a good discussion of volatile vs. atomics here:

https://stackoverflow.com/questions/8819095/concurrency-atomic-and-volatile-in-c11-memory-model