Skip to content

Commit 4af0c1b

Browse files
FznamznonFznamznon
authored
[SYCL] Add testing for arrays as free function kernel parameters (#15411)
Arrays cannot be passed by value as function parameters, so they won’t be supported as free function parameters. To pass an array the simplest way is to wrap it in a struct and pass the struct. This PR adds tests for the case
1 parent ddd23ad commit 4af0c1b

File tree

4 files changed

+172
-0
lines changed

4 files changed

+172
-0
lines changed

‎clang/test/CodeGenSYCL/free_function_int_header.cpp‎

Lines changed: 52 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -76,6 +76,26 @@ __attribute__((sycl_device))
7676

7777
template void ff_6(Agg S1, Derived S2, int);
7878

79+
constexpr int TestArrSize = 3;
80+
81+
template <int ArrSize>
82+
struct KArgWithPtrArray {
83+
int *data[ArrSize];
84+
int start[ArrSize];
85+
int end[ArrSize];
86+
constexpr int getArrSize() { return ArrSize; }
87+
};
88+
89+
template <int ArrSize>
90+
[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
91+
void ff_7(KArgWithPtrArray<ArrSize> KArg) {
92+
for (int j = 0; j < ArrSize; j++)
93+
for (int i = KArg.start[j]; i <= KArg.end[j]; i++)
94+
KArg.data[j][i] = KArg.start[j] + KArg.end[j];
95+
}
96+
97+
template void ff_7(KArgWithPtrArray<TestArrSize> KArg);
98+
7999
// CHECK: const char* const kernel_names[] = {
80100
// CHECK-NEXT: {{.*}}__sycl_kernel_ff_2Piii
81101
// CHECK-NEXT: {{.*}}__sycl_kernel_ff_2Piiii
@@ -84,6 +104,7 @@ template void ff_6(Agg S1, Derived S2, int);
84104
// CHECK-NEXT: {{.*}}__sycl_kernel_ff_3IdEvPT_S0_S0_
85105
// CHECK-NEXT: {{.*}}__sycl_kernel_ff_410NoPointers8Pointers3Agg
86106
// CHECK-NEXT: {{.*}}__sycl_kernel_ff_6I3Agg7DerivedEvT_T0_i
107+
// CHECK-NEXT: {{.*}}__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE
87108
// CHECK-NEXT: ""
88109
// CHECK-NEXT: };
89110

@@ -124,6 +145,9 @@ template void ff_6(Agg S1, Derived S2, int);
124145
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 40, 32 },
125146
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 72 },
126147

148+
// CHECK: //--- _Z18__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE
149+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 48, 0 },
150+
127151
// CHECK: { kernel_param_kind_t::kind_invalid, -987654321, -987654321 },
128152
// CHECK-NEXT: };
129153

@@ -249,6 +273,26 @@ template void ff_6(Agg S1, Derived S2, int);
249273
// CHECK-NEXT: static constexpr bool value = true;
250274
// CHECK-NEXT: };
251275
// CHECK-NEXT: }
276+
//
277+
// CHECK: Definition of _Z18__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE as a free function kernel
278+
279+
// CHECK: Forward declarations of kernel and its argument types:
280+
// CHECK: template <int ArrSize> struct KArgWithPtrArray;
281+
//
282+
// CHECK: template <int ArrSize> void ff_7(KArgWithPtrArray<ArrSize> KArg);
283+
// CHECK-NEXT: static constexpr auto __sycl_shim8() {
284+
// CHECK-NEXT: return (void (*)(struct KArgWithPtrArray<3>))ff_7<3>;
285+
// CHECK-NEXT: }
286+
// CHECK-NEXT: namespace sycl {
287+
// CHECK-NEXT: template <>
288+
// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim8()> {
289+
// CHECK-NEXT: static constexpr bool value = true;
290+
// CHECK-NEXT: };
291+
// CHECK-NEXT: template <>
292+
// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim8()> {
293+
// CHECK-NEXT: static constexpr bool value = true;
294+
// CHECK-NEXT: };
295+
// CHECK-NEXT: }
252296

253297
// CHECK: #include <sycl/kernel_bundle.hpp>
254298

@@ -307,3 +351,11 @@ template void ff_6(Agg S1, Derived S2, int);
307351
// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_6I3Agg7DerivedEvT_T0_i"});
308352
// CHECK-NEXT: }
309353
// CHECK-NEXT: }
354+
355+
// CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE
356+
// CHECK-NEXT: namespace sycl {
357+
// CHECK-NEXT: template <>
358+
// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim8()>() {
359+
// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE"});
360+
// CHECK-NEXT: }
361+
// CHECK-NEXT: }

‎clang/test/CodeGenSYCL/free_function_kernel_params.cpp‎

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -26,10 +26,33 @@ __attribute__((sycl_device))
2626
void ff_4(NoPointers S1, Pointers S2, Agg S3) {
2727
}
2828

29+
constexpr int TestArrSize = 3;
30+
31+
template <int ArrSize>
32+
struct KArgWithPtrArray {
33+
int *data[ArrSize];
34+
int start[ArrSize];
35+
int end[ArrSize];
36+
constexpr int getArrSize() { return ArrSize; }
37+
};
38+
39+
template <int ArrSize>
40+
[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
41+
void ff_6(KArgWithPtrArray<ArrSize> KArg) {
42+
for (int j = 0; j < ArrSize; j++)
43+
for (int i = KArg.start[j]; i <= KArg.end[j]; i++)
44+
KArg.data[j][i] = KArg.start[j] + KArg.end[j];
45+
}
46+
47+
template void ff_6(KArgWithPtrArray<TestArrSize> KArg);
48+
2949
// CHECK: %struct.NoPointers = type { i32 }
3050
// CHECK: %struct.Pointers = type { ptr addrspace(4), ptr addrspace(4) }
3151
// CHECK: %struct.Agg = type { %struct.NoPointers, i32, ptr addrspace(4), %struct.Pointers }
3252
// CHECK: %struct.__generated_Pointers = type { ptr addrspace(1), ptr addrspace(1) }
3353
// CHECK: %struct.__generated_Agg = type { %struct.NoPointers, i32, ptr addrspace(1), %struct.__generated_Pointers.0 }
3454
// CHECK: %struct.__generated_Pointers.0 = type { ptr addrspace(1), ptr addrspace(1) }
55+
// CHECK: %struct.__generated_KArgWithPtrArray = type { [3 x ptr addrspace(1)], [3 x i32], [3 x i32] }
56+
// CHECK: %struct.KArgWithPtrArray = type { [3 x ptr addrspace(4)], [3 x i32], [3 x i32] }
3557
// 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)
58+
// 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‎

Lines changed: 40 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,40 @@
1+
// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -ast-dump \
2+
// RUN: %s -o - | FileCheck %s
3+
4+
// This test checks parameter rewriting for free functions with parameters
5+
// of type struct with array and array of pointers.
6+
7+
#include "sycl.hpp"
8+
9+
constexpr int TestArrSize = 3;
10+
11+
template <int ArrSize>
12+
struct KArgWithPtrArray {
13+
int *data[ArrSize];
14+
int start[ArrSize];
15+
int end[ArrSize];
16+
constexpr int getArrSize() { return ArrSize; }
17+
};
18+
19+
template <int ArrSize>
20+
[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
21+
void ff_6(KArgWithPtrArray<ArrSize> KArg) {
22+
for (int j = 0; j < ArrSize; j++)
23+
for (int i = KArg.start[j]; i <= KArg.end[j]; i++)
24+
KArg.data[j][i] = KArg.start[j] + KArg.end[j];
25+
}
26+
27+
template void ff_6(KArgWithPtrArray<TestArrSize> KArg);
28+
29+
// CHECK: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (__generated_KArgWithPtrArray)'
30+
// CHECK-NEXT: ParmVarDecl {{.*}} used __arg_KArg '__generated_KArgWithPtrArray'
31+
// CHECK-NEXT: CompoundStmt
32+
// CHECK-NEXT: CallExpr
33+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'void (*)(KArgWithPtrArray<3>)' <FunctionToPointerDecay>
34+
// CHECK-NEXT: DeclRefExpr {{.*}} 'void (KArgWithPtrArray<3>)' lvalue Function {{.*}} 'ff_6' 'void (KArgWithPtrArray<3>)'
35+
// CHECK-NEXT: CXXConstructExpr {{.*}} 'KArgWithPtrArray<3>' 'void (const KArgWithPtrArray<3> &) noexcept'
36+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const KArgWithPtrArray<3>' lvalue <NoOp>
37+
// CHECK-NEXT: UnaryOperator {{.*}} 'KArgWithPtrArray<3>' lvalue prefix '*' cannot overflow
38+
// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'KArgWithPtrArray<3> *' reinterpret_cast<KArgWithPtrArray<3> *> <BitCast>
39+
// CHECK-NEXT: UnaryOperator {{.*}} '__generated_KArgWithPtrArray *' prefix '&' cannot overflow
40+
// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_KArgWithPtrArray' lvalue ParmVar {{.*}} '__arg_KArg' '__generated_KArgWithPtrArray'

‎sycl/test-e2e/KernelAndProgram/free_function_kernels.cpp‎

Lines changed: 57 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -371,6 +371,62 @@ bool test_5(queue Queue) {
371371
return PassA && PassB;
372372
}
373373

374+
constexpr int TestArrSize = 3;
375+
376+
template <int ArrSize>
377+
struct KArgWithPtrArray {
378+
int *data[ArrSize];
379+
int start[ArrSize];
380+
int end[ArrSize];
381+
constexpr int getArrSize() { return ArrSize; }
382+
};
383+
384+
template <int ArrSize>
385+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
386+
(ext::oneapi::experimental::single_task_kernel))
387+
void ff_6(KArgWithPtrArray<ArrSize> KArg) {
388+
for (int j = 0; j < ArrSize; j++)
389+
for (int i = KArg.start[j]; i <= KArg.end[j]; i++)
390+
KArg.data[j][i] = KArg.start[j] + KArg.end[j];
391+
}
392+
393+
template void ff_6(KArgWithPtrArray<TestArrSize> KArg);
394+
395+
bool test_6(queue Queue) {
396+
constexpr int Range = 10;
397+
KArgWithPtrArray<TestArrSize> KArg;
398+
for (int i = 0; i < TestArrSize; ++i) {
399+
KArg.data[i] = malloc_shared<int>(Range, Queue);
400+
memset(KArg.data[i], 0, Range * sizeof(int));
401+
KArg.start[i]= 3;
402+
KArg.end[i] = 5;
403+
}
404+
int Result[Range] = {0, 0, 0, 8, 8, 8, 0, 0, 0, 0};
405+
range<1> R1{Range};
406+
407+
bool Pass = true;
408+
#ifndef __SYCL_DEVICE_ONLY__
409+
kernel_bundle Bundle =
410+
get_kernel_bundle<bundle_state::executable>(Queue.get_context());
411+
kernel_id Kernel_id = ext::oneapi::experimental::get_kernel_id<(
412+
void (*)(KArgWithPtrArray<TestArrSize>))ff_6<TestArrSize>>();
413+
kernel Kernel = Bundle.get_kernel(Kernel_id);
414+
Queue.submit([&](handler &Handler) {
415+
Handler.set_arg(0, KArg);
416+
Handler.single_task(Kernel);
417+
});
418+
Queue.wait();
419+
for (int i = 0; i < TestArrSize; ++i) {
420+
Pass &= checkUSM(KArg.data[i], Range, Result);
421+
std::cout << "Test 6, array: " << i << (Pass ? " PASS" : " FAIL")
422+
<< std::endl;
423+
free(KArg.data[i], Queue);
424+
}
425+
426+
#endif
427+
return Pass;
428+
}
429+
374430
int main() {
375431
queue Queue;
376432

@@ -381,6 +437,7 @@ int main() {
381437
Pass &= test_3(Queue);
382438
Pass &= test_4(Queue);
383439
Pass &= test_5(Queue);
440+
Pass &= test_6(Queue);
384441

385442
return Pass ? 0 : 1;
386443
}

0 commit comments

Comments
 (0)