[clang] d023995 - AMDGPU: Simplify EmitAMDGPUBuiltinExpr for load transposes, NFC (#86707)
via cfe-commits
cfe-commits at lists.llvm.org
Tue Mar 26 17:51:07 PDT 2024
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: {
More information about the cfe-commits
mailing list