================
@@ -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());
----------------
yxsamliu wrote:
we also accept 64 bit integer types. They are all implicitly converted to 32
bit types (double to float, long to int). The situation is similar to the
situation of double passed to a function accepting float, or long passed to a
function accepting int, which causes implicit conversion with a warning.
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
