Skip to content

Commit c2a0db0

Browse files
elizabethandrewselizabethandrews
authored
[SYCL] Allow non-conforming lambda syntax for work-group-size-hint attribute (#6976)
We missed this attribute when support for non-conforming lambda syntax was originally implemented. Adding to maintain consistency. Please note - now that clang supports attributes on lambdas, discussions are underway on removing this non-conforming syntax support for all SYCL attributes in the future. Signed-off-by: Elizabeth Andrews <[email protected]>
1 parent eb9cb9e commit c2a0db0

File tree

4 files changed

+86
-2
lines changed

4 files changed

+86
-2
lines changed

‎clang/include/clang/Basic/Attr.td

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3427,6 +3427,7 @@ def WorkGroupSizeHint : InheritableAttr {
34273427
}
34283428
}];
34293429
let Documentation = [WorkGroupSizeHintAttrDocs];
3430+
let SupportsNonconformingLambdaSyntax = 1;
34303431
}
34313432

34323433
def InitPriority : InheritableAttr, TargetSpecificAttr<TargetSupportsInitPriority> {

‎clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp

Lines changed: 36 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,7 @@
33
// Tests for IR of [[intel::scheduler_target_fmax_mhz()]], [[intel::num_simd_work_items()]],
44
// [[intel::no_global_work_offset()]], [[intel::max_global_work_dim()]], [[sycl::reqd_sub_group_size()]],
55
// [[sycl::reqd_work_group_size()]], [[intel::kernel_args_restrict]], [[intel::max_work_group_size()]],
6-
// and [[intel::sycl_explicit_simd]] function attributes in SYCL 2020.
6+
// [[sycl::work_group_size_hint()]] and [[intel::sycl_explicit_simd]] function attributes in SYCL 2020.
77

88
#include "sycl.hpp"
99

@@ -144,6 +144,19 @@ class Functor10 {
144144
}
145145
};
146146

147+
class Foo11 {
148+
public:
149+
[[sycl::work_group_size_hint(1, 2, 3)]] void operator()() const {}
150+
};
151+
152+
template <int SIZE, int SIZE1, int SIZE2>
153+
class Functor11 {
154+
public:
155+
[[sycl::work_group_size_hint(SIZE, SIZE1, SIZE2)]] void operator()() const {}
156+
};
157+
158+
[[sycl::work_group_size_hint(1, 2, 3)]] void foo11() {}
159+
147160
int main() {
148161
q.submit([&](handler &h) {
149162
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 !kernel_arg_buffer_location ![[NUM:[0-9]+]] !scheduler_target_fmax_mhz ![[NUM1:[0-9]+]]
@@ -320,6 +333,27 @@ int main() {
320333
// CHECK: define {{.*}}spir_func void @{{.*}}(ptr addrspace(4) noalias noundef align 1 dereferenceable_or_null(1) %this) #4 align 2
321334
h.single_task<class kernel_name34>(
322335
[]() [[intel::kernel_args_restrict]]{});
336+
337+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name35() #0 {{.*}} !work_group_size_hint ![[NUM123:[0-9]+]]
338+
Foo11 boo11;
339+
h.single_task<class kernel_name35>(boo11);
340+
341+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name36() #0 {{.*}} !work_group_size_hint ![[NUM123]]
342+
Functor11<1, 2, 3> f11;
343+
h.single_task<class kernel_name36>(f11);
344+
345+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name37() #0 {{.*}} !work_group_size_hint ![[NUM123]]
346+
h.single_task<class kernel_name37>(
347+
[]() [[sycl::work_group_size_hint(1, 2, 3)]]{});
348+
349+
// Test attribute is not propagated.
350+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name38()
351+
// CHECK-NOT: !work_group_size_hint
352+
// CHECK-SAME: {
353+
// CHECK: define dso_local spir_func void @_Z5foo11v()
354+
h.single_task<class kernel_name38>(
355+
[]() { foo11(); });
356+
323357
});
324358
return 0;
325359
}
@@ -333,3 +367,4 @@ int main() {
333367
// CHECK: ![[NUM32]] = !{i32 16, i32 16, i32 32}
334368
// CHECK: ![[NUM88]] = !{i32 8, i32 8, i32 8}
335369
// CHECK: ![[NUM22]] = !{i32 2, i32 2, i32 2}
370+
// CHECK: ![[NUM123]] = !{i32 1, i32 2, i32 3}

‎clang/test/CodeGenSYCL/no_opaque_check-direct-attribute-propagation.cpp

Lines changed: 36 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,7 @@
33
// Tests for IR of [[intel::scheduler_target_fmax_mhz()]], [[intel::num_simd_work_items()]],
44
// [[intel::no_global_work_offset()]], [[intel::max_global_work_dim()]], [[sycl::reqd_sub_group_size()]],
55
// [[sycl::reqd_work_group_size()]], [[intel::kernel_args_restrict]], [[intel::max_work_group_size()]],
6-
// and [[intel::sycl_explicit_simd]] function attributes in SYCL 2020.
6+
// [[sycl::work_group_size_hint()]] and [[intel::sycl_explicit_simd]] function attributes in SYCL 2020.
77

88
#include "sycl.hpp"
99

@@ -144,6 +144,19 @@ class Functor10 {
144144
}
145145
};
146146

147+
class Foo11 {
148+
public:
149+
[[sycl::work_group_size_hint(1, 2, 3)]] void operator()() const {}
150+
};
151+
152+
template <int SIZE, int SIZE1, int SIZE2>
153+
class Functor11 {
154+
public:
155+
[[sycl::work_group_size_hint(SIZE, SIZE1, SIZE2)]] void operator()() const {}
156+
};
157+
158+
[[sycl::work_group_size_hint(1, 2, 3)]] void foo11() {}
159+
147160
int main() {
148161
q.submit([&](handler &h) {
149162
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 !kernel_arg_buffer_location ![[NUM:[0-9]+]] !scheduler_target_fmax_mhz ![[NUM1:[0-9]+]]
@@ -320,6 +333,27 @@ int main() {
320333
// CHECK: define {{.*}}spir_func void @{{.*}}(%class.anon{{.*}} addrspace(4)* noalias noundef align 1 dereferenceable_or_null(1) %this) #4 align 2
321334
h.single_task<class kernel_name34>(
322335
[]() [[intel::kernel_args_restrict]]{});
336+
337+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name35() #0 {{.*}} !work_group_size_hint ![[NUM123:[0-9]+]]
338+
Foo11 boo11;
339+
h.single_task<class kernel_name35>(boo11);
340+
341+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name36() #0 {{.*}} !work_group_size_hint ![[NUM123]]
342+
Functor11<1, 2, 3> f11;
343+
h.single_task<class kernel_name36>(f11);
344+
345+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name37() #0 {{.*}} !work_group_size_hint ![[NUM123]]
346+
h.single_task<class kernel_name37>(
347+
[]() [[sycl::work_group_size_hint(1, 2, 3)]]{});
348+
349+
// Test attribute is not propagated.
350+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name38()
351+
// CHECK-NOT: !work_group_size_hint
352+
// CHECK-SAME: {
353+
// CHECK: define dso_local spir_func void @_Z5foo11v()
354+
h.single_task<class kernel_name38>(
355+
[]() { foo11(); });
356+
323357
});
324358
return 0;
325359
}
@@ -333,3 +367,4 @@ int main() {
333367
// CHECK: ![[NUM32]] = !{i32 16, i32 16, i32 32}
334368
// CHECK: ![[NUM88]] = !{i32 8, i32 8, i32 8}
335369
// CHECK: ![[NUM22]] = !{i32 2, i32 2, i32 2}
370+
// CHECK: ![[NUM123]] = !{i32 1, i32 2, i32 3}

‎clang/test/SemaSYCL/check-work-group-size-hint-device.cpp

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -166,6 +166,19 @@ void invoke() {
166166
// CHECK: FunctionDecl {{.*}} {{.*}}kernel_3
167167
// CHECK-NOT: WorkGroupSizeHintAttr
168168

169+
h.single_task<class kernel_name4>([]() [[sycl::work_group_size_hint(4,4,4)]] {});
170+
// CHECK: FunctionDecl {{.*}}kernel_name4
171+
// CHECK: WorkGroupSizeHintAttr {{.*}}
172+
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
173+
// CHECK-NEXT: value: Int 4
174+
// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}}
175+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
176+
// CHECK-NEXT: value: Int 4
177+
// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}}
178+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
179+
// CHECK-NEXT: value: Int 4
180+
// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}}
181+
169182
});
170183

171184
// FIXME: Add tests with the C++23 lambda attribute syntax.

0 commit comments

Comments
 (0)