llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-amdgpu Author: Matt Arsenault (arsenm) <details> <summary>Changes</summary> These cannot be 0. --- Full diff: https://github.com/llvm/llvm-project/pull/113038.diff 2 Files Affected: - (modified) clang/lib/CodeGen/CGBuiltin.cpp (+6) - (modified) clang/test/CodeGenOpenCL/builtins-amdgcn.cl (+2-1) ``````````diff diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index 28f28c70b5ae52..69a7dfc2433ae8 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -18538,6 +18538,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(), std::nullopt)); return LD; diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index bf5f2971cf118c..be6cee5e9217bf 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 %.sink -// 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) } `````````` </details> https://github.com/llvm/llvm-project/pull/113038 _______________________________________________ llvm-branch-commits mailing list llvm-branch-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
