Skip to content

Commit 4fbfa90

Browse files
ranapratap55ranapratap55
committed
[WIP][AMDGPU] Support for type inferring image load/store builtins for AMDGPU
1 parent 48db3fd commit 4fbfa90

File tree

8 files changed

+2582
-0
lines changed

8 files changed

+2582
-0
lines changed

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

Lines changed: 74 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -885,5 +885,79 @@ TARGET_BUILTIN(__builtin_amdgcn_cooperative_atomic_store_16x8B, "vV2i*V2iIicC*",
885885
TARGET_BUILTIN(__builtin_amdgcn_cooperative_atomic_load_8x16B, "V4iV4i*IicC*", "nc", "gfx1250-insts,wavefrontsize32")
886886
TARGET_BUILTIN(__builtin_amdgcn_cooperative_atomic_store_8x16B, "vV4i*V4iIicC*", "nc", "gfx1250-insts,wavefrontsize32")
887887

888+
//===----------------------------------------------------------------------===//
889+
// Image builtins
890+
//===----------------------------------------------------------------------===//
891+
TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f32_i32, "V4fiiV8iii", "nc", "image-insts")
892+
TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f16_i32, "V4hiiV8iii", "nc", "image-insts")
893+
TARGET_BUILTIN(__builtin_amdgcn_image_load_1darray_v4f32_i32, "V4fiiiV8iii", "nc", "image-insts")
894+
TARGET_BUILTIN(__builtin_amdgcn_image_load_1darray_v4f16_i32, "V4hiiiV8iii", "nc", "image-insts")
895+
TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_f32_i32, "fiiiV8iii", "nc", "image-insts")
896+
TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_v4f32_i32, "V4fiiiV8iii", "nc", "image-insts")
897+
TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_v4f16_i32, "V4hiiiV8iii", "nc", "image-insts")
898+
TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_f32_i32, "fiiiiV8iii", "nc", "image-insts")
899+
TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_v4f32_i32, "V4fiiiiV8iii", "nc", "image-insts")
900+
TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_v4f16_i32, "V4hiiiiV8iii", "nc", "image-insts")
901+
TARGET_BUILTIN(__builtin_amdgcn_image_load_3d_v4f32_i32, "V4fiiiiV8iii", "nc", "image-insts")
902+
TARGET_BUILTIN(__builtin_amdgcn_image_load_3d_v4f16_i32, "V4hiiiiV8iii", "nc", "image-insts")
903+
TARGET_BUILTIN(__builtin_amdgcn_image_load_cube_v4f32_i32, "V4fiiiiV8iii", "nc", "image-insts")
904+
TARGET_BUILTIN(__builtin_amdgcn_image_load_cube_v4f16_i32, "V4hiiiiV8iii", "nc", "image-insts")
905+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1d_v4f32_i32, "V4fiiiV8iii", "nc", "image-insts")
906+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1d_v4f16_i32, "V4hiiiV8iii", "nc", "image-insts")
907+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1darray_v4f32_i32, "V4fiiiiV8iii", "nc", "image-insts")
908+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1darray_v4f16_i32, "V4hiiiiV8iii", "nc", "image-insts")
909+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_f32_i32, "fiiiiV8iii", "nc", "image-insts")
910+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_v4f32_i32, "V4fiiiiV8iii", "nc", "image-insts")
911+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_v4f16_i32, "V4hiiiiV8iii", "nc", "image-insts")
912+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_f32_i32, "fiiiiiV8iii", "nc", "image-insts")
913+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_v4f32_i32, "V4fiiiiiV8iii", "nc", "image-insts")
914+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_v4f16_i32, "V4hiiiiiV8iii", "nc", "image-insts")
915+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_3d_v4f32_i32, "V4fiiiiiV8iii", "nc", "image-insts")
916+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_3d_v4f16_i32, "V4hiiiiiV8iii", "nc", "image-insts")
917+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_cube_v4f32_i32, "V4fiiiiiV8iii", "nc", "image-insts")
918+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_cube_v4f16_i32, "V4hiiiiiV8iii", "nc", "image-insts")
919+
TARGET_BUILTIN(__builtin_amdgcn_image_store_1d_v4f32_i32, "vV4fiiV8iii", "nc", "image-insts")
920+
TARGET_BUILTIN(__builtin_amdgcn_image_store_1d_v4f16_i32, "vV4hiiV8iii", "nc", "image-insts")
921+
TARGET_BUILTIN(__builtin_amdgcn_image_store_1darray_v4f32_i32, "vV4fiiiV8iii", "nc", "image-insts")
922+
TARGET_BUILTIN(__builtin_amdgcn_image_store_1darray_v4f16_i32, "vV4hiiiV8iii", "nc", "image-insts")
923+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_f32_i32, "vfiiiV8iii", "nc", "image-insts")
924+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_v4f32_i32, "vV4fiiiV8iii", "nc", "image-insts")
925+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_v4f16_i32, "vV4hiiiV8iii", "nc", "image-insts")
926+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_f32_i32, "vfiiiiV8iii", "nc", "image-insts")
927+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_v4f32_i32, "vV4fiiiiV8iii", "nc", "image-insts")
928+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_v4f16_i32, "vV4hiiiiV8iii", "nc", "image-insts")
929+
TARGET_BUILTIN(__builtin_amdgcn_image_store_3d_v4f32_i32, "vV4fiiiiV8iii", "nc", "image-insts")
930+
TARGET_BUILTIN(__builtin_amdgcn_image_store_3d_v4f16_i32, "vV4hiiiiV8iii", "nc", "image-insts")
931+
TARGET_BUILTIN(__builtin_amdgcn_image_store_cube_v4f32_i32, "vV4fiiiiV8iii", "nc", "image-insts")
932+
TARGET_BUILTIN(__builtin_amdgcn_image_store_cube_v4f16_i32, "vV4hiiiiV8iii", "nc", "image-insts")
933+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1d_v4f32_i32, "vV4fiiiV8iii", "nc", "image-insts")
934+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1d_v4f16_i32, "vV4hiiiV8iii", "nc", "image-insts")
935+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1darray_v4f32_i32, "vV4fiiiiV8iii", "nc", "image-insts")
936+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1darray_v4f16_i32, "vV4hiiiiV8iii", "nc", "image-insts")
937+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_f32_i32, "vfiiiiV8iii", "nc", "image-insts")
938+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_v4f32_i32, "vV4fiiiiV8iii", "nc", "image-insts")
939+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_v4f16_i32, "vV4hiiiiV8iii", "nc", "image-insts")
940+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_f32_i32, "vfiiiiiV8iii", "nc", "image-insts")
941+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_v4f32_i32, "vV4fiiiiiV8iii", "nc", "image-insts")
942+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_v4f16_i32, "vV4hiiiiiV8iii", "nc", "image-insts")
943+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f32_i32, "vV4fiiiiiV8iii", "nc", "image-insts")
944+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f16_i32, "vV4hiiiiiV8iii", "nc", "image-insts")
945+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f32_i32, "vV4fiiiiiV8iii", "nc", "image-insts")
946+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f16_i32, "vV4hiiiiiV8iii", "nc", "image-insts")
947+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1d_v4f32_f32, "V4fifV8iV4ibii", "nc", "image-insts")
948+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1d_v4f16_f32, "V4hifV8iV4ibii", "nc", "image-insts")
949+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1darray_v4f32_f32, "V4fiffV8iV4ibii", "nc", "image-insts")
950+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1darray_v4f16_f32, "V4hiffV8iV4ibii", "nc", "image-insts")
951+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_f32_f32, "fiffV8iV4ibii", "nc", "image-insts")
952+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_v4f32_f32, "V4fiffV8iV4ibii", "nc", "image-insts")
953+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_v4f16_f32, "V4hiffV8iV4ibii", "nc", "image-insts")
954+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_f32_f32, "fifffV8iV4ibii", "nc", "image-insts")
955+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_v4f32_f32, "V4fifffV8iV4ibii", "nc", "image-insts")
956+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_v4f16_f32, "V4hifffV8iV4ibii", "nc", "image-insts")
957+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_3d_v4f32_f32, "V4fifffV8iV4ibii", "nc", "image-insts")
958+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_3d_v4f16_f32, "V4hifffV8iV4ibii", "nc", "image-insts")
959+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_cube_v4f32_f32, "V4fifffV8iV4ibii", "nc", "image-insts")
960+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_cube_v4f16_f32, "V4hifffV8iV4ibii", "nc", "image-insts")
961+
888962
#undef BUILTIN
889963
#undef TARGET_BUILTIN

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

Lines changed: 145 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -181,6 +181,21 @@ static Value *EmitAMDGCNBallotForExec(CodeGenFunction &CGF, const CallExpr *E,
181181
return Call;
182182
}
183183

184+
llvm::CallInst *EmitAMDGCNImageOverloadedReturnType(clang::CodeGen::CodeGenFunction &CGF,
185+
const clang::CallExpr *E,
186+
unsigned IntrinsicID,
187+
bool IsImageStore) {
188+
clang::SmallVector<llvm::Value *, 10> Args;
189+
for (unsigned I = 0; I < E->getNumArgs(); ++I)
190+
Args.push_back(CGF.EmitScalarExpr(E->getArg(I)));
191+
192+
llvm::Type *RetTy = CGF.ConvertType(E->getType());
193+
if (IsImageStore)
194+
RetTy = CGF.VoidTy;
195+
llvm::CallInst *Call = CGF.Builder.CreateIntrinsic(RetTy, IntrinsicID, Args);
196+
return Call;
197+
}
198+
184199
// Emit an intrinsic that has 1 float or double operand, and 1 integer.
185200
static Value *emitFPIntBuiltin(CodeGenFunction &CGF,
186201
const CallExpr *E,
@@ -937,6 +952,136 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
937952

938953
return Builder.CreateInsertElement(I0, A, 1);
939954
}
955+
case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
956+
case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:
957+
return EmitAMDGCNImageOverloadedReturnType(
958+
*this, E, Intrinsic::amdgcn_image_load_1d, false);
959+
case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
960+
case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:
961+
return EmitAMDGCNImageOverloadedReturnType(
962+
*this, E, Intrinsic::amdgcn_image_load_1darray, false);
963+
case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32:
964+
case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32:
965+
case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32:
966+
return EmitAMDGCNImageOverloadedReturnType(
967+
*this, E, Intrinsic::amdgcn_image_load_2d, false);
968+
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32:
969+
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32:
970+
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32:
971+
return EmitAMDGCNImageOverloadedReturnType(
972+
*this, E, Intrinsic::amdgcn_image_load_2darray, false);
973+
case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
974+
case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
975+
return EmitAMDGCNImageOverloadedReturnType(
976+
*this, E, Intrinsic::amdgcn_image_load_3d, false);
977+
case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
978+
case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
979+
return EmitAMDGCNImageOverloadedReturnType(
980+
*this, E, Intrinsic::amdgcn_image_load_cube, false);
981+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
982+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:
983+
return EmitAMDGCNImageOverloadedReturnType(
984+
*this, E, Intrinsic::amdgcn_image_load_mip_1d, false);
985+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:
986+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:
987+
return EmitAMDGCNImageOverloadedReturnType(
988+
*this, E, Intrinsic::amdgcn_image_load_mip_1darray, false);
989+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32:
990+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:
991+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32:
992+
return EmitAMDGCNImageOverloadedReturnType(
993+
*this, E, Intrinsic::amdgcn_image_load_mip_2d, false);
994+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32:
995+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32:
996+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32:
997+
return EmitAMDGCNImageOverloadedReturnType(
998+
*this, E, Intrinsic::amdgcn_image_load_mip_2darray, false);
999+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
1000+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
1001+
return EmitAMDGCNImageOverloadedReturnType(
1002+
*this, E, Intrinsic::amdgcn_image_load_mip_3d, false);
1003+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
1004+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32:
1005+
return EmitAMDGCNImageOverloadedReturnType(
1006+
*this, E, Intrinsic::amdgcn_image_load_mip_cube, false);
1007+
case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:
1008+
case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32:
1009+
return EmitAMDGCNImageOverloadedReturnType(
1010+
*this, E, Intrinsic::amdgcn_image_store_1d, true);
1011+
case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:
1012+
case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:
1013+
return EmitAMDGCNImageOverloadedReturnType(
1014+
*this, E, Intrinsic::amdgcn_image_store_1darray, true);
1015+
case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32:
1016+
case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32:
1017+
case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32:
1018+
return EmitAMDGCNImageOverloadedReturnType(
1019+
*this, E, Intrinsic::amdgcn_image_store_2d, true);
1020+
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32:
1021+
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32:
1022+
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32:
1023+
return EmitAMDGCNImageOverloadedReturnType(
1024+
*this, E, Intrinsic::amdgcn_image_store_2darray, true);
1025+
case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:
1026+
case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:
1027+
return EmitAMDGCNImageOverloadedReturnType(
1028+
*this, E, Intrinsic::amdgcn_image_store_3d, true);
1029+
case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:
1030+
case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:
1031+
return EmitAMDGCNImageOverloadedReturnType(
1032+
*this, E, Intrinsic::amdgcn_image_store_cube, true);
1033+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:
1034+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32:
1035+
return EmitAMDGCNImageOverloadedReturnType(
1036+
*this, E, Intrinsic::amdgcn_image_store_mip_1d, true);
1037+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:
1038+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:
1039+
return EmitAMDGCNImageOverloadedReturnType(
1040+
*this, E, Intrinsic::amdgcn_image_store_mip_1darray, true);
1041+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:
1042+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:
1043+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32:
1044+
return EmitAMDGCNImageOverloadedReturnType(
1045+
*this, E, Intrinsic::amdgcn_image_store_mip_2d, true);
1046+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32:
1047+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32:
1048+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32:
1049+
return EmitAMDGCNImageOverloadedReturnType(
1050+
*this, E, Intrinsic::amdgcn_image_store_mip_2darray, true);
1051+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:
1052+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:
1053+
return EmitAMDGCNImageOverloadedReturnType(
1054+
*this, E, Intrinsic::amdgcn_image_store_mip_3d, true);
1055+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
1056+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32:
1057+
return EmitAMDGCNImageOverloadedReturnType(
1058+
*this, E, Intrinsic::amdgcn_image_store_mip_cube, true);
1059+
case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f32_f32:
1060+
case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f16_f32:
1061+
return EmitAMDGCNImageOverloadedReturnType(
1062+
*this, E, Intrinsic::amdgcn_image_sample_1d, false);
1063+
case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f32_f32:
1064+
case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f16_f32:
1065+
return EmitAMDGCNImageOverloadedReturnType(
1066+
*this, E, Intrinsic::amdgcn_image_sample_1darray, false);
1067+
case AMDGPU::BI__builtin_amdgcn_image_sample_2d_f32_f32:
1068+
case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f32_f32:
1069+
case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f16_f32:
1070+
return EmitAMDGCNImageOverloadedReturnType(
1071+
*this, E, Intrinsic::amdgcn_image_sample_2d, false);
1072+
case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_f32_f32:
1073+
case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f32_f32:
1074+
case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f16_f32:
1075+
return EmitAMDGCNImageOverloadedReturnType(
1076+
*this, E, Intrinsic::amdgcn_image_sample_2darray, false);
1077+
case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f32_f32:
1078+
case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f16_f32:
1079+
return EmitAMDGCNImageOverloadedReturnType(
1080+
*this, E, Intrinsic::amdgcn_image_sample_3d, false);
1081+
case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f32_f32:
1082+
case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f16_f32:
1083+
return EmitAMDGCNImageOverloadedReturnType(
1084+
*this, E, Intrinsic::amdgcn_image_sample_cube, false);
9401085
case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4:
9411086
case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4: {
9421087
llvm::FixedVectorType *VT = FixedVectorType::get(Builder.getInt32Ty(), 8);

0 commit comments

Comments
 (0)