================
@@ -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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
