[Test][SYCL] Add test to verify SPIR kernel argument names (#5819)

This is a follow-up to PR#5772. This test verifies that
user specified names are retained in SPIR kernel argument.
Please note that the names were already retained (prior to
PR#5772) when kernel was specified as a functor. PR#5722 added
code to ensure the same behavior when kernel is specified using
lambda.

Signed-off-by: Elizabeth Andrews <[email protected]>

Author: elizabethandrews

clang/test/CodeGenSYCL/save-user-names.cpp

@@ -0,0 +1,74 @@
+// RUN:  %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+
+// Test to verify that user specified names are retained in SPIR
+// kernel argument names.
+
+#include "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