C11 Atomics: Race condition in atomic_fence

We have identified an issue in the atomic_fence test - specifically `release atomic synchronizes-with acquire fence, local atomic_int, scope_work_group` in the case described here.

@kpet pointed out [Internal Bugzilla 15814](https://www.khronos.org/members/login/bugzilla/show_bug.cgi?id=15814) which appears to be the same underlying issue - that instances _may_ make independent forward progress.

The high level summary is that we think that it is possible for the test to fail if interleaving threads execute their instructions in a particular order which is valid, but leads to an incorrect test failure.

Suppose we have 4 instances 0, 1, 2, and 3 which can each make independent forward progress (with a workgroup size of 4).

The following timeline is possible showing how multiple instances accessing the same address races (showing operations from instances 0, 1, and 2):
```C
localValues[myId*1+myValue] = myValue; // Instance 0 writes localValues[1] = 1
atomic_store_explicit(&destMemory[myId], myValue, memory_order_release, memory_scope_work_group); // Instance 0 writes destMemory[0] = 1
hisAtomicValue = atomic_load_explicit(&destMemory[hisId], memory_order_relaxed, memory_scope_work_group); // Instance 0 reads destMemory[3] which is 0 as instance 3 may not have written it yet
localValues[myId*1+myValue] = myValue; // Instance 2 writes localValues[3] = 1.
atomic_work_item_fence(CLK_LOCAL_MEM_FENCE, memory_order_acquire, memory_scope_work_group); // Instance 0
hisValue = localValues[hisId*1+hisAtomicValue]; // Instance 0 Reads localValues[3] which is 1
if(hisAtomicValue != hisValue) // Instance 0 - At this point hisAtomicValue != hisValue so test failure
```
We do not yet have a proposal for fixing the kernel. Instead, looking to spark some discussion around the issue and test itself so that we can come up with a proposal to fix (whilst maintaining the original intent).

For reference, the entire kernel is:
```C
__kernel void test_atomic_kernel(uint threadCount, uint numDestItems, __global int *finalDest, __global int *oldValues, volatile __local atomic_int *destMemory, __local int *localValues)
{
uint tid = get_global_id(0);

// initialize atomics not reachable from host (first thread
// is doing this, other threads are waiting on barrier)
if (get_local_id(0) == 0)
for (uint dstItemIdx = 0; dstItemIdx < numDestItems; dstItemIdx++)
{
atomic_store_explicit(destMemory + dstItemIdx,
finalDest[dstItemIdx],
memory_order_relaxed,
memory_scope_work_group);
}
barrier(CLK_LOCAL_MEM_FENCE); //Redundant

// Copy input reference values into local memory
localValues[get_local_id(0)] = oldValues[tid];
barrier(CLK_LOCAL_MEM_FENCE);

size_t myId = get_local_id(0), hisId = get_local_size(0) - 1 - myId;
int myValue = 0, hisAtomicValue, hisValue;
do
{
myValue++;
localValues[myId * 1 + myValue] = myValue;
atomic_store_explicit(&destMemory[myId], myValue, memory_order_release, memory_scope_work_group);
hisAtomicValue = atomic_load_explicit(&destMemory[hisId], memory_order_relaxed, memory_scope_work_group);
atomic_work_item_fence(CLK_LOCAL_MEM_FENCE, memory_order_acquire, memory_scope_work_group);
hisValue = localValues[hisId * 1 + hisAtomicValue];
hisId = (hisId + 1) % get_local_size(0);
} while (hisAtomicValue == hisValue && myValue < 1 - 1); // Always false, never loops
if (hisAtomicValue != hisValue)
{ // fail
atomic_store(&destMemory[myId], myValue - 1);
hisId = (hisId + get_local_size(0) - 1) % get_local_size(0);
if (myValue + 1 < 1)
localValues[myId * 1 + myValue + 1] = hisId;
if (myValue + 2 < 1)
localValues[myId * 1 + myValue + 2] = hisAtomicValue;
if (myValue + 3 < 1)
localValues[myId * 1 + myValue + 3] = hisValue;
}

// Copy local reference values into output array
barrier(CLK_LOCAL_MEM_FENCE);
oldValues[tid] = localValues[get_local_id(0)];

// Copy final values to host reachable buffer
barrier(CLK_LOCAL_MEM_FENCE); //Redundant
if (get_local_id(0) == 0) // first thread in workgroup
for (uint dstItemIdx = 0; dstItemIdx < numDestItems; dstItemIdx++)

finalDest[dstItemIdx] =
atomic_load_explicit(destMemory + dstItemIdx,
memory_order_relaxed,
memory_scope_work_group);
}
```
There is another issue with OoB local memory access, but will directly file a PR for that as its trivial.

Start with the atomic_fence test and the supplied test_atomic_kernel, then trace the described interleaving among independently progressing work-items. Review the referenced Bugzilla issue and discussion to determine whether the failure is valid. Done means an agreed proposal that fixes the test while preserving its original intent; the separate out-of-bounds issue is already intended for a PR.

Written by the indexing model from the issue text.