@@ -18544,31 +18544,19 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
18544
18544
case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8i16: {
18545
18545
18546
18546
Intrinsic::ID IID;
18547
- llvm::Type *ArgTy;
18548
18547
switch (BuiltinID) {
18549
18548
case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_i32:
18550
- ArgTy = llvm::Type::getInt32Ty(getLLVMContext());
18551
- IID = Intrinsic::amdgcn_global_load_tr_b64;
18552
- break;
18553
18549
case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_v2i32:
18554
- ArgTy = llvm::FixedVectorType::get(
18555
- llvm::Type::getInt32Ty(getLLVMContext()), 2);
18556
18550
IID = Intrinsic::amdgcn_global_load_tr_b64;
18557
18551
break;
18558
18552
case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4i16:
18559
- ArgTy = llvm::FixedVectorType::get(
18560
- llvm::Type::getInt16Ty(getLLVMContext()), 4);
18561
- IID = Intrinsic::amdgcn_global_load_tr_b128;
18562
- break;
18563
18553
case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8i16:
18564
- ArgTy = llvm::FixedVectorType::get(
18565
- llvm::Type::getInt16Ty(getLLVMContext()), 8);
18566
18554
IID = Intrinsic::amdgcn_global_load_tr_b128;
18567
18555
break;
18568
18556
}
18569
-
18557
+ llvm::Type *LoadTy = ConvertType(E->getType());
18570
18558
llvm::Value *Addr = EmitScalarExpr(E->getArg(0));
18571
- llvm::Function *F = CGM.getIntrinsic(IID, {ArgTy });
18559
+ llvm::Function *F = CGM.getIntrinsic(IID, {LoadTy });
18572
18560
return Builder.CreateCall(F, {Addr});
18573
18561
}
18574
18562
case AMDGPU::BI__builtin_amdgcn_get_fpenv: {
0 commit comments