[clang] 481170c - [Clang][CodeGen] Use poison instead of undef for extra argument in __builtin_amdgcn_mov_dpp [NFC]
Nuno Lopes via cfe-commits
cfe-commits at lists.llvm.org
Tue Dec 6 04:41:02 PST 2022
Author: Manuel Brito
Date: 2022-12-06T12:40:33Z
New Revision: 481170cb551124ed89353dd55a89b1d9971e9a6d
URL: https://github.com/llvm/llvm-project/commit/481170cb551124ed89353dd55a89b1d9971e9a6d
DIFF: https://github.com/llvm/llvm-project/commit/481170cb551124ed89353dd55a89b1d9971e9a6d.diff
LOG: [Clang][CodeGen] Use poison instead of undef for extra argument in __builtin_amdgcn_mov_dpp [NFC]
Differential Revision: https://reviews.llvm.org/D138755
Added:
Modified:
clang/lib/CodeGen/CGBuiltin.cpp
clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
Removed:
################################################################################
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 0138111758657..d40910b76b599 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -16905,7 +16905,7 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
Args.push_back(EmitScalarExpr(E->getArg(I)));
assert(Args.size() == 5 || Args.size() == 6);
if (Args.size() == 5)
- Args.insert(Args.begin(), llvm::UndefValue::get(Args[0]->getType()));
+ Args.insert(Args.begin(), llvm::PoisonValue::get(Args[0]->getType()));
Function *F =
CGM.getIntrinsic(Intrinsic::amdgcn_update_dpp, Args[0]->getType());
return Builder.CreateCall(F, Args);
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
index daba1b972787e..c7437d7c1e881 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
@@ -101,7 +101,7 @@ void test_s_dcache_wb()
}
// CHECK-LABEL: @test_mov_dpp
-// CHECK: call i32 @llvm.amdgcn.update.dpp.i32(i32 undef, i32 %src, i32 0, i32 0, i32 0, i1 false)
+// CHECK: call i32 @llvm.amdgcn.update.dpp.i32(i32 poison, i32 %src, i32 0, i32 0, i32 0, i1 false)
void test_mov_dpp(global int* out, int src)
{
*out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0, false);
More information about the cfe-commits
mailing list