[clang] [llvm] [AMDGPU] Extend __builtin_amdgcn_ds_bpermute argument types (PR #153501)

[Shilei Tian via cfe-commits]

================
@@ -159,6 +159,119 @@ Value *EmitAMDGPUGridSize(CodeGenFunction &CGF, unsigned Index) {
                   llvm::MDNode::get(CGF.getLLVMContext(), {}));
   return LD;
 }
+// Lowers __builtin_amdgcn_ds_bpermute to the corresponding LLVM intrinsic with
+// careful bit-level coercions of operands and result to match Clang types.
+llvm::Value *emitAMDGCNDsBpermute(clang::CodeGen::CodeGenFunction &Cgf,
+                                  const clang::CallExpr *Call) {
+  auto &Builder = Cgf.Builder;
+  auto &Cgm = Cgf.CGM;
+  const llvm::DataLayout &Dl = Cgm.getDataLayout();
+
+  llvm::Type *I32Ty = Builder.getInt32Ty();
+
+  auto GetBitWidth = [&](llvm::Type *Ty) -> unsigned {
+    return Dl.getTypeSizeInBits(Ty).getFixedValue();
+  };
+
+  // Coerces arbitrary scalar/vector/pointer to i32 by preserving value/bit
+  // semantics where applicable.
+  auto ToI32Bits = [&](llvm::Value *Val, clang::QualType Qt) -> llvm::Value * {
+    llvm::Type *Ty = Val->getType();
+
+    if (Ty->isIntegerTy()) {
+      unsigned BitWidth = Ty->getIntegerBitWidth();
+      if (BitWidth < 32) {
----------------
shiltian wrote:

maybe `<=`?

https://github.com/llvm/llvm-project/pull/153501