https://github.com/ranapratap55 created https://github.com/llvm/llvm-project/pull/166720
None >From 039276238f48949ab7254a05dbe851507ebf107c Mon Sep 17 00:00:00 2001 From: N Rana Pratap Reddy <[email protected]> Date: Thu, 6 Nov 2025 12:10:50 +0530 Subject: [PATCH] [AMDGPU] Modifies builtin def to take _Float16 for HIP/C++ and __fp16 for openCL --- clang/include/clang/Basic/BuiltinsAMDGPU.def | 60 +++++++++---------- .../builtins-image-load-param-gfx1100-err.cl | 2 + .../builtins-image-load-param-gfx942-err.cl | 2 + .../builtins-image-store-param-gfx1100-err.cl | 2 + .../builtins-image-store-param-gfx942-err.cl | 2 + 5 files changed, 38 insertions(+), 30 deletions(-) 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))); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
