Skip to content

Commit ab18d68

Browse files
ranapratap55ranapratap55
committed
[AMDGPU] Adds builtins for image store and sema checking for image store
1 parent 7168d19 commit ab18d68

File tree

8 files changed

+1073
-24
lines changed

8 files changed

+1073
-24
lines changed

‎clang/include/clang/Basic/BuiltinsAMDGPU.def‎

Lines changed: 30 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -636,7 +636,7 @@ TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_bf16_f32, "V2yV2yfUiIb", "nc", "f32-to-f1
636636
TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_f16_f32, "V2hV2hfUiIb", "nc", "f32-to-f16bf16-cvt-sr-insts")
637637

638638
//===----------------------------------------------------------------------===//
639-
// Image load builtins
639+
// Image load/store builtins
640640
//===----------------------------------------------------------------------===//
641641
TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f32_i32, "V4fiiV8iii", "nc", "")
642642
TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f16_i32, "V4hiiV8iii", "nc", "")
@@ -667,5 +667,34 @@ TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_3d_v4f16_i32, "V4hiiiiiV8iii", "n
667667
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_cube_v4f32_i32, "V4fiiiiiV8iii", "nc", "")
668668
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_cube_v4f16_i32, "V4hiiiiiV8iii", "nc", "")
669669

670+
TARGET_BUILTIN(__builtin_amdgcn_image_store_1d_v4f32_i32, "vV4fiiV8iii", "nc", "")
671+
TARGET_BUILTIN(__builtin_amdgcn_image_store_1d_v4f16_i32, "vV4hiiV8iii", "nc", "")
672+
TARGET_BUILTIN(__builtin_amdgcn_image_store_1darray_v4f32_i32, "vV4fiiiV8iii", "nc", "")
673+
TARGET_BUILTIN(__builtin_amdgcn_image_store_1darray_v4f16_i32, "vV4hiiiV8iii", "nc", "")
674+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_f32_i32, "vfiiiV8iii", "nc", "")
675+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_v4f32_i32, "vV4fiiiV8iii", "nc", "")
676+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_v4f16_i32, "vV4hiiiV8iii", "nc", "")
677+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_f32_i32, "vfiiiiV8iii", "nc", "")
678+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_v4f32_i32, "vV4fiiiiV8iii", "nc", "")
679+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_v4f16_i32, "vV4hiiiiV8iii", "nc", "")
680+
TARGET_BUILTIN(__builtin_amdgcn_image_store_3d_v4f32_i32, "vV4fiiiiV8iii", "nc", "")
681+
TARGET_BUILTIN(__builtin_amdgcn_image_store_3d_v4f16_i32, "vV4hiiiiV8iii", "nc", "")
682+
TARGET_BUILTIN(__builtin_amdgcn_image_store_cube_v4f32_i32, "vV4fiiiiV8iii", "nc", "")
683+
TARGET_BUILTIN(__builtin_amdgcn_image_store_cube_v4f16_i32, "vV4hiiiiV8iii", "nc", "")
684+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1d_v4f32_i32, "vV4fiiiV8iii", "nc", "")
685+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1d_v4f16_i32, "vV4hiiiV8iii", "nc", "")
686+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1darray_v4f32_i32, "vV4fiiiiV8iii", "nc", "")
687+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1darray_v4f16_i32, "vV4hiiiiV8iii", "nc", "")
688+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_f32_i32, "vfiiiiV8iii", "nc", "")
689+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_v4f32_i32, "vV4fiiiiV8iii", "nc", "")
690+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_v4f16_i32, "vV4hiiiiV8iii", "nc", "")
691+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_f32_i32, "vfiiiiiV8iii", "nc", "")
692+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_v4f32_i32, "vV4fiiiiiV8iii", "nc", "")
693+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_v4f16_i32, "vV4hiiiiiV8iii", "nc", "")
694+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f32_i32, "vV4fiiiiiV8iii", "nc", "")
695+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f16_i32, "vV4hiiiiiV8iii", "nc", "")
696+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f32_i32, "vV4fiiiiiV8iii", "nc", "")
697+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f16_i32, "vV4hiiiiiV8iii", "nc", "")
698+
670699
#undef BUILTIN
671700
#undef TARGET_BUILTIN

‎clang/include/clang/Sema/SemaAMDGPU.h‎

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -29,8 +29,6 @@ class SemaAMDGPU : public SemaBase {
2929
bool checkMovDPPFunctionCall(CallExpr *TheCall, unsigned NumArgs,
3030
unsigned NumDataArgs);
3131

32-
bool checkImageImmArgFunctionCall(CallExpr *TheCall, unsigned ArgCount);
33-
3432
/// Create an AMDGPUWavesPerEUAttr attribute.
3533
AMDGPUFlatWorkGroupSizeAttr *
3634
CreateAMDGPUFlatWorkGroupSizeAttr(const AttributeCommonInfo &CI, Expr *Min,

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

Lines changed: 168 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -754,11 +754,11 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
754754
llvm::Value *Dmask = EmitScalarExpr(E->getArg(0));
755755
llvm::Value *S = EmitScalarExpr(E->getArg(1));
756756
llvm::Value *T = EmitScalarExpr(E->getArg(2));
757-
llvm::Value *Slice;
758-
llvm::Value *Mip;
759-
llvm::Value *Rsrc;
760-
llvm::Value *Tfe;
761-
llvm::Value *Cpol;
757+
llvm::Value *Slice = nullptr;
758+
llvm::Value *Mip = nullptr;
759+
llvm::Value *Rsrc = nullptr;
760+
llvm::Value *Tfe = nullptr;
761+
llvm::Value *Cpol = nullptr;
762762

763763
SmallVector<Value *, 10> ArgTys;
764764

@@ -788,11 +788,9 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
788788

789789
ArgTys = {Dmask, S, Slice, Rsrc, Tfe, Cpol};
790790
IID = Intrinsic::amdgcn_image_load_1darray;
791-
switch (BuiltinID) {
792-
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
793-
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:
791+
if (BuiltinID == AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32 ||
792+
BuiltinID == AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32) {
794793
IID = Intrinsic::amdgcn_image_load_mip_1d;
795-
break;
796794
}
797795
Call = Builder.CreateIntrinsic(RetTy, IID, ArgTys);
798796
break;
@@ -883,6 +881,167 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
883881

884882
return Call;
885883
}
884+
case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:
885+
case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:
886+
case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32:
887+
case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:
888+
case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32:
889+
case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32:
890+
case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32:
891+
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32:
892+
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32:
893+
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32:
894+
case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:
895+
case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:
896+
case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:
897+
case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:
898+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:
899+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32:
900+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:
901+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:
902+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:
903+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:
904+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32:
905+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32:
906+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32:
907+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32:
908+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:
909+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:
910+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
911+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32: {
912+
llvm::Type *RetTy = VoidTy;
913+
llvm::Value *Vdata = EmitScalarExpr(E->getArg(0));
914+
llvm::Value *Dmask = EmitScalarExpr(E->getArg(1));
915+
llvm::Value *S = EmitScalarExpr(E->getArg(2));
916+
llvm::Value *T = EmitScalarExpr(E->getArg(3));
917+
llvm::Value *Slice = nullptr;
918+
llvm::Value *Mip = nullptr;
919+
llvm::Value *Rsrc = nullptr;
920+
llvm::Value *Tfe = nullptr;
921+
llvm::Value *Cpol = nullptr;
922+
923+
SmallVector<Value *, 10> ArgTys;
924+
925+
Intrinsic::ID IID;
926+
llvm::CallInst *Call;
927+
928+
switch (BuiltinID) {
929+
case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:
930+
case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32: {
931+
Rsrc = EmitScalarExpr(E->getArg(3));
932+
Tfe = EmitScalarExpr(E->getArg(4));
933+
Cpol = EmitScalarExpr(E->getArg(5));
934+
935+
ArgTys = {Vdata, Dmask, S, Rsrc, Tfe, Cpol};
936+
IID = Intrinsic::amdgcn_image_store_1d;
937+
Call = Builder.CreateIntrinsic(RetTy, IID, ArgTys);
938+
break;
939+
}
940+
case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:
941+
case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:
942+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:
943+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32: {
944+
Slice = EmitScalarExpr(E->getArg(3));
945+
Rsrc = EmitScalarExpr(E->getArg(4));
946+
Tfe = EmitScalarExpr(E->getArg(5));
947+
Cpol = EmitScalarExpr(E->getArg(6));
948+
949+
ArgTys = {Vdata, Dmask, S, Slice, Rsrc, Tfe, Cpol};
950+
IID = Intrinsic::amdgcn_image_store_1darray;
951+
if (BuiltinID ==
952+
AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32 ||
953+
BuiltinID ==
954+
AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32) {
955+
IID = Intrinsic::amdgcn_image_store_mip_1d;
956+
}
957+
Call = Builder.CreateIntrinsic(RetTy, IID, ArgTys);
958+
break;
959+
}
960+
case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32:
961+
case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32:
962+
case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32: {
963+
Rsrc = EmitScalarExpr(E->getArg(4));
964+
Tfe = EmitScalarExpr(E->getArg(5));
965+
Cpol = EmitScalarExpr(E->getArg(6));
966+
967+
ArgTys = {Vdata, Dmask, S, T, Rsrc, Tfe, Cpol};
968+
IID = Intrinsic::amdgcn_image_store_2d;
969+
Call = Builder.CreateIntrinsic(RetTy, IID, ArgTys);
970+
break;
971+
}
972+
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32:
973+
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32:
974+
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32:
975+
case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:
976+
case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:
977+
case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:
978+
case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:
979+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:
980+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:
981+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:
982+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:
983+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32: {
984+
Slice = EmitScalarExpr(E->getArg(4));
985+
Rsrc = EmitScalarExpr(E->getArg(5));
986+
Tfe = EmitScalarExpr(E->getArg(6));
987+
Cpol = EmitScalarExpr(E->getArg(7));
988+
989+
ArgTys = {Vdata, Dmask, S, T, Slice, Rsrc, Tfe, Cpol};
990+
IID = Intrinsic::amdgcn_image_store_2darray;
991+
992+
switch (BuiltinID) {
993+
case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:
994+
case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:
995+
IID = Intrinsic::amdgcn_image_store_3d;
996+
break;
997+
case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:
998+
case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:
999+
IID = Intrinsic::amdgcn_image_store_cube;
1000+
break;
1001+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:
1002+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:
1003+
IID = Intrinsic::amdgcn_image_store_mip_1darray;
1004+
break;
1005+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:
1006+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:
1007+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32:
1008+
IID = Intrinsic::amdgcn_image_store_mip_2d;
1009+
break;
1010+
}
1011+
Call = Builder.CreateIntrinsic(RetTy, IID, ArgTys);
1012+
break;
1013+
}
1014+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32:
1015+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32:
1016+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32:
1017+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:
1018+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:
1019+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
1020+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32: {
1021+
Slice = EmitScalarExpr(E->getArg(4));
1022+
Mip = EmitScalarExpr(E->getArg(5));
1023+
Rsrc = EmitScalarExpr(E->getArg(6));
1024+
Tfe = EmitScalarExpr(E->getArg(7));
1025+
Cpol = EmitScalarExpr(E->getArg(8));
1026+
1027+
ArgTys = {Vdata, Dmask, S, T, Slice, Mip, Rsrc, Tfe, Cpol};
1028+
IID = Intrinsic::amdgcn_image_store_mip_2darray;
1029+
switch (BuiltinID) {
1030+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:
1031+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:
1032+
IID = Intrinsic::amdgcn_image_store_mip_3d;
1033+
break;
1034+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
1035+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32:
1036+
IID = Intrinsic::amdgcn_image_store_mip_cube;
1037+
break;
1038+
}
1039+
Call = Builder.CreateIntrinsic(RetTy, IID, ArgTys);
1040+
break;
1041+
}
1042+
}
1043+
return Call;
1044+
}
8861045
case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4:
8871046
case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4: {
8881047
llvm::FixedVectorType *VT = FixedVectorType::get(Builder.getInt32Ty(), 8);

‎clang/lib/Sema/SemaAMDGPU.cpp‎

Lines changed: 48 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -112,8 +112,54 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
112112
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
113113
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32: {
114114
unsigned ArgCount = TheCall->getNumArgs() - 1;
115-
116-
return checkImageImmArgFunctionCall(TheCall, ArgCount);
115+
llvm::APSInt Result;
116+
bool isImmArg =
117+
(!(SemaRef.BuiltinConstantArg(TheCall, 0, Result)) &&
118+
!(SemaRef.BuiltinConstantArg(TheCall, ArgCount, Result)) &&
119+
!(SemaRef.BuiltinConstantArg(TheCall, (ArgCount - 1), Result)))
120+
? false
121+
: true;
122+
123+
return isImmArg;
124+
}
125+
case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:
126+
case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:
127+
case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32:
128+
case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:
129+
case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32:
130+
case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32:
131+
case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32:
132+
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32:
133+
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32:
134+
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32:
135+
case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:
136+
case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:
137+
case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:
138+
case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:
139+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:
140+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32:
141+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:
142+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:
143+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:
144+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:
145+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32:
146+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32:
147+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32:
148+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32:
149+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:
150+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:
151+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
152+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32: {
153+
unsigned ArgCount = TheCall->getNumArgs() - 1;
154+
llvm::APSInt Result;
155+
bool isImmArg =
156+
(!(SemaRef.BuiltinConstantArg(TheCall, 1, Result)) &&
157+
!(SemaRef.BuiltinConstantArg(TheCall, ArgCount, Result)) &&
158+
!(SemaRef.BuiltinConstantArg(TheCall, (ArgCount - 1), Result)))
159+
? false
160+
: true;
161+
162+
return isImmArg;
117163
}
118164
default:
119165
return false;
@@ -160,16 +206,6 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
160206
return false;
161207
}
162208

163-
bool SemaAMDGPU::checkImageImmArgFunctionCall(CallExpr *TheCall,
164-
unsigned ArgCount) {
165-
llvm::APSInt Result;
166-
if (!(SemaRef.BuiltinConstantArg(TheCall, 0, Result)) &&
167-
!(SemaRef.BuiltinConstantArg(TheCall, ArgCount, Result)) &&
168-
!(SemaRef.BuiltinConstantArg(TheCall, (ArgCount - 1), Result)))
169-
return false;
170-
return true;
171-
}
172-
173209
bool SemaAMDGPU::checkMovDPPFunctionCall(CallExpr *TheCall, unsigned NumArgs,
174210
unsigned NumDataArgs) {
175211
assert(NumDataArgs <= 2);

‎clang/test/CodeGen/builtins-image-load-2d-f32.c‎ renamed to ‎clang/test/CodeGen/builtins-image-load.c‎

File renamed without changes.

0 commit comments

Comments
 (0)