False positive "Read-write" data race with barrier-protected global memory access

[gabriele-costa]

I am running the following (minimal) example

int sread(__global unsigned int *L) {
  barrier(CLK_GLOBAL_MEM_FENCE);
  int r = *L;
  barrier(CLK_GLOBAL_MEM_FENCE);
  return r;
}

__kernel void manager(__global unsigned int *L) {
  int val = sread(L);
  while(val < 0xFFFF) { 
    val = sread(L);
    *L = val + 1;
  }
}

Using the --data-races option. I obtain the following result

Read-write data race at global memory address 0x1000000000000
	Kernel: manager
	
	First entity:  Global(1,0,0) Local(0,0,0) Group(1,0,0)
	  %2 = load i32, i32 addrspace(1)* %1, align 4, !dbg !16
	At line 6 (column 11) of input.cl:
	  int r = *L;
	
	Second entity: Global(0,0,0) Local(0,0,0) Group(0,0,0)
	  store i32 %add, i32 addrspace(1)* %7, align 4, !dbg !31
	At line 15 (column 8) of input.cl:
	  *L = val + 1;
	

Read-write data race at global memory address 0x1000000000000
	Kernel: manager
	
	First entity:  Global(1,0,0) Local(0,0,0) Group(1,0,0)
	  store i32 %add, i32 addrspace(1)* %7, align 4, !dbg !31
	At line 15 (column 8) of input.cl:
	  *L = val + 1;
	
	Second entity: Global(0,0,0) Local(0,0,0) Group(0,0,0)
	  %2 = load i32, i32 addrspace(1)* %1, align 4, !dbg !16
	At line 6 (column 11) of input.cl:
	  int r = *L;

However, according to my understanding of the barrier mechanism, this should not happen since the function sread "protects" the read operation with two barriers. Am I missing something or is this a bug?

Thank you

PS. Code in attachment

issue-min.zip

[jrprice]

The name of the barrier function in OpenCL 1.2 is a little misleading, since it only operates within a work-group (in OpenCL 2.0 it was renamed to work_group_barrier to try and make this clearer).

In OpenCL 1.x, there is no mechanism to synchronise global memory access between different work-groups. You can see from the data-race messages that Oclgrind is producing that there are two different work-groups involved in the race.

[gabriele-costa]

Thank you, that is exactly what was going on. Actually I was misled by the flag CLK_GLOBAL_MEM_FENCE that, implicitly, made me think that barriers can prevent data races in global memory.