From 481170cb551124ed89353dd55a89b1d9971e9a6d Mon Sep 17 00:00:00 2001 From: Manuel Brito Date: Tue, 6 Dec 2022 12:39:38 +0000 Subject: [PATCH] [Clang][CodeGen] Use poison instead of undef for extra argument in __builtin_amdgcn_mov_dpp [NFC] Differential Revision: https://reviews.llvm.org/D138755 --- clang/lib/CodeGen/CGBuiltin.cpp | 2 +- clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index 0138111..d40910b 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 daba1b9..c7437d7 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); -- 2.7.4