From 0d3bfaa417ddb0914254ddfc43e4d19ac5b7e59a Mon Sep 17 00:00:00 2001 From: Jinsong Ji Date: Fri, 27 Jun 2025 12:18:18 -0700 Subject: [PATCH] [SYCL] Ensure the metadata type is correct We use getInt to generate MD, we might endup getting wrong type in IR. eg: -1 will be generated as i64 type. --- clang/lib/CodeGen/CodeGenFunction.cpp | 24 +++++++++---------- .../test/CodeGenSYCL/reqd-work-group-size.cpp | 5 ++++ 2 files changed, 17 insertions(+), 12 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 30748d694cd09..cd3387b5ac3e1 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -715,13 +715,13 @@ void CodeGenFunction::EmitKernelMetadata(const FunctionDecl *FD, // On SYCL target the dimensions are reversed if present. if (ZDimVal) - AttrMDArgs.push_back( - llvm::ConstantAsMetadata::get(Builder.getInt(*ZDimVal))); + AttrMDArgs.push_back(llvm::ConstantAsMetadata::get( + Builder.getInt32(ZDimVal->getLimitedValue()))); if (YDimVal) - AttrMDArgs.push_back( - llvm::ConstantAsMetadata::get(Builder.getInt(*YDimVal))); - AttrMDArgs.push_back( - llvm::ConstantAsMetadata::get(Builder.getInt(*XDimVal))); + AttrMDArgs.push_back(llvm::ConstantAsMetadata::get( + Builder.getInt32(YDimVal->getLimitedValue()))); + AttrMDArgs.push_back(llvm::ConstantAsMetadata::get( + Builder.getInt32(XDimVal->getLimitedValue()))); for (auto i = AttrMDArgs.size(); i < 3; ++i) AttrMDArgs.push_back( @@ -752,17 +752,17 @@ void CodeGenFunction::EmitKernelMetadata(const FunctionDecl *FD, llvm::APInt NumDims(32, 1); // X // On SYCL target the dimensions are reversed if present. if (ZDimVal) { - AttrMDArgs.push_back( - llvm::ConstantAsMetadata::get(Builder.getInt(*ZDimVal))); + AttrMDArgs.push_back(llvm::ConstantAsMetadata::get( + Builder.getInt32(ZDimVal->getLimitedValue()))); ++NumDims; } if (YDimVal) { - AttrMDArgs.push_back( - llvm::ConstantAsMetadata::get(Builder.getInt(*YDimVal))); + AttrMDArgs.push_back(llvm::ConstantAsMetadata::get( + Builder.getInt32(YDimVal->getLimitedValue()))); ++NumDims; } - AttrMDArgs.push_back( - llvm::ConstantAsMetadata::get(Builder.getInt(*XDimVal))); + AttrMDArgs.push_back(llvm::ConstantAsMetadata::get( + Builder.getInt32(XDimVal->getLimitedValue()))); for (auto i = NumDims.getZExtValue(); i < 3; ++i) AttrMDArgs.push_back( diff --git a/clang/test/CodeGenSYCL/reqd-work-group-size.cpp b/clang/test/CodeGenSYCL/reqd-work-group-size.cpp index 542655a94ac3a..824347dc9c485 100644 --- a/clang/test/CodeGenSYCL/reqd-work-group-size.cpp +++ b/clang/test/CodeGenSYCL/reqd-work-group-size.cpp @@ -102,6 +102,9 @@ int main() { h.single_task( []() [[sycl::reqd_work_group_size(1)]]{}); + // Test invalid group size + h.single_task( + []() [[sycl::reqd_work_group_size(18446744073709551615UL)]]{}); }); return 0; } @@ -122,6 +125,7 @@ int main() { // CHECK: define {{.*}} void @{{.*}}kernel_name21() #0 {{.*}} !work_group_num_dim ![[NDRWGS1D:[0-9]+]] !reqd_work_group_size ![[WGSIZE2D2_or_WGSIZE1D8]] // CHECK: define {{.*}} void @{{.*}}kernel_name22() #0 {{.*}} !work_group_num_dim ![[NDRWGS1D:[0-9]+]] !reqd_work_group_size ![[WGSIZE1D22:[0-9]+]] // CHECK: define {{.*}} void @{{.*}}kernel_name24() #0 {{.*}} !work_group_num_dim ![[NDRWGS1D:[0-9]+]] !reqd_work_group_size ![[WGSIZE1D2:[0-9]+]] +// CHECK: define {{.*}} void @{{.*}}kernel_name26() #0 {{.*}} !work_group_num_dim ![[NDRWGS1D:[0-9]+]] !reqd_work_group_size ![[WGSIZE1D26:[0-9]+]] // CHECK: ![[NDRWGS3D]] = !{i32 3} // CHECK: ![[WGSIZE3D32]] = !{i32 16, i32 16, i32 32} @@ -137,3 +141,4 @@ int main() { // CHECK: ![[WGSIZE1D32]] = !{i32 32, i32 1, i32 1} // CHECK: ![[WGSIZE1D22]] = !{i32 2, i32 1, i32 1} // CHECK: ![[WGSIZE1D2]] = !{i32 1, i32 1, i32 1} +// CHECK: ![[WGSIZE1D26]] = !{i32 -1, i32 1, i32 1}