Skip to content

Commit d6b2395

Browse files
ranapratap55ranapratap55
committed
[AMDGPU] Minor changes
1 parent f512dae commit d6b2395

File tree

2 files changed

+28
-36
lines changed

2 files changed

+28
-36
lines changed

‎clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp‎

Lines changed: 26 additions & 32 deletions
Original file line numberDiff line numberDiff line change
@@ -163,18 +163,12 @@ static Value *EmitAMDGCNBallotForExec(CodeGenFunction &CGF, const CallExpr *E,
163163
return Call;
164164
}
165165

166-
template <unsigned N>
167166
llvm::CallInst *EmitAMDGCNImageOverloadedReturnType(clang::CodeGen::CodeGenFunction &CGF,
168167
const clang::CallExpr *E,
169168
unsigned IntrinsicID,
170169
bool IsImageStore) {
171-
static_assert(N, "expect non-empty argument");
172-
173-
assert(E->getNumArgs() == N &&
174-
"Argument count mismatch with builtin definition");
175-
176-
clang::SmallVector<llvm::Value *, N> Args;
177-
for (unsigned I = 0; I < N; ++I)
170+
clang::SmallVector<llvm::Value *, 10> Args;
171+
for (unsigned I = 0; I < E->getNumArgs(); ++I)
178172
Args.push_back(CGF.EmitScalarExpr(E->getArg(I)));
179173

180174
llvm::Type *RetTy = CGF.ConvertType(E->getType());
@@ -709,107 +703,107 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
709703
}
710704
case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
711705
case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:
712-
return EmitAMDGCNImageOverloadedReturnType<5>(
706+
return EmitAMDGCNImageOverloadedReturnType(
713707
*this, E, Intrinsic::amdgcn_image_load_1d, false);
714708
case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
715709
case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:
716-
return EmitAMDGCNImageOverloadedReturnType<6>(
710+
return EmitAMDGCNImageOverloadedReturnType(
717711
*this, E, Intrinsic::amdgcn_image_load_1darray, false);
718712
case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32:
719713
case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32:
720714
case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32:
721-
return EmitAMDGCNImageOverloadedReturnType<6>(
715+
return EmitAMDGCNImageOverloadedReturnType(
722716
*this, E, Intrinsic::amdgcn_image_load_2d, false);
723717
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32:
724718
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32:
725719
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32:
726-
return EmitAMDGCNImageOverloadedReturnType<7>(
720+
return EmitAMDGCNImageOverloadedReturnType(
727721
*this, E, Intrinsic::amdgcn_image_load_2darray, false);
728722
case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
729723
case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
730-
return EmitAMDGCNImageOverloadedReturnType<7>(
724+
return EmitAMDGCNImageOverloadedReturnType(
731725
*this, E, Intrinsic::amdgcn_image_load_3d, false);
732726
case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
733727
case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
734-
return EmitAMDGCNImageOverloadedReturnType<7>(
728+
return EmitAMDGCNImageOverloadedReturnType(
735729
*this, E, Intrinsic::amdgcn_image_load_cube, false);
736730
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
737731
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:
738-
return EmitAMDGCNImageOverloadedReturnType<6>(
732+
return EmitAMDGCNImageOverloadedReturnType(
739733
*this, E, Intrinsic::amdgcn_image_load_mip_1d, false);
740734
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:
741735
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:
742-
return EmitAMDGCNImageOverloadedReturnType<7>(
736+
return EmitAMDGCNImageOverloadedReturnType(
743737
*this, E, Intrinsic::amdgcn_image_load_mip_1darray, false);
744738
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32:
745739
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:
746740
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32:
747-
return EmitAMDGCNImageOverloadedReturnType<7>(
741+
return EmitAMDGCNImageOverloadedReturnType(
748742
*this, E, Intrinsic::amdgcn_image_load_mip_2d, false);
749743
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32:
750744
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32:
751745
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32:
752-
return EmitAMDGCNImageOverloadedReturnType<8>(
746+
return EmitAMDGCNImageOverloadedReturnType(
753747
*this, E, Intrinsic::amdgcn_image_load_mip_2darray, false);
754748
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
755749
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
756-
return EmitAMDGCNImageOverloadedReturnType<8>(
750+
return EmitAMDGCNImageOverloadedReturnType(
757751
*this, E, Intrinsic::amdgcn_image_load_mip_3d, false);
758752
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
759753
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32:
760-
return EmitAMDGCNImageOverloadedReturnType<8>(
754+
return EmitAMDGCNImageOverloadedReturnType(
761755
*this, E, Intrinsic::amdgcn_image_load_mip_cube, false);
762756
case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:
763757
case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32:
764-
return EmitAMDGCNImageOverloadedReturnType<6>(
758+
return EmitAMDGCNImageOverloadedReturnType(
765759
*this, E, Intrinsic::amdgcn_image_store_1d, true);
766760
case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:
767761
case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:
768-
return EmitAMDGCNImageOverloadedReturnType<7>(
762+
return EmitAMDGCNImageOverloadedReturnType(
769763
*this, E, Intrinsic::amdgcn_image_store_1darray, true);
770764
case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32:
771765
case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32:
772766
case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32:
773-
return EmitAMDGCNImageOverloadedReturnType<7>(
767+
return EmitAMDGCNImageOverloadedReturnType(
774768
*this, E, Intrinsic::amdgcn_image_store_2d, true);
775769
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32:
776770
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32:
777771
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32:
778-
return EmitAMDGCNImageOverloadedReturnType<8>(
772+
return EmitAMDGCNImageOverloadedReturnType(
779773
*this, E, Intrinsic::amdgcn_image_store_2darray, true);
780774
case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:
781775
case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:
782-
return EmitAMDGCNImageOverloadedReturnType<8>(
776+
return EmitAMDGCNImageOverloadedReturnType(
783777
*this, E, Intrinsic::amdgcn_image_store_3d, true);
784778
case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:
785779
case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:
786-
return EmitAMDGCNImageOverloadedReturnType<8>(
780+
return EmitAMDGCNImageOverloadedReturnType(
787781
*this, E, Intrinsic::amdgcn_image_store_cube, true);
788782
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:
789783
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32:
790-
return EmitAMDGCNImageOverloadedReturnType<7>(
784+
return EmitAMDGCNImageOverloadedReturnType(
791785
*this, E, Intrinsic::amdgcn_image_store_mip_1d, true);
792786
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:
793787
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:
794-
return EmitAMDGCNImageOverloadedReturnType<8>(
788+
return EmitAMDGCNImageOverloadedReturnType(
795789
*this, E, Intrinsic::amdgcn_image_store_mip_1darray, true);
796790
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:
797791
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:
798792
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32:
799-
return EmitAMDGCNImageOverloadedReturnType<8>(
793+
return EmitAMDGCNImageOverloadedReturnType(
800794
*this, E, Intrinsic::amdgcn_image_store_mip_2d, true);
801795
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32:
802796
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32:
803797
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32:
804-
return EmitAMDGCNImageOverloadedReturnType<9>(
798+
return EmitAMDGCNImageOverloadedReturnType(
805799
*this, E, Intrinsic::amdgcn_image_store_mip_2darray, true);
806800
case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:
807801
case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:
808-
return EmitAMDGCNImageOverloadedReturnType<9>(
802+
return EmitAMDGCNImageOverloadedReturnType(
809803
*this, E, Intrinsic::amdgcn_image_store_mip_3d, true);
810804
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
811805
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32:
812-
return EmitAMDGCNImageOverloadedReturnType<9>(
806+
return EmitAMDGCNImageOverloadedReturnType(
813807
*this, E, Intrinsic::amdgcn_image_store_mip_cube, true);
814808
case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4:
815809
case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4: {

‎clang/lib/Sema/SemaAMDGPU.cpp‎

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -113,9 +113,8 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
113113
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
114114
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32: {
115115
StringRef FeatureList(getASTContext().BuiltinInfo.getRequiredFeatures(BuiltinID));
116-
bool HasImageInsts = Builtin::evaluateRequiredTargetFeatures("image-insts", CallerFeatureMap);
117116
if (!Builtin::evaluateRequiredTargetFeatures(
118-
FeatureList, CallerFeatureMap) && !HasImageInsts){
117+
FeatureList, CallerFeatureMap)){
119118
Diag(TheCall->getBeginLoc(), diag::err_builtin_needs_feature)
120119
<< FD->getDeclName() << FeatureList;
121120
return false;
@@ -157,9 +156,8 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
157156
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
158157
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32: {
159158
StringRef FeatureList(getASTContext().BuiltinInfo.getRequiredFeatures(BuiltinID));
160-
bool HasImageInsts = Builtin::evaluateRequiredTargetFeatures("image-insts", CallerFeatureMap);
161159
if (!Builtin::evaluateRequiredTargetFeatures(
162-
FeatureList, CallerFeatureMap) && !HasImageInsts){
160+
FeatureList, CallerFeatureMap)){
163161
Diag(TheCall->getBeginLoc(), diag::err_builtin_needs_feature)
164162
<< FD->getDeclName() << FeatureList;
165163
return false;

0 commit comments

Comments
 (0)