================
@@ -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) {
+ if (Qt->isSignedIntegerType())
+ return Builder.CreateSExt(Val, I32Ty);
+ else
+ return Builder.CreateZExt(Val, I32Ty);
+ } else
+ return Builder.CreateZExtOrTrunc(Val, I32Ty);
+ }
+
+ if (Ty->isPointerTy()) {
+ unsigned PtrBits = DL.getPointerSizeInBits(Ty->getPointerAddressSpace());
+ llvm::Type *IntPtrTy = Builder.getIntNTy(PtrBits);
+ llvm::Value *AsInt = Builder.CreatePtrToInt(Val, IntPtrTy);
+ return Builder.CreateZExtOrTrunc(AsInt, I32Ty);
+ }
+
+ unsigned Bits = GetBitWidth(Ty);
+ llvm::Type *IntN = Builder.getIntNTy(Bits);
+ llvm::Value *AsInt = Builder.CreateBitCast(Val, IntN);
+ return Builder.CreateZExtOrTrunc(AsInt, I32Ty);
+ };
+
+ // Bit-preserving resize/cast between arbitrary source and destination LLVM
+ // types.
+ auto BitCoerceTo = [&](llvm::Value *Val, llvm::Type *DstTy) -> llvm::Value *
{
+ llvm::Type *SrcTy = Val->getType();
+ if (SrcTy == DstTy)
+ return Val;
+
+ unsigned SrcBits = DL.getTypeSizeInBits(SrcTy).getFixedValue();
+ unsigned DstBits = DL.getTypeSizeInBits(DstTy).getFixedValue();
+
+ if (SrcTy->isIntegerTy() && DstTy->isIntegerTy())
+ return Builder.CreateZExtOrTrunc(Val, DstTy);
+
+ if (SrcBits == DstBits)
+ return Builder.CreateBitCast(Val, DstTy);
+
+ llvm::Type *IntSrcTy = Builder.getIntNTy(SrcBits);
+ llvm::Value *AsInt = Val;
+ if (SrcTy->isPointerTy())
+ AsInt = Builder.CreatePtrToInt(Val, IntSrcTy);
+ else if (!SrcTy->isIntegerTy())
+ AsInt = Builder.CreateBitCast(Val, IntSrcTy);
+
+ llvm::Type *IntDstTy = Builder.getIntNTy(DstBits);
+ llvm::Value *Resized = Builder.CreateZExtOrTrunc(AsInt, IntDstTy);
+
+ if (DstTy->isPointerTy())
+ return Builder.CreateIntToPtr(Resized, DstTy);
+
+ return Builder.CreateBitCast(Resized, DstTy);
+ };
+
+ llvm::Value *Index = CGF.EmitScalarExpr(Call->getArg(0));
+ llvm::Value *Source = CGF.EmitScalarExpr(Call->getArg(1));
+
+ llvm::Type *ReturnTy = CGF.ConvertType(Call->getType());
+
+ llvm::Value *IndexI32 = ToI32Bits(Index, Call->getArg(0)->getType());
+
+ llvm::Value *SourceForIntrinsic;
+ llvm::Type *SourceTy = Source->getType();
+
+ if (SourceTy->isDoubleTy()) {
+ llvm::Value *AsFloat = Builder.CreateFPTrunc(Source, Builder.getFloatTy());
----------------
shiltian wrote:
Do we really want to do this? Why is double handed here but other types are not?
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
