From f42de80151b4d3c981c61a6d526026d79adc4036 Mon Sep 17 00:00:00 2001 From: Wenju He Date: Fri, 26 Jun 2026 01:56:16 +0200 Subject: [PATCH 1/2] [SYCL][NFC] Fix check-work-group-attributes-match.cpp for AMDGPU b3ca5fb8a7b5 introduced amdgpu-flat-work-group-size attributes with per-kernel values, causing each kernel to get a distinct attribute group (#0, #2, #4) on amdgcn targets. Relax the attribute group checks from the hardcoded #0 to #{{[0-9]+}}. Fixes CodeGenSYCL/check-work-group-attributes-match.cpp CMPLRLLVM-76303 Co-Authored-By: Claude Sonnet 4.6 --- .../test/CodeGenSYCL/check-work-group-attributes-match.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/test/CodeGenSYCL/check-work-group-attributes-match.cpp b/clang/test/CodeGenSYCL/check-work-group-attributes-match.cpp index 8ecc995dc566b..bbc89bf601a22 100644 --- a/clang/test/CodeGenSYCL/check-work-group-attributes-match.cpp +++ b/clang/test/CodeGenSYCL/check-work-group-attributes-match.cpp @@ -15,17 +15,17 @@ int main() { queue q; q.submit([&](handler &h) { - // CHECK: define {{.*}} void @{{.*}}kernel_1d() #0 {{.*}} !work_group_size_hint ![[WGSH1D:[0-9]+]]{{.*}} !work_group_num_dim ![[NDRWGS1D:[0-9]+]]{{.*}} !reqd_work_group_size ![[WGSH1D]] + // CHECK: define {{.*}} void @{{.*}}kernel_1d() #{{[0-9]+}} {{.*}} !work_group_size_hint ![[WGSH1D:[0-9]+]]{{.*}} !work_group_num_dim ![[NDRWGS1D:[0-9]+]]{{.*}} !reqd_work_group_size ![[WGSH1D]] h.single_task([]() [[sycl::work_group_size_hint(8)]] [[sycl::reqd_work_group_size(8)]] {}); }); q.submit([&](handler &h) { - // CHECK: define {{.*}} void @{{.*}}kernel_2d() #0 {{.*}} !work_group_size_hint ![[WGSH2D:[0-9]+]]{{.*}} !work_group_num_dim ![[NDRWGS2D:[0-9]+]]{{.*}} !reqd_work_group_size ![[WGSH2D:[0-9]+]]{{.*}} + // CHECK: define {{.*}} void @{{.*}}kernel_2d() #{{[0-9]+}} {{.*}} !work_group_size_hint ![[WGSH2D:[0-9]+]]{{.*}} !work_group_num_dim ![[NDRWGS2D:[0-9]+]]{{.*}} !reqd_work_group_size ![[WGSH2D:[0-9]+]]{{.*}} h.single_task([]() [[sycl::work_group_size_hint(8, 16)]] [[sycl::reqd_work_group_size(8, 16)]] {}); }); q.submit([&](handler &h) { - // CHECK: define {{.*}} void @{{.*}}kernel_3d() #0 {{.*}} !work_group_size_hint ![[WG3D:[0-9]+]]{{.*}} !work_group_num_dim ![[NDRWGS3D:[0-9]+]]{{.*}} !reqd_work_group_size ![[WG3D]] + // CHECK: define {{.*}} void @{{.*}}kernel_3d() #{{[0-9]+}} {{.*}} !work_group_size_hint ![[WG3D:[0-9]+]]{{.*}} !work_group_num_dim ![[NDRWGS3D:[0-9]+]]{{.*}} !reqd_work_group_size ![[WG3D]] h.single_task([]() [[sycl::work_group_size_hint(8, 16, 32)]] [[sycl::reqd_work_group_size(8, 16, 32)]] {}); }); } From 3157aae8f957666533a2bd509a831841be4b4bd8 Mon Sep 17 00:00:00 2001 From: Wenju He Date: Fri, 26 Jun 2026 01:58:16 +0200 Subject: [PATCH 2/2] [[#]] --- .../test/CodeGenSYCL/check-work-group-attributes-match.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/test/CodeGenSYCL/check-work-group-attributes-match.cpp b/clang/test/CodeGenSYCL/check-work-group-attributes-match.cpp index bbc89bf601a22..1a32bfbe54d28 100644 --- a/clang/test/CodeGenSYCL/check-work-group-attributes-match.cpp +++ b/clang/test/CodeGenSYCL/check-work-group-attributes-match.cpp @@ -15,17 +15,17 @@ int main() { queue q; q.submit([&](handler &h) { - // CHECK: define {{.*}} void @{{.*}}kernel_1d() #{{[0-9]+}} {{.*}} !work_group_size_hint ![[WGSH1D:[0-9]+]]{{.*}} !work_group_num_dim ![[NDRWGS1D:[0-9]+]]{{.*}} !reqd_work_group_size ![[WGSH1D]] + // CHECK: define {{.*}} void @{{.*}}kernel_1d() #[[#]] {{.*}} !work_group_size_hint ![[WGSH1D:[0-9]+]]{{.*}} !work_group_num_dim ![[NDRWGS1D:[0-9]+]]{{.*}} !reqd_work_group_size ![[WGSH1D]] h.single_task([]() [[sycl::work_group_size_hint(8)]] [[sycl::reqd_work_group_size(8)]] {}); }); q.submit([&](handler &h) { - // CHECK: define {{.*}} void @{{.*}}kernel_2d() #{{[0-9]+}} {{.*}} !work_group_size_hint ![[WGSH2D:[0-9]+]]{{.*}} !work_group_num_dim ![[NDRWGS2D:[0-9]+]]{{.*}} !reqd_work_group_size ![[WGSH2D:[0-9]+]]{{.*}} + // CHECK: define {{.*}} void @{{.*}}kernel_2d() #[[#]] {{.*}} !work_group_size_hint ![[WGSH2D:[0-9]+]]{{.*}} !work_group_num_dim ![[NDRWGS2D:[0-9]+]]{{.*}} !reqd_work_group_size ![[WGSH2D:[0-9]+]]{{.*}} h.single_task([]() [[sycl::work_group_size_hint(8, 16)]] [[sycl::reqd_work_group_size(8, 16)]] {}); }); q.submit([&](handler &h) { - // CHECK: define {{.*}} void @{{.*}}kernel_3d() #{{[0-9]+}} {{.*}} !work_group_size_hint ![[WG3D:[0-9]+]]{{.*}} !work_group_num_dim ![[NDRWGS3D:[0-9]+]]{{.*}} !reqd_work_group_size ![[WG3D]] + // CHECK: define {{.*}} void @{{.*}}kernel_3d() #[[#]] {{.*}} !work_group_size_hint ![[WG3D:[0-9]+]]{{.*}} !work_group_num_dim ![[NDRWGS3D:[0-9]+]]{{.*}} !reqd_work_group_size ![[WG3D]] h.single_task([]() [[sycl::work_group_size_hint(8, 16, 32)]] [[sycl::reqd_work_group_size(8, 16, 32)]] {}); }); }