[SYCL][CUDA] Add unit test for local arguments alignment (#608)

This issue was solved in intel/llvm#5113, local kernel arguments have to
be aligned to the type size.

Author: npmiller

SYCL/Regression/local-arg-align.cpp

@@ -0,0 +1,66 @@
+// 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;
+}