[SYCL] Add testing for arrays as free function kernel parameters

clang/test/CodeGenSYCL/free_function_int_header.cpp

@@ -76,6 +76,26 @@ __attribute__((sycl_device))
 
 template void ff_6(Agg S1, Derived S2, int);
 
+constexpr int TestArrSize = 3;
+
+template <int ArrSize>
+struct KArgWithPtrArray {
+  int *data[ArrSize];
+  int start[ArrSize];
+  int end[ArrSize];
+  constexpr int getArrSize() { return ArrSize; }
+};
+
+template <int ArrSize>
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
+void ff_7(KArgWithPtrArray<ArrSize> KArg) {
+  for (int j = 0; j < ArrSize; j++)
+    for (int i = KArg.start[j]; i <= KArg.end[j]; i++)
+      KArg.data[j][i] = KArg.start[j] + KArg.end[j];
+}
+
+template void ff_7(KArgWithPtrArray<TestArrSize> KArg);
+
 // CHECK:      const char* const kernel_names[] = {
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_2Piii
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_2Piiii
@@ -84,6 +104,7 @@ template void ff_6(Agg S1, Derived S2, int);
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_3IdEvPT_S0_S0_
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_410NoPointers8Pointers3Agg
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_6I3Agg7DerivedEvT_T0_i
+// CHECK-NEXT:   {{.*}}__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE
 // CHECK-NEXT:   ""
 // CHECK-NEXT: };
 
@@ -124,6 +145,9 @@ template void ff_6(Agg S1, Derived S2, int);
 // CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 40, 32 },
 // CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 72 },
 
+// CHECK:  //--- _Z18__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 48, 0 },
+
 // CHECK:        { kernel_param_kind_t::kind_invalid, -987654321, -987654321 },
 // CHECK-NEXT: };
 
@@ -249,6 +273,26 @@ template void ff_6(Agg S1, Derived S2, int);
 // CHECK-NEXT:   static constexpr bool value = true;
 // CHECK-NEXT: };
 // CHECK-NEXT: }
+//
+// CHECK: Definition of _Z18__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE as a free function kernel
+
+// CHECK: Forward declarations of kernel and its argument types:
+// CHECK: template <int ArrSize> struct KArgWithPtrArray;
+//
+// CHECK: template <int ArrSize> void ff_7(KArgWithPtrArray<ArrSize> KArg);
+// CHECK-NEXT: static constexpr auto __sycl_shim8() {
+// CHECK-NEXT:   return (void (*)(struct KArgWithPtrArray<3>))ff_7<3>;
+// CHECK-NEXT: }
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: template <>
+// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim8()> {
+// CHECK-NEXT:   static constexpr bool value = true;
+// CHECK-NEXT: };
+// CHECK-NEXT: template <>
+// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim8()> {
+// CHECK-NEXT:   static constexpr bool value = true;
+// CHECK-NEXT: };
+// CHECK-NEXT: }
 
 // CHECK: #include <sycl/kernel_bundle.hpp>
 
@@ -307,3 +351,11 @@ template void ff_6(Agg S1, Derived S2, int);
 // CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_6I3Agg7DerivedEvT_T0_i"});
 // CHECK-NEXT: }
 // CHECK-NEXT: }
+
+// CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: template <>
+// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim8()>() {
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE"});
+// CHECK-NEXT: }
+// CHECK-NEXT: }

clang/test/CodeGenSYCL/free_function_kernel_params.cpp

@@ -26,10 +26,33 @@ __attribute__((sycl_device))
 void ff_4(NoPointers S1, Pointers S2, Agg S3) {
 }
 
+constexpr int TestArrSize = 3;
+
+template <int ArrSize>
+struct KArgWithPtrArray {
+  int *data[ArrSize];
+  int start[ArrSize];
+  int end[ArrSize];
+  constexpr int getArrSize() { return ArrSize; }
+};
+
+template <int ArrSize>
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
+void ff_6(KArgWithPtrArray<ArrSize> KArg) {
+  for (int j = 0; j < ArrSize; j++)
+    for (int i = KArg.start[j]; i <= KArg.end[j]; i++)
+      KArg.data[j][i] = KArg.start[j] + KArg.end[j];
+}
+
+template void ff_6(KArgWithPtrArray<TestArrSize> KArg);
+
 // CHECK: %struct.NoPointers = type { i32 }
 // CHECK: %struct.Pointers = type { ptr addrspace(4), ptr addrspace(4) }
 // CHECK: %struct.Agg = type { %struct.NoPointers, i32, ptr addrspace(4), %struct.Pointers }
 // CHECK: %struct.__generated_Pointers = type { ptr addrspace(1), ptr addrspace(1) }
 // CHECK: %struct.__generated_Agg = type { %struct.NoPointers, i32, ptr addrspace(1), %struct.__generated_Pointers.0 }
 // CHECK: %struct.__generated_Pointers.0 = type { ptr addrspace(1), ptr addrspace(1) }
+// CHECK: %struct.__generated_KArgWithPtrArray = type { [3 x ptr addrspace(1)], [3 x i32], [3 x i32] }
+// CHECK: %struct.KArgWithPtrArray = type { [3 x ptr addrspace(4)], [3 x i32], [3 x i32] }
 // CHECK: define dso_local spir_kernel void @{{.*}}__sycl_kernel{{.*}}(ptr noundef byval(%struct.NoPointers) align 4 %__arg_S1, ptr noundef byval(%struct.__generated_Pointers) align 8 %__arg_S2, ptr noundef byval(%struct.__generated_Agg) align 8 %__arg_S3)
+// CHECK: define dso_local spir_kernel void @{{.*}}__sycl_kernel_ff_6{{.*}}(ptr noundef byval(%struct.__generated_KArgWithPtrArray) align 8 %__arg_KArg)

clang/test/SemaSYCL/free_function_array_kernel_param.cpp

@@ -0,0 +1,40 @@
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -ast-dump \
+// RUN: %s -o - | FileCheck %s
+
+// This test checks parameter rewriting for free functions with parameters
+// of type struct with array and array of pointers.
+
+#include "sycl.hpp"
+
+constexpr int TestArrSize = 3;
+
+template <int ArrSize>
+struct KArgWithPtrArray {
+  int *data[ArrSize];
+  int start[ArrSize];
+  int end[ArrSize];
+  constexpr int getArrSize() { return ArrSize; }
+};
+
+template <int ArrSize>
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
+void ff_6(KArgWithPtrArray<ArrSize> KArg) {
+  for (int j = 0; j < ArrSize; j++)
+    for (int i = KArg.start[j]; i <= KArg.end[j]; i++)
+      KArg.data[j][i] = KArg.start[j] + KArg.end[j];
+}
+
+template void ff_6(KArgWithPtrArray<TestArrSize> KArg);
+
+// CHECK: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (__generated_KArgWithPtrArray)'
+// CHECK-NEXT: ParmVarDecl {{.*}} used __arg_KArg '__generated_KArgWithPtrArray'
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: CallExpr
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'void (*)(KArgWithPtrArray<3>)' <FunctionToPointerDecay>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'void (KArgWithPtrArray<3>)' lvalue Function {{.*}} 'ff_6' 'void (KArgWithPtrArray<3>)'
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'KArgWithPtrArray<3>' 'void (const KArgWithPtrArray<3> &) noexcept'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const KArgWithPtrArray<3>' lvalue <NoOp>
+// CHECK-NEXT: UnaryOperator {{.*}} 'KArgWithPtrArray<3>' lvalue prefix '*' cannot overflow
+// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'KArgWithPtrArray<3> *' reinterpret_cast<KArgWithPtrArray<3> *> <BitCast>
+// CHECK-NEXT: UnaryOperator {{.*}} '__generated_KArgWithPtrArray *' prefix '&' cannot overflow
+// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_KArgWithPtrArray' lvalue ParmVar {{.*}} '__arg_KArg' '__generated_KArgWithPtrArray'

sycl/test-e2e/KernelAndProgram/free_function_kernels.cpp

@@ -371,6 +371,62 @@ bool test_5(queue Queue) {
   return PassA && PassB;
 }
 
+constexpr int TestArrSize = 3;
+
+template <int ArrSize>
+struct KArgWithPtrArray {
+  int *data[ArrSize];
+  int start[ArrSize];
+  int end[ArrSize];
+  constexpr int getArrSize() { return ArrSize; }
+};
+
+template <int ArrSize>
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
+    (ext::oneapi::experimental::single_task_kernel))
+void ff_6(KArgWithPtrArray<ArrSize> KArg) {
+  for (int j = 0; j < ArrSize; j++)
+    for (int i = KArg.start[j]; i <= KArg.end[j]; i++)
+      KArg.data[j][i] = KArg.start[j] + KArg.end[j];
+}
+
+template void ff_6(KArgWithPtrArray<TestArrSize> KArg);
+
+bool test_6(queue Queue) {
+  constexpr int Range = 10;
+  KArgWithPtrArray<TestArrSize> KArg;
+  for (int i = 0; i < TestArrSize; ++i) {
+    KArg.data[i] = malloc_shared<int>(Range, Queue);
+    memset(KArg.data[i], 0, Range * sizeof(int));
+    KArg.start[i]= 3;
+    KArg.end[i] = 5;
+  }
+  int Result[Range] = {0, 0, 0, 8, 8, 8, 0, 0, 0, 0};
+  range<1> R1{Range};
+
+  bool Pass = true;
+#ifndef __SYCL_DEVICE_ONLY__
+  kernel_bundle Bundle =
+      get_kernel_bundle<bundle_state::executable>(Queue.get_context());
+  kernel_id Kernel_id = ext::oneapi::experimental::get_kernel_id<(
+      void (*)(KArgWithPtrArray<TestArrSize>))ff_6<TestArrSize>>();
+  kernel Kernel = Bundle.get_kernel(Kernel_id);
+  Queue.submit([&](handler &Handler) {
+    Handler.set_arg(0, KArg);
+    Handler.single_task(Kernel);
+  });
+  Queue.wait();
+  for (int i = 0; i < TestArrSize; ++i) {
+    Pass &= checkUSM(KArg.data[i], Range, Result);
+    std::cout << "Test 6, array: " << i << (Pass ? " PASS" : " FAIL")
+              << std::endl;
+    free(KArg.data[i], Queue);
+  }
+
+#endif
+  return Pass;
+}
+
 int main() {
   queue Queue;
 
@@ -381,6 +437,7 @@ int main() {
   Pass &= test_3(Queue);
   Pass &= test_4(Queue);
   Pass &= test_5(Queue);
+  Pass &= test_6(Queue);
 
   return Pass ? 0 : 1;
 }