Author: Matt Arsenault Date: 2024-11-05T12:47:04-08:00 New Revision: 0c60573d1c2d19133d84da092b240f32e0574be5
URL: https://github.com/llvm/llvm-project/commit/0c60573d1c2d19133d84da092b240f32e0574be5 DIFF: https://github.com/llvm/llvm-project/commit/0c60573d1c2d19133d84da092b240f32e0574be5.diff LOG: clang/AMDGPU: Emit grid size builtins with range metadata (#113038) These cannot be 0. Added: Modified: clang/lib/CodeGen/CGBuiltin.cpp clang/test/CodeGenOpenCL/builtins-amdgcn.cl Removed: ################################################################################ diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index fb731413d8c9ad..82770a75af23e4 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -18671,6 +18671,12 @@ Value *EmitAMDGPUGridSize(CodeGenFunction &CGF, unsigned Index) { auto *GEP = CGF.Builder.CreateGEP(CGF.Int8Ty, DP, Offset); auto *LD = CGF.Builder.CreateLoad( Address(GEP, CGF.Int32Ty, CharUnits::fromQuantity(4))); + + llvm::MDBuilder MDB(CGF.getLLVMContext()); + + // Known non-zero. + LD->setMetadata(llvm::LLVMContext::MD_range, + MDB.createRange(APInt(32, 1), APInt::getZero(32))); LD->setMetadata(llvm::LLVMContext::MD_invariant_load, llvm::MDNode::get(CGF.getLLVMContext(), {})); return LD; diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index 9132cc8a717e0f..3bc6107b7fd40d 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -639,7 +639,7 @@ void test_get_workgroup_size(int d, global int *out) // CHECK-LABEL: @test_get_grid_size( // CHECK: {{.*}}call align 4 dereferenceable(64){{.*}} ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() // CHECK: getelementptr inbounds i8, ptr addrspace(4) %{{.*}}, i64 %{{.+}} -// CHECK: load i32, ptr addrspace(4) %{{.*}}, align 4, !invariant.load +// CHECK: load i32, ptr addrspace(4) %{{.*}}, align 4, !range [[$GRID_RANGE:![0-9]+]], !invariant.load void test_get_grid_size(int d, global int *out) { switch (d) { @@ -896,5 +896,6 @@ void test_set_fpenv(unsigned long env) { __builtin_amdgcn_set_fpenv(env); } +// CHECK-DAG: [[$GRID_RANGE]] = !{i32 1, i32 0} // CHECK-DAG: [[$WS_RANGE]] = !{i16 1, i16 1025} // CHECK-DAG: attributes #[[$NOUNWIND_READONLY]] = { convergent mustprogress nocallback nofree nounwind willreturn memory(none) } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
