Fix local memory out of bounds issue in atomic_fence (replaces PR #1285) (#1437)

* Fix local memory out of bounds in atomic_fence

In the error condition, the atomic_fence kernel can illegally access local memory addresses.

In this snippet, localValues is in the local address space and provided as a kernel argument. Its size is effectively get_local_size(0) * sizeof(int). The stores to localValues lead to OoB accesses.

  size_t myId = get_local_id(0);

  ...

  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;
  }

* Fix formatting

* Fix formatting again

* Formatting

test_conformance/c11_atomics/common.h

@@ -1360,8 +1360,10 @@ int CBasicTest<HostAtomicType, HostDataType>::ExecuteSingleTest(
         {
             error =
                 clSetKernelArg(kernel, argInd++,
-                               LocalRefValues() ? typeSize * CurrentGroupSize()
-                                       * NumNonAtomicVariablesPerThread()
+                               LocalRefValues() ? typeSize
+                                       * ((CurrentGroupSize()
+                                           * NumNonAtomicVariablesPerThread())
+                                          + 4)
                                                 : 1,
                                NULL);
             test_error(error, "Unable to set indexed kernel argument");