[SYCL][Test] Add integration test to verify SPIR kernel argument names

sycl/test/CMakeLists.txt

@@ -34,6 +34,7 @@ list(APPEND SYCL_TEST_DEPS
   get_device_count_by_type
   llvm-config
   llvm-cxxdump
+  llvm-dis
   llvm-readobj
   )
 

sycl/test/basic_tests/sycl-kernel-save-user-names.cpp

@@ -0,0 +1,80 @@
+// RUN: %clangxx -fsycl -fsycl-device-only -fno-sycl-early-optimizations -o %t.bc %s
+// RUN: sycl-post-link %t.bc -spec-const=default -o %t.table
+// RUN: llvm-spirv -o %t.spv -spirv-max-version=1.3 -spirv-ext=+all %t.bc
+// RUN: llvm-spirv -o %t.rev.bc -r %t.spv
+// RUN: llvm-dis %t.rev.bc -o=- | FileCheck %s
+
+// Test to verify that user specified names are retained in SPIR kernel argument
+// names. (It is a copy of clang/test/CodeGenSYCL/save-user-names.cpp with just
+// additional compilation steps).
+
+#include <CL/sycl.hpp>
+
+struct NestedSimple {
+  int NestedSimpleField;
+};
+
+struct NestedComplex {
+  int NestedComplexField;
+  cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> NestedAccField;
+};
+
+struct KernelFunctor {
+  int IntField;
+  cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> AccField1;
+  cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> AccField2;
+  NestedSimple NestedSimpleObj;
+  NestedComplex NestedComplexObj;
+  void operator()() const {}
+};
+
+int main() {
+  cl::sycl::queue q;
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    KernelFunctor FunctorObj;
+    cgh.single_task<class Kernel1>(FunctorObj);
+  });
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    int Data;
+    NestedSimple NestedSimpleObj;
+    NestedComplex NestedComplexObj;
+    cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> CapturedAcc1,
+        CapturedAcc2;
+    cgh.single_task<class Kernel2>([=]() {
+      Data;
+      CapturedAcc1;
+      CapturedAcc2;
+      NestedSimpleObj;
+      NestedComplexObj;
+    });
+  });
+
+  return 0;
+}
+
+// Check kernel parameters generated when kernel is defined as Functor
+
+// NOTE: Accessor fields have 4 corresponding openCL kernel arguments. When
+// the compiler generates the openCL kernel arguments, they are generated
+// with the same name (i.e the user-specified name for accessor). Since LLVM
+// IR cannot have same names, a number is appended in IR.
+//
+// CHECK: define {{.*}}spir_kernel void @{{.*}}Kernel1
+// CHECK-SAME: %_arg_IntField
+// CHECK-SAME: %_arg_AccField1{{.*}}%_arg_AccField11{{.*}}%_arg_AccField12{{.*}}%_arg_AccField13
+// CHECK-SAME: %_arg_AccField2{{.*}}%_arg_AccField24{{.*}}%_arg_AccField25{{.*}}%_arg_AccField26
+// CHECK-SAME: %_arg_NestedSimpleObj
+// CHECK-SAME: %_arg_NestedComplexField
+// CHECK-SAME: %_arg_NestedAccField{{.*}}%_arg_NestedAccField7{{.*}}%_arg_NestedAccField8{{.*}}%_arg_NestedAccField9
+
+// Check kernel parameters generated when kernel is defined as Lambda
+//
+// CHECK: define {{.*}}spir_kernel void @{{.*}}Kernel2
+// CHECK-SAME: %_arg_Data
+// CHECK-SAME: %_arg_CapturedAcc1{{.*}}%_arg_CapturedAcc11{{.*}}%_arg_CapturedAcc12{{.*}}%_arg_CapturedAcc13
+// CHECK-SAME: %_arg_CapturedAcc2{{.*}}%_arg_CapturedAcc24{{.*}}%_arg_CapturedAcc25{{.*}}%_arg_CapturedAcc26
+// CHECK-SAME: %_arg_NestedSimpleObj
+// CHECK-SAME: %_arg_NestedComplexField
+// CHECK-SAME: %_arg_NestedAccField{{.*}}%_arg_NestedAccField7{{.*}}%_arg_NestedAccField8{{.*}}%_arg_NestedAccField9