Skip to content

Commit 21c6829

Browse files
ranapratap55ranapratap55
committed
[AMDGPU] Adds builtins for image store and sema checking for image store
1 parent 90065ee commit 21c6829

File tree

7 files changed

+78
-875
lines changed

7 files changed

+78
-875
lines changed

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

Lines changed: 29 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -959,5 +959,34 @@ TARGET_BUILTIN(__builtin_amdgcn_image_sample_3d_v4f16_f32, "V4hifffV8iV4ibii", "
959959
TARGET_BUILTIN(__builtin_amdgcn_image_sample_cube_v4f32_f32, "V4fifffV8iV4ibii", "nc", "image-insts")
960960
TARGET_BUILTIN(__builtin_amdgcn_image_sample_cube_v4f16_f32, "V4hifffV8iV4ibii", "nc", "image-insts")
961961

962+
TARGET_BUILTIN(__builtin_amdgcn_image_store_1d_v4f32_i32, "vV4fiiV8iii", "nc", "")
963+
TARGET_BUILTIN(__builtin_amdgcn_image_store_1d_v4f16_i32, "vV4hiiV8iii", "nc", "")
964+
TARGET_BUILTIN(__builtin_amdgcn_image_store_1darray_v4f32_i32, "vV4fiiiV8iii", "nc", "")
965+
TARGET_BUILTIN(__builtin_amdgcn_image_store_1darray_v4f16_i32, "vV4hiiiV8iii", "nc", "")
966+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_f32_i32, "vfiiiV8iii", "nc", "")
967+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_v4f32_i32, "vV4fiiiV8iii", "nc", "")
968+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_v4f16_i32, "vV4hiiiV8iii", "nc", "")
969+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_f32_i32, "vfiiiiV8iii", "nc", "")
970+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_v4f32_i32, "vV4fiiiiV8iii", "nc", "")
971+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_v4f16_i32, "vV4hiiiiV8iii", "nc", "")
972+
TARGET_BUILTIN(__builtin_amdgcn_image_store_3d_v4f32_i32, "vV4fiiiiV8iii", "nc", "")
973+
TARGET_BUILTIN(__builtin_amdgcn_image_store_3d_v4f16_i32, "vV4hiiiiV8iii", "nc", "")
974+
TARGET_BUILTIN(__builtin_amdgcn_image_store_cube_v4f32_i32, "vV4fiiiiV8iii", "nc", "")
975+
TARGET_BUILTIN(__builtin_amdgcn_image_store_cube_v4f16_i32, "vV4hiiiiV8iii", "nc", "")
976+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1d_v4f32_i32, "vV4fiiiV8iii", "nc", "")
977+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1d_v4f16_i32, "vV4hiiiV8iii", "nc", "")
978+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1darray_v4f32_i32, "vV4fiiiiV8iii", "nc", "")
979+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1darray_v4f16_i32, "vV4hiiiiV8iii", "nc", "")
980+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_f32_i32, "vfiiiiV8iii", "nc", "")
981+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_v4f32_i32, "vV4fiiiiV8iii", "nc", "")
982+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_v4f16_i32, "vV4hiiiiV8iii", "nc", "")
983+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_f32_i32, "vfiiiiiV8iii", "nc", "")
984+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_v4f32_i32, "vV4fiiiiiV8iii", "nc", "")
985+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_v4f16_i32, "vV4hiiiiiV8iii", "nc", "")
986+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f32_i32, "vV4fiiiiiV8iii", "nc", "")
987+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f16_i32, "vV4hiiiiiV8iii", "nc", "")
988+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f32_i32, "vV4fiiiiiV8iii", "nc", "")
989+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f16_i32, "vV4hiiiiiV8iii", "nc", "")
990+
962991
#undef BUILTIN
963992
#undef TARGET_BUILTIN

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

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

34-
bool checkImageImmArgFunctionCall(CallExpr *TheCall, unsigned ArgCount);
35-
3634
/// Create an AMDGPUWavesPerEUAttr attribute.
3735
AMDGPUFlatWorkGroupSizeAttr *
3836
CreateAMDGPUFlatWorkGroupSizeAttr(const AttributeCommonInfo &CI, Expr *Min,

‎clang/lib/Sema/SemaAMDGPU.cpp‎

Lines changed: 48 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -240,8 +240,54 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
240240
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
241241
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32: {
242242
unsigned ArgCount = TheCall->getNumArgs() - 1;
243-
244-
return checkImageImmArgFunctionCall(TheCall, ArgCount);
243+
llvm::APSInt Result;
244+
bool isImmArg =
245+
(!(SemaRef.BuiltinConstantArg(TheCall, 0, Result)) &&
246+
!(SemaRef.BuiltinConstantArg(TheCall, ArgCount, Result)) &&
247+
!(SemaRef.BuiltinConstantArg(TheCall, (ArgCount - 1), Result)))
248+
? false
249+
: true;
250+
251+
return isImmArg;
252+
}
253+
case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:
254+
case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:
255+
case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32:
256+
case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:
257+
case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32:
258+
case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32:
259+
case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32:
260+
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32:
261+
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32:
262+
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32:
263+
case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:
264+
case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:
265+
case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:
266+
case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:
267+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:
268+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32:
269+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:
270+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:
271+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:
272+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:
273+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32:
274+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32:
275+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32:
276+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32:
277+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:
278+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:
279+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
280+
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32: {
281+
unsigned ArgCount = TheCall->getNumArgs() - 1;
282+
llvm::APSInt Result;
283+
bool isImmArg =
284+
(!(SemaRef.BuiltinConstantArg(TheCall, 1, Result)) &&
285+
!(SemaRef.BuiltinConstantArg(TheCall, ArgCount, Result)) &&
286+
!(SemaRef.BuiltinConstantArg(TheCall, (ArgCount - 1), Result)))
287+
? false
288+
: true;
289+
290+
return isImmArg;
245291
}
246292
default:
247293
return false;

0 commit comments

Comments
 (0)