[PATCH] D138755: [Clang][CodeGen] Use poison instead of undef for extra argument in __builtin_amdgcn_mov_dpp [NFC]

[Manuel Brito via Phabricator via cfe-commits]

ManuelJBrito created this revision.
ManuelJBrito added a project: AMDGPU.
Herald added subscribers: kosarev, jvesely.
Herald added a project: All.
ManuelJBrito requested review of this revision.
Herald added a project: clang.
Herald added a subscriber: cfe-commits.

Use poison instead of undef when emitting llvm.amdgcn.update.dpp for 
__builtin_amdgcn_mov_dpp.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D138755

Files:
  clang/lib/CodeGen/CGBuiltin.cpp
  clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl


Index: clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
===================================================================
--- clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
+++ clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
@@ -101,7 +101,7 @@
 }
 
 // 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);
Index: clang/lib/CodeGen/CGBuiltin.cpp
===================================================================
--- clang/lib/CodeGen/CGBuiltin.cpp
+++ clang/lib/CodeGen/CGBuiltin.cpp
@@ -16893,7 +16893,7 @@
       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);



Index: clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
===================================================================
--- clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
+++ clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
@@ -101,7 +101,7 @@
 }
 
 // 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);
Index: clang/lib/CodeGen/CGBuiltin.cpp
===================================================================
--- clang/lib/CodeGen/CGBuiltin.cpp
+++ clang/lib/CodeGen/CGBuiltin.cpp
@@ -16893,7 +16893,7 @@
       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);

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits