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

[Matt Arsenault via llvm-commits]

================
@@ -159,6 +159,242 @@ Value *EmitAMDGPUGridSize(CodeGenFunction &CGF, unsigned Index) {
                   llvm::MDNode::get(CGF.getLLVMContext(), {}));
   return LD;
 }
+
+// Emits LLVM IR to lower a generic AMDGCN ds_bpermute over arbitrary payload
+// types. Assumes DataLayout is accurate; index is coerced to i32; payload is
+// split/coerced to 32-bit words.
+llvm::Value *emitAMDGCNDsBpermute(clang::CodeGen::CodeGenFunction &CGF,
+                                  const clang::CallExpr *Call,
+                                  ReturnValueSlot Dest) {
+  auto &B = CGF.Builder;
+  auto &CGM = CGF.CGM;
+  const llvm::DataLayout &DL = CGM.getDataLayout();
+
+  llvm::Type *I8 = B.getInt8Ty();
+  llvm::Type *I32 = B.getInt32Ty();
+
+  auto C32 = [&](uint32_t V) { return llvm::ConstantInt::get(I32, V); };
+
+  // Size/bitwidth and coercion helpers for arbitrary first-class types.
+  auto GetBitWidth = [&](llvm::Type *Ty) -> unsigned {
+    return DL.getTypeSizeInBits(Ty).getFixedValue();
----------------
arsenm wrote:

Remove this lambda?

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