llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-amdgpu Author: Rana Pratap Reddy (ranapratap55) <details> <summary>Changes</summary> For amdgcn_image_load/store_* builtins, using 'e' in the builtin def so that it will take `_Float16` for HIP/C++ and `__fp16` for OpenCL. --- Full diff: https://github.com/llvm/llvm-project/pull/166720.diff 5 Files Affected: - (modified) clang/include/clang/Basic/BuiltinsAMDGPU.def (+30-30) - (modified) clang/test/SemaOpenCL/builtins-image-load-param-gfx1100-err.cl (+2) - (modified) clang/test/SemaOpenCL/builtins-image-load-param-gfx942-err.cl (+2) - (modified) clang/test/SemaOpenCL/builtins-image-store-param-gfx1100-err.cl (+2) - (modified) clang/test/SemaOpenCL/builtins-image-store-param-gfx942-err.cl (+2) ``````````diff diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 36cb527a9c806..4d265df9b4b5f 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -898,75 +898,75 @@ TARGET_BUILTIN(__builtin_amdgcn_cooperative_atomic_store_8x16B, "vV4i*V4iIicC*", // Image builtins //===----------------------------------------------------------------------===// TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f32_i32, "V4fiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f16_i32, "V4hiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f16_i32, "V4eiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_1darray_v4f32_i32, "V4fiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_1darray_v4f16_i32, "V4hiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_1darray_v4f16_i32, "V4eiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_f32_i32, "fiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_v4f32_i32, "V4fiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_v4f16_i32, "V4hiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_v4f16_i32, "V4eiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_f32_i32, "fiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_v4f16_i32, "V4eiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_3d_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_3d_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_3d_v4f16_i32, "V4eiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_cube_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_cube_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_cube_v4f16_i32, "V4eiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1d_v4f32_i32, "V4fiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1d_v4f16_i32, "V4hiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1d_v4f16_i32, "V4eiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1darray_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1darray_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1darray_v4f16_i32, "V4eiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_f32_i32, "fiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_v4f16_i32, "V4eiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_f32_i32, "fiiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_v4f32_i32, "V4fiiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_v4f16_i32, "V4hiiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_v4f16_i32, "V4eiiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_3d_v4f32_i32, "V4fiiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_3d_v4f16_i32, "V4hiiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_3d_v4f16_i32, "V4eiiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_cube_v4f32_i32, "V4fiiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_cube_v4f16_i32, "V4hiiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_cube_v4f16_i32, "V4eiiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_1d_v4f32_i32, "vV4fiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_1d_v4f16_i32, "vV4hiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_1d_v4f16_i32, "vV4eiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_1darray_v4f32_i32, "vV4fiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_1darray_v4f16_i32, "vV4hiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_1darray_v4f16_i32, "vV4eiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_f32_i32, "vfiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_v4f32_i32, "vV4fiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_v4f16_i32, "vV4hiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_v4f16_i32, "vV4eiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_f32_i32, "vfiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_v4f16_i32, "vV4eiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_3d_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_3d_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_3d_v4f16_i32, "vV4eiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_cube_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_cube_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_cube_v4f16_i32, "vV4eiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1d_v4f32_i32, "vV4fiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1d_v4f16_i32, "vV4hiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1d_v4f16_i32, "vV4eiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1darray_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1darray_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1darray_v4f16_i32, "vV4eiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_f32_i32, "vfiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_v4f16_i32, "vV4eiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_f32_i32, "vfiiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_v4f32_i32, "vV4fiiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_v4f16_i32, "vV4hiiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_v4f16_i32, "vV4eiiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f32_i32, "vV4fiiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f16_i32, "vV4hiiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f16_i32, "vV4eiiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f32_i32, "vV4fiiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f16_i32, "vV4hiiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f16_i32, "vV4eiiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_sample_1d_v4f32_f32, "V4fifQtV4ibii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_sample_1d_v4f16_f32, "V4hifQtV4ibii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_sample_1d_v4f16_f32, "V4eifQtV4ibii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_sample_1darray_v4f32_f32, "V4fiffQtV4ibii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_sample_1darray_v4f16_f32, "V4hiffQtV4ibii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_sample_1darray_v4f16_f32, "V4eiffQtV4ibii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_f32_f32, "fiffQtV4ibii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_v4f32_f32, "V4fiffQtV4ibii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_v4f16_f32, "V4hiffQtV4ibii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_v4f16_f32, "V4eiffQtV4ibii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_f32_f32, "fifffQtV4ibii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_v4f32_f32, "V4fifffQtV4ibii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_v4f16_f32, "V4hifffQtV4ibii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_v4f16_f32, "V4eifffQtV4ibii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_sample_3d_v4f32_f32, "V4fifffQtV4ibii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_sample_3d_v4f16_f32, "V4hifffQtV4ibii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_sample_3d_v4f16_f32, "V4eifffQtV4ibii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_sample_cube_v4f32_f32, "V4fifffQtV4ibii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_sample_cube_v4f16_f32, "V4hifffQtV4ibii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_sample_cube_v4f16_f32, "V4eifffQtV4ibii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_sample_lz_1d_v4f32_f32, "V4fifQtV4ibii", "nc", "extended-image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_sample_lz_1d_v4f16_f32, "V4eifQtV4ibii", "nc", "extended-image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_sample_lz_1darray_v4f32_f32, "V4fiffQtV4ibii", "nc", "extended-image-insts") diff --git a/clang/test/SemaOpenCL/builtins-image-load-param-gfx1100-err.cl b/clang/test/SemaOpenCL/builtins-image-load-param-gfx1100-err.cl index 8f609dcbd34f2..b8b629f91623a 100644 --- a/clang/test/SemaOpenCL/builtins-image-load-param-gfx1100-err.cl +++ b/clang/test/SemaOpenCL/builtins-image-load-param-gfx1100-err.cl @@ -1,6 +1,8 @@ // RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1100 -S -verify=expected -o - %s // REQUIRES: amdgpu-registered-target +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + typedef int int4 __attribute__((ext_vector_type(4))); typedef float float4 __attribute__((ext_vector_type(4))); typedef half half4 __attribute__((ext_vector_type(4))); diff --git a/clang/test/SemaOpenCL/builtins-image-load-param-gfx942-err.cl b/clang/test/SemaOpenCL/builtins-image-load-param-gfx942-err.cl index b8780024f1076..dbf2693592f43 100644 --- a/clang/test/SemaOpenCL/builtins-image-load-param-gfx942-err.cl +++ b/clang/test/SemaOpenCL/builtins-image-load-param-gfx942-err.cl @@ -1,6 +1,8 @@ // RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx942 -verify -S -o - %s // REQUIRES: amdgpu-registered-target +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + typedef int int4 __attribute__((ext_vector_type(4))); typedef float float4 __attribute__((ext_vector_type(4))); typedef half half4 __attribute__((ext_vector_type(4))); diff --git a/clang/test/SemaOpenCL/builtins-image-store-param-gfx1100-err.cl b/clang/test/SemaOpenCL/builtins-image-store-param-gfx1100-err.cl index 4f6347e1c5286..f28d2ca34b77c 100644 --- a/clang/test/SemaOpenCL/builtins-image-store-param-gfx1100-err.cl +++ b/clang/test/SemaOpenCL/builtins-image-store-param-gfx1100-err.cl @@ -1,6 +1,8 @@ // RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1100 -S -verify=expected -o - %s // REQUIRES: amdgpu-registered-target +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + typedef float float4 __attribute__((ext_vector_type(4))); typedef half half4 __attribute__((ext_vector_type(4))); diff --git a/clang/test/SemaOpenCL/builtins-image-store-param-gfx942-err.cl b/clang/test/SemaOpenCL/builtins-image-store-param-gfx942-err.cl index d0085e5403b5f..39ca3b7ecfa6f 100644 --- a/clang/test/SemaOpenCL/builtins-image-store-param-gfx942-err.cl +++ b/clang/test/SemaOpenCL/builtins-image-store-param-gfx942-err.cl @@ -1,6 +1,8 @@ // RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx942 -S -verify -o - %s // REQUIRES: amdgpu-registered-target +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + typedef float float4 __attribute__((ext_vector_type(4))); typedef half half4 __attribute__((ext_vector_type(4))); `````````` </details> https://github.com/llvm/llvm-project/pull/166720 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
