Skip to content

Commit 1e9cbed

Browse files
ranapratap55ranapratap55
committed
[AMDGPU] Extends builtin support for amdgcn_image_sample and adds sema checking tests
1 parent d6b2395 commit 1e9cbed

File tree

5 files changed

+564
-2
lines changed

5 files changed

+564
-2
lines changed

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

Lines changed: 15 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/store builtins
639+
// Image builtins
640640
//===----------------------------------------------------------------------===//
641641
TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f32_i32, "V4fiiV8iii", "nc", "image-insts")
642642
TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f16_i32, "V4hiiV8iii", "nc", "image-insts")
@@ -694,6 +694,20 @@ TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f32_i32, "vV4fiiiiiV8iii",
694694
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f16_i32, "vV4hiiiiiV8iii", "nc", "image-insts")
695695
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f32_i32, "vV4fiiiiiV8iii", "nc", "image-insts")
696696
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f16_i32, "vV4hiiiiiV8iii", "nc", "image-insts")
697+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1d_v4f32_f32, "V4fifV8iV4ibii", "nc", "image-insts")
698+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1d_v4f16_f32, "V4hifV8iV4ibii", "nc", "image-insts")
699+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1darray_v4f32_f32, "V4fiffV8iV4ibii", "nc", "image-insts")
700+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1darray_v4f16_f32, "V4hiffV8iV4ibii", "nc", "image-insts")
701+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_f32_f32, "fiffV8iV4ibii", "nc", "image-insts")
702+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_v4f32_f32, "V4fiffV8iV4ibii", "nc", "image-insts")
703+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_v4f16_f32, "V4hiffV8iV4ibii", "nc", "image-insts")
704+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_f32_f32, "fifffV8iV4ibii", "nc", "image-insts")
705+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_v4f32_f32, "V4fifffV8iV4ibii", "nc", "image-insts")
706+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_v4f16_f32, "V4hifffV8iV4ibii", "nc", "image-insts")
707+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_3d_v4f32_f32, "V4fifffV8iV4ibii", "nc", "image-insts")
708+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_3d_v4f16_f32, "V4hifffV8iV4ibii", "nc", "image-insts")
709+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_cube_v4f32_f32, "V4fifffV8iV4ibii", "nc", "image-insts")
710+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_cube_v4f16_f32, "V4hifffV8iV4ibii", "nc", "image-insts")
697711

698712
#undef BUILTIN
699713
#undef TARGET_BUILTIN

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

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -805,6 +805,32 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
805805
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32:
806806
return EmitAMDGCNImageOverloadedReturnType(
807807
*this, E, Intrinsic::amdgcn_image_store_mip_cube, true);
808+
case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f32_f32:
809+
case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f16_f32:
810+
return EmitAMDGCNImageOverloadedReturnType(
811+
*this, E, Intrinsic::amdgcn_image_sample_1d, false);
812+
case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f32_f32:
813+
case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f16_f32:
814+
return EmitAMDGCNImageOverloadedReturnType(
815+
*this, E, Intrinsic::amdgcn_image_sample_1darray, false);
816+
case AMDGPU::BI__builtin_amdgcn_image_sample_2d_f32_f32:
817+
case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f32_f32:
818+
case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f16_f32:
819+
return EmitAMDGCNImageOverloadedReturnType(
820+
*this, E, Intrinsic::amdgcn_image_sample_2d, false);
821+
case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_f32_f32:
822+
case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f32_f32:
823+
case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f16_f32:
824+
return EmitAMDGCNImageOverloadedReturnType(
825+
*this, E, Intrinsic::amdgcn_image_sample_2darray, false);
826+
case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f32_f32:
827+
case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f16_f32:
828+
return EmitAMDGCNImageOverloadedReturnType(
829+
*this, E, Intrinsic::amdgcn_image_sample_3d, false);
830+
case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f32_f32:
831+
case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f16_f32:
832+
return EmitAMDGCNImageOverloadedReturnType(
833+
*this, E, Intrinsic::amdgcn_image_sample_cube, false);
808834
case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4:
809835
case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4: {
810836
llvm::FixedVectorType *VT = FixedVectorType::get(Builder.getInt32Ty(), 8);

‎clang/lib/Sema/SemaAMDGPU.cpp‎

Lines changed: 15 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -111,7 +111,21 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
111111
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
112112
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
113113
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
114-
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32: {
114+
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32:
115+
case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f32_f32:
116+
case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f32_f32:
117+
case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f16_f32:
118+
case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f16_f32:
119+
case AMDGPU::BI__builtin_amdgcn_image_sample_2d_f32_f32:
120+
case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f32_f32:
121+
case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f16_f32:
122+
case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_f32_f32:
123+
case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f32_f32:
124+
case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f16_f32:
125+
case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f32_f32:
126+
case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f16_f32:
127+
case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f32_f32:
128+
case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f16_f32: {
115129
StringRef FeatureList(getASTContext().BuiltinInfo.getRequiredFeatures(BuiltinID));
116130
if (!Builtin::evaluateRequiredTargetFeatures(
117131
FeatureList, CallerFeatureMap)){

0 commit comments

Comments
 (0)