OpenCL* for CPU
Ask questions and share information on Intel® SDK for OpenCL™ Applications and OpenCL™ implementations for Intel® CPU.
This forum covers OpenCL* for CPU only. OpenCL* for GPU questions can be asked in the GPU Compute Software forum. Intel® FPGA SDK for OpenCL™ questions can be ask in the FPGA Intel® High Level Design forum.
1718 Discussions

Caching and volatile variable with intel integrated GPU


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.


0 Kudos
1 Reply

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:

0 Kudos