Skip to content

Commit 7168d19

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

File tree

6 files changed

+1105
-35
lines changed

6 files changed

+1105
-35
lines changed

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

Lines changed: 30 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -636,9 +636,36 @@ 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 builtins
640-
//===----------------------------------------------------------------------===//
641-
BUILTIN(__builtin_amdgcn_image_load_2d_f32_i32, "fiiV8i", "n")
639+
// Image load builtins
640+
//===----------------------------------------------------------------------===//
641+
TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f32_i32, "V4fiiV8iii", "nc", "")
642+
TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f16_i32, "V4hiiV8iii", "nc", "")
643+
TARGET_BUILTIN(__builtin_amdgcn_image_load_1darray_v4f32_i32, "V4fiiiV8iii", "nc", "")
644+
TARGET_BUILTIN(__builtin_amdgcn_image_load_1darray_v4f16_i32, "V4hiiiV8iii", "nc", "")
645+
TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_f32_i32, "fiiiV8iii", "nc", "")
646+
TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_v4f32_i32, "V4fiiiV8iii", "nc", "")
647+
TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_v4f16_i32, "V4hiiiV8iii", "nc", "")
648+
TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_f32_i32, "fiiiiV8iii", "nc", "")
649+
TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_v4f32_i32, "V4fiiiiV8iii", "nc", "")
650+
TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_v4f16_i32, "V4hiiiiV8iii", "nc", "")
651+
TARGET_BUILTIN(__builtin_amdgcn_image_load_3d_v4f32_i32, "V4fiiiiV8iii", "nc", "")
652+
TARGET_BUILTIN(__builtin_amdgcn_image_load_3d_v4f16_i32, "V4hiiiiV8iii", "nc", "")
653+
TARGET_BUILTIN(__builtin_amdgcn_image_load_cube_v4f32_i32, "V4fiiiiV8iii", "nc", "")
654+
TARGET_BUILTIN(__builtin_amdgcn_image_load_cube_v4f16_i32, "V4hiiiiV8iii", "nc", "")
655+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1d_v4f32_i32, "V4fiiiV8iii", "nc", "")
656+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1d_v4f16_i32, "V4hiiiV8iii", "nc", "")
657+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1darray_v4f32_i32, "V4fiiiiV8iii", "nc", "")
658+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1darray_v4f16_i32, "V4hiiiiV8iii", "nc", "")
659+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_f32_i32, "fiiiiV8iii", "nc", "")
660+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_v4f32_i32, "V4fiiiiV8iii", "nc", "")
661+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_v4f16_i32, "V4hiiiiV8iii", "nc", "")
662+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_f32_i32, "fiiiiiV8iii", "nc", "")
663+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_v4f32_i32, "V4fiiiiiV8iii", "nc", "")
664+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_v4f16_i32, "V4hiiiiiV8iii", "nc", "")
665+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_3d_v4f32_i32, "V4fiiiiiV8iii", "nc", "")
666+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_3d_v4f16_i32, "V4hiiiiiV8iii", "nc", "")
667+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_cube_v4f32_i32, "V4fiiiiiV8iii", "nc", "")
668+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_cube_v4f16_i32, "V4hiiiiiV8iii", "nc", "")
642669

643670
#undef BUILTIN
644671
#undef TARGET_BUILTIN

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

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

32+
bool checkImageImmArgFunctionCall(CallExpr *TheCall, unsigned ArgCount);
33+
3234
/// Create an AMDGPUWavesPerEUAttr attribute.
3335
AMDGPUFlatWorkGroupSizeAttr *
3436
CreateAMDGPUFlatWorkGroupSizeAttr(const AttributeCommonInfo &CI, Expr *Min,

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

Lines changed: 196 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -683,27 +683,203 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
683683

684684
return Builder.CreateInsertElement(I0, A, 1);
685685
}
686-
case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32: {
687-
llvm::Type *RetTy = llvm::Type::getFloatTy(Builder.getContext());
688-
llvm::Type *IntTy = llvm::IntegerType::get(Builder.getContext(), 32u);
689-
690-
llvm::Value *imm0 = llvm::ConstantInt::get(IntTy, 1);
691-
llvm::Value *arg0 = EmitScalarExpr(E->getArg(0));
692-
llvm::Value *arg1 = EmitScalarExpr(E->getArg(1));
693-
llvm::Value *arg2 = EmitScalarExpr(E->getArg(2));
694-
llvm::Value *imm1 = llvm::ConstantInt::get(IntTy, 0);
695-
llvm::Value *imm2 = llvm::ConstantInt::get(IntTy, 0);
696-
697-
SmallVector<Value *, 6> ArgTys;
698-
ArgTys.push_back(imm0);
699-
ArgTys.push_back(arg0);
700-
ArgTys.push_back(arg1);
701-
ArgTys.push_back(arg2);
702-
ArgTys.push_back(imm1);
703-
ArgTys.push_back(imm2);
686+
case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
687+
case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
688+
case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:
689+
case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:
690+
case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32:
691+
case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32:
692+
case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32:
693+
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32:
694+
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32:
695+
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32:
696+
case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
697+
case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
698+
case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
699+
case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
700+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
701+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:
702+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:
703+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:
704+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32:
705+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:
706+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32:
707+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32:
708+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32:
709+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32:
710+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
711+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
712+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
713+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32: {
714+
llvm::Type *RetTy = nullptr;
715+
switch (BuiltinID) {
716+
case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32:
717+
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32:
718+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32:
719+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32:
720+
RetTy = llvm::Type::getFloatTy(Builder.getContext());
721+
break;
722+
case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
723+
case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
724+
case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32:
725+
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32:
726+
case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
727+
case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
728+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
729+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:
730+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:
731+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32:
732+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
733+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
734+
RetTy =
735+
FixedVectorType::get(llvm::Type::getFloatTy(Builder.getContext()), 4);
736+
break;
737+
case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:
738+
case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:
739+
case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32:
740+
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32:
741+
case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
742+
case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
743+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:
744+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:
745+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32:
746+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32:
747+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
748+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32:
749+
RetTy =
750+
FixedVectorType::get(llvm::Type::getHalfTy(Builder.getContext()), 4);
751+
break;
752+
}
704753

705-
llvm::CallInst *Call =
706-
Builder.CreateIntrinsic(RetTy, Intrinsic::amdgcn_image_load_2d, ArgTys);
754+
llvm::Value *Dmask = EmitScalarExpr(E->getArg(0));
755+
llvm::Value *S = EmitScalarExpr(E->getArg(1));
756+
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;
762+
763+
SmallVector<Value *, 10> ArgTys;
764+
765+
Intrinsic::ID IID;
766+
llvm::CallInst *Call;
767+
768+
switch (BuiltinID) {
769+
case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
770+
case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32: {
771+
Rsrc = EmitScalarExpr(E->getArg(2));
772+
Tfe = EmitScalarExpr(E->getArg(3));
773+
Cpol = EmitScalarExpr(E->getArg(4));
774+
775+
ArgTys = {Dmask, S, Rsrc, Tfe, Cpol};
776+
IID = Intrinsic::amdgcn_image_load_1d;
777+
Call = Builder.CreateIntrinsic(RetTy, IID, ArgTys);
778+
break;
779+
}
780+
case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
781+
case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:
782+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
783+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32: {
784+
Slice = EmitScalarExpr(E->getArg(2));
785+
Rsrc = EmitScalarExpr(E->getArg(3));
786+
Tfe = EmitScalarExpr(E->getArg(4));
787+
Cpol = EmitScalarExpr(E->getArg(5));
788+
789+
ArgTys = {Dmask, S, Slice, Rsrc, Tfe, Cpol};
790+
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:
794+
IID = Intrinsic::amdgcn_image_load_mip_1d;
795+
break;
796+
}
797+
Call = Builder.CreateIntrinsic(RetTy, IID, ArgTys);
798+
break;
799+
}
800+
case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32:
801+
case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32:
802+
case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32: {
803+
Rsrc = EmitScalarExpr(E->getArg(3));
804+
Tfe = EmitScalarExpr(E->getArg(4));
805+
Cpol = EmitScalarExpr(E->getArg(5));
806+
807+
ArgTys = {Dmask, S, T, Rsrc, Tfe, Cpol};
808+
IID = Intrinsic::amdgcn_image_load_2d;
809+
Call = Builder.CreateIntrinsic(RetTy, IID, ArgTys);
810+
break;
811+
}
812+
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32:
813+
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32:
814+
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32:
815+
case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
816+
case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
817+
case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
818+
case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
819+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:
820+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:
821+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32:
822+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:
823+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32: {
824+
Slice = EmitScalarExpr(E->getArg(3));
825+
Rsrc = EmitScalarExpr(E->getArg(4));
826+
Tfe = EmitScalarExpr(E->getArg(5));
827+
Cpol = EmitScalarExpr(E->getArg(6));
828+
829+
ArgTys = {Dmask, S, T, Slice, Rsrc, Tfe, Cpol};
830+
IID = Intrinsic::amdgcn_image_load_2darray;
831+
832+
switch (BuiltinID) {
833+
case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
834+
case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
835+
IID = Intrinsic::amdgcn_image_load_3d;
836+
break;
837+
case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
838+
case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
839+
IID = Intrinsic::amdgcn_image_load_cube;
840+
break;
841+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:
842+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:
843+
IID = Intrinsic::amdgcn_image_load_mip_1darray;
844+
break;
845+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32:
846+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:
847+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32:
848+
IID = Intrinsic::amdgcn_image_load_mip_2d;
849+
break;
850+
}
851+
Call = Builder.CreateIntrinsic(RetTy, IID, ArgTys);
852+
break;
853+
}
854+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32:
855+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32:
856+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32:
857+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
858+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
859+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
860+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32: {
861+
Slice = EmitScalarExpr(E->getArg(3));
862+
Mip = EmitScalarExpr(E->getArg(4));
863+
Rsrc = EmitScalarExpr(E->getArg(5));
864+
Tfe = EmitScalarExpr(E->getArg(6));
865+
Cpol = EmitScalarExpr(E->getArg(7));
866+
867+
ArgTys = {Dmask, S, T, Slice, Mip, Rsrc, Tfe, Cpol};
868+
IID = Intrinsic::amdgcn_image_load_mip_2darray;
869+
switch (BuiltinID) {
870+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
871+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
872+
IID = Intrinsic::amdgcn_image_load_mip_3d;
873+
break;
874+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
875+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32:
876+
IID = Intrinsic::amdgcn_image_load_mip_cube;
877+
break;
878+
}
879+
Call = Builder.CreateIntrinsic(RetTy, IID, ArgTys);
880+
break;
881+
}
882+
}
707883

708884
return Call;
709885
}

‎clang/lib/Sema/SemaAMDGPU.cpp‎

Lines changed: 42 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -83,6 +83,38 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
8383
case AMDGPU::BI__builtin_amdgcn_update_dpp: {
8484
return checkMovDPPFunctionCall(TheCall, 6, 2);
8585
}
86+
case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
87+
case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
88+
case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:
89+
case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:
90+
case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32:
91+
case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32:
92+
case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32:
93+
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32:
94+
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32:
95+
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32:
96+
case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
97+
case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
98+
case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
99+
case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
100+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
101+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:
102+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:
103+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:
104+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32:
105+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:
106+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32:
107+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32:
108+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32:
109+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32:
110+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
111+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
112+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
113+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32: {
114+
unsigned ArgCount = TheCall->getNumArgs() - 1;
115+
116+
return checkImageImmArgFunctionCall(TheCall, ArgCount);
117+
}
86118
default:
87119
return false;
88120
}
@@ -128,6 +160,16 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
128160
return false;
129161
}
130162

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+
131173
bool SemaAMDGPU::checkMovDPPFunctionCall(CallExpr *TheCall, unsigned NumArgs,
132174
unsigned NumDataArgs) {
133175
assert(NumDataArgs <= 2);

0 commit comments

Comments
 (0)