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

[ManuelJBrito]

Summary

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

[ManuelJBrito]

Couldn't find documentation for this intrinsic, can you shed a light on whether or not this change is OK?

[ManuelJBrito]

Is this change OK? It seems to me that here undef is just a placeholder. Was hoping that some AMDGPU folks could review it.

[arsenm]

I think this is OK. It’s use of uninitialized value, those uninitialized values just happen to come from inactive lanes.

[ManuelJBrito]

Thanks for the review arsenm!

clang/lib/CodeGen/CGBuiltin.cpp

case AMDGPU::BI__builtin_amdgcn_mov_dpp8:
     return emitBinaryBuiltin(*this, E, Intrinsic::amdgcn_mov_dpp8);
   case AMDGPU::BI__builtin_amdgcn_mov_dpp:
   case AMDGPU::BI__builtin_amdgcn_update_dpp: {
     llvm::SmallVector<llvm::Value *, 6> Args;
     for (unsigned I = 0; I != E->getNumArgs(); ++I)
       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);
   }
   case AMDGPU::BI__builtin_amdgcn_div_fixup:
   case AMDGPU::BI__builtin_amdgcn_div_fixupf:
   case AMDGPU::BI__builtin_amdgcn_div_fixuph:
     return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_div_fixup);

clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl

 // CHECK-LABEL: @test_s_dcache_wb()
 // CHECK: call void @llvm.amdgcn.s.dcache.wb()
 void test_s_dcache_wb()
 {
   __builtin_amdgcn_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);
 }
 
 // CHECK-LABEL: @test_update_dpp
 // CHECK: call i32 @llvm.amdgcn.update.dpp.i32(i32 %arg1, i32 %arg2, i32 0, i32 0, i32 0, i1 false)
 void test_update_dpp(global int* out, int arg1, int arg2)