Author: Changpeng Fang Date: 2024-03-26T17:51:03-07:00 New Revision: d023995ae2cc6ab968ec305ea7b6f11b6ac2e78f
URL: https://github.com/llvm/llvm-project/commit/d023995ae2cc6ab968ec305ea7b6f11b6ac2e78f DIFF: https://github.com/llvm/llvm-project/commit/d023995ae2cc6ab968ec305ea7b6f11b6ac2e78f.diff LOG: AMDGPU: Simplify EmitAMDGPUBuiltinExpr for load transposes, NFC (#86707) We should not manually get the types of the loading data. Instead, we can get the types from the intrinsics directly. Added: Modified: clang/lib/CodeGen/CGBuiltin.cpp Removed: ################################################################################ diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index 46a815155e7b87..3cfdb261a0eac0 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -18544,31 +18544,19 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8i16: { Intrinsic::ID IID; - llvm::Type *ArgTy; switch (BuiltinID) { case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_i32: - ArgTy = llvm::Type::getInt32Ty(getLLVMContext()); - IID = Intrinsic::amdgcn_global_load_tr_b64; - break; case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_v2i32: - ArgTy = llvm::FixedVectorType::get( - llvm::Type::getInt32Ty(getLLVMContext()), 2); IID = Intrinsic::amdgcn_global_load_tr_b64; break; case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4i16: - ArgTy = llvm::FixedVectorType::get( - llvm::Type::getInt16Ty(getLLVMContext()), 4); - IID = Intrinsic::amdgcn_global_load_tr_b128; - break; case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8i16: - ArgTy = llvm::FixedVectorType::get( - llvm::Type::getInt16Ty(getLLVMContext()), 8); IID = Intrinsic::amdgcn_global_load_tr_b128; break; } - + llvm::Type *LoadTy = ConvertType(E->getType()); llvm::Value *Addr = EmitScalarExpr(E->getArg(0)); - llvm::Function *F = CGM.getIntrinsic(IID, {ArgTy}); + llvm::Function *F = CGM.getIntrinsic(IID, {LoadTy}); return Builder.CreateCall(F, {Addr}); } case AMDGPU::BI__builtin_amdgcn_get_fpenv: { _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
