From c2aca7d8e6a6ec2162a1c68b127409aa9931974d Mon Sep 17 00:00:00 2001
From: paulfradgley <39525348+paulfradgley@users.noreply.github.com>
Date: Tue, 14 Jun 2022 16:47:06 +0100
Subject: [PATCH] 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 | 6 ++++--
 1 file changed, 4 insertions(+), 2 deletions(-)

diff --git a/test_conformance/c11_atomics/common.h b/test_conformance/c11_atomics/common.h
index 42fe32b6..5bb9e5b7 100644
--- a/test_conformance/c11_atomics/common.h
+++ b/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");