From b0f099fbffc6095a615062628a943addd214964f Mon Sep 17 00:00:00 2001
From: Nicolas Miller <nicolas.miller@codeplay.com>
Date: Fri, 10 Dec 2021 17:46:48 +0000
Subject: [PATCH 1/2] [SYCL][CUDA] Add unit test for local arguments alignment

This issue was solved in intel/llvm#5113, local kernel arguments have to
be aligned to the type size.
---
 SYCL/Regression/local-arg-align.cpp | 64 +++++++++++++++++++++++++++++
 1 file changed, 64 insertions(+)
 create mode 100644 SYCL/Regression/local-arg-align.cpp

diff --git a/SYCL/Regression/local-arg-align.cpp b/SYCL/Regression/local-arg-align.cpp
new file mode 100644
index 0000000000..2a5ce5ca70
--- /dev/null
+++ b/SYCL/Regression/local-arg-align.cpp
@@ -0,0 +1,64 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+//
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+//==-- local-arg-align.cpp - Test for local argument alignmnent ------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <CL/sycl.hpp>
+#include <iostream>
+
+using namespace sycl;
+
+// This test is a simple unit test to ensure that local kernel arguments are
+// properly aligned.
+int main(int argc, char *argv[]) {
+  queue q;
+  buffer<size_t, 1> res(2);
+
+  q.submit([&](sycl::handler &h) {
+     // Use two local buffers, one with an int and one with a double4
+     accessor<cl_int, 1, access::mode::read_write, access::target::local> a(1, h);
+     accessor<double4, 1, access::mode::read_write, access::target::local> b(1, h);
+
+     auto ares = res.get_access<access::mode::read_write>(h);
+
+     // Manually capture kernel arguments to ensure an order with the int
+     // argument first and the double4 argument second. If the two arguments are
+     // simply laid out consecutively, the double4 argument will not be
+     // correctly aligned.
+     h.parallel_for(1, [a, b, ares](sycl::id<1> i) {
+       // Get the addresses of the two local buffers
+       ares[0] = (size_t)&a[0];
+       ares[1] = (size_t)&b[0];
+     });
+   }).wait_and_throw();
+
+  auto hres = res.get_access<access::mode::read_write>();
+
+  int ret = 0;
+  // Check that the addresses are aligned as expected
+  if (hres[0] % sizeof(cl_int) != 0) {
+    std::cout
+        << "Error: incorrect alignment for argument a, required alignment: "
+        << sizeof(cl_int) << ", address: " << (void *)hres[0] << std::endl;
+    ret = -1;
+  }
+
+  if (hres[1] % sizeof(double4) != 0) {
+    std::cout
+        << "Error: incorrect alignment for argument b, required alignment: "
+        << sizeof(double4) << ", address: " << (void *)hres[1] << std::endl;
+    ret = -1;
+  }
+
+  return ret;
+}

From c8a8c4cae686ba3b17c55a2f7f62a7f09e9e5c4f Mon Sep 17 00:00:00 2001
From: Nicolas Miller <nicolas.miller@codeplay.com>
Date: Fri, 10 Dec 2021 18:28:17 +0000
Subject: [PATCH 2/2] [SYCL][CUDA] Fix formatting

---
 SYCL/Regression/local-arg-align.cpp | 6 ++++--
 1 file changed, 4 insertions(+), 2 deletions(-)

diff --git a/SYCL/Regression/local-arg-align.cpp b/SYCL/Regression/local-arg-align.cpp
index 2a5ce5ca70..9cc07785b2 100644
--- a/SYCL/Regression/local-arg-align.cpp
+++ b/SYCL/Regression/local-arg-align.cpp
@@ -26,8 +26,10 @@ int main(int argc, char *argv[]) {
 
   q.submit([&](sycl::handler &h) {
      // Use two local buffers, one with an int and one with a double4
-     accessor<cl_int, 1, access::mode::read_write, access::target::local> a(1, h);
-     accessor<double4, 1, access::mode::read_write, access::target::local> b(1, h);
+     accessor<cl_int, 1, access::mode::read_write, access::target::local> a(1,
+                                                                            h);
+     accessor<double4, 1, access::mode::read_write, access::target::local> b(1,
+                                                                             h);
 
      auto ares = res.get_access<access::mode::read_write>(h);