llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-amdgpu Author: Matt Arsenault (arsenm) <details> <summary>Changes</summary> I believe this is dead since clover was removed from mesa. Creating all of the targets as aliases was also busted, but may not have mattered due to how incomplete the backend is. --- Patch is 42.76 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/181976.diff 33 Files Affected: - (modified) libclc/CMakeLists.txt (+2-7) - (modified) libclc/clc/lib/amdgpu/math/clc_sqrt.cl (+1-7) - (removed) libclc/clc/lib/r600/SOURCES (-4) - (removed) libclc/clc/lib/r600/math/clc_fma.cl (-14) - (removed) libclc/clc/lib/r600/math/clc_fma.inc (-16) - (removed) libclc/clc/lib/r600/math/clc_native_rsqrt.cl (-18) - (removed) libclc/clc/lib/r600/math/clc_rsqrt.cl (-27) - (removed) libclc/clc/lib/r600/math/clc_sw_fma.cl (-174) - (removed) libclc/opencl/lib/r600/SOURCES (-8) - (removed) libclc/opencl/lib/r600/SOURCES_3.9 (-15) - (removed) libclc/opencl/lib/r600/image/get_image_attributes_impl.ll (-95) - (removed) libclc/opencl/lib/r600/image/get_image_channel_data_type.cl (-19) - (removed) libclc/opencl/lib/r600/image/get_image_channel_order.cl (-19) - (removed) libclc/opencl/lib/r600/image/get_image_depth.cl (-15) - (removed) libclc/opencl/lib/r600/image/get_image_dim.cl (-17) - (removed) libclc/opencl/lib/r600/image/get_image_height.cl (-19) - (removed) libclc/opencl/lib/r600/image/get_image_width.cl (-19) - (removed) libclc/opencl/lib/r600/image/read_image_impl.ll (-54) - (removed) libclc/opencl/lib/r600/image/read_imagef.cl (-22) - (removed) libclc/opencl/lib/r600/image/read_imagei.cl (-31) - (removed) libclc/opencl/lib/r600/image/read_imageui.cl (-31) - (removed) libclc/opencl/lib/r600/image/write_image_impl.ll (-60) - (removed) libclc/opencl/lib/r600/image/write_imagef.cl (-16) - (removed) libclc/opencl/lib/r600/image/write_imagei.cl (-16) - (removed) libclc/opencl/lib/r600/image/write_imageui.cl (-16) - (removed) libclc/opencl/lib/r600/synchronization/barrier.cl (-16) - (removed) libclc/opencl/lib/r600/workitem/get_global_offset.cl (-18) - (removed) libclc/opencl/lib/r600/workitem/get_global_size.cl (-26) - (removed) libclc/opencl/lib/r600/workitem/get_group_id.cl (-22) - (removed) libclc/opencl/lib/r600/workitem/get_local_id.cl (-22) - (removed) libclc/opencl/lib/r600/workitem/get_local_size.cl (-26) - (removed) libclc/opencl/lib/r600/workitem/get_num_groups.cl (-26) - (removed) libclc/opencl/lib/r600/workitem/get_work_dim.cl (-16) ``````````diff diff --git a/libclc/CMakeLists.txt b/libclc/CMakeLists.txt index 6430c4bf9c2be..77c3dca51d3d5 100644 --- a/libclc/CMakeLists.txt +++ b/libclc/CMakeLists.txt @@ -24,7 +24,6 @@ set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS opencl/lib/clspv/SOURCES; opencl/lib/generic/SOURCES; opencl/lib/ptx-nvidiacl/SOURCES; - opencl/lib/r600/SOURCES; opencl/lib/spirv/SOURCES; # CLC internal libraries clc/lib/generic/SOURCES; @@ -32,7 +31,6 @@ set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS clc/lib/amdgpu/SOURCES; clc/lib/clspv/SOURCES; clc/lib/ptx-nvidiacl/SOURCES; - clc/lib/r600/SOURCES; clc/lib/spirv/SOURCES; ) @@ -151,7 +149,6 @@ set( LIBCLC_TARGETS_ALL amdgcn-amd-amdhsa-llvm clspv-- clspv64-- - r600-- nvptx64-- nvptx64--nvidiacl nvptx64-nvidia-cuda @@ -198,7 +195,6 @@ list( SORT LIBCLC_TARGETS_TO_BUILD ) include_directories( ${LLVM_INCLUDE_DIRS} ) # Setup arch devices -set( r600--_devices cedar cypress barts cayman ) set( amdgcn--_devices none ) set( amdgcn-mesa-mesa3d_devices none ) set( amdgcn-amd-amdhsa-llvm_devices none ) @@ -253,7 +249,6 @@ set_source_files_properties( ${CMAKE_CURRENT_SOURCE_DIR}/clc/lib/amdgpu/math/clc_native_exp2.cl ${CMAKE_CURRENT_SOURCE_DIR}/clc/lib/amdgpu/math/clc_native_exp.cl ${CMAKE_CURRENT_SOURCE_DIR}/clc/lib/amdgpu/math/clc_native_log10.cl - ${CMAKE_CURRENT_SOURCE_DIR}/clc/lib/r600/math/clc_native_rsqrt.cl # OpenCL builtins ${CMAKE_CURRENT_SOURCE_DIR}/opencl/lib/generic/math/native_cos.cl ${CMAKE_CURRENT_SOURCE_DIR}/opencl/lib/generic/math/native_divide.cl @@ -283,7 +278,7 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} ) set( opencl_dirs ) - if( ${ARCH} STREQUAL r600 OR ${ARCH} STREQUAL amdgcn ) + if( ${ARCH} STREQUAL amdgcn ) list( APPEND opencl_dirs amdgpu ) endif() @@ -414,7 +409,7 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} ) # maps to the private address space. set ( private_addrspace_val 0 ) set ( generic_addrspace_val 0 ) - if( ARCH STREQUAL amdgcn OR ARCH STREQUAL r600 OR ARCH STREQUAL amdgcn-amdhsa ) + if( ARCH STREQUAL amdgcn-amdhsa ) set ( private_addrspace_val 5 ) endif() if( ARCH STREQUAL spirv OR ARCH STREQUAL spirv64) diff --git a/libclc/clc/lib/amdgpu/math/clc_sqrt.cl b/libclc/clc/lib/amdgpu/math/clc_sqrt.cl index 143afa5ece56d..3d13b3edfe48c 100644 --- a/libclc/clc/lib/amdgpu/math/clc_sqrt.cl +++ b/libclc/clc/lib/amdgpu/math/clc_sqrt.cl @@ -16,19 +16,13 @@ #pragma OPENCL EXTENSION cl_khr_fp64 : enable -#ifdef __AMDGCN__ -#define __clc_builtin_rsq __builtin_amdgcn_rsq -#else -#define __clc_builtin_rsq __builtin_r600_recipsqrt_ieee -#endif - _CLC_OVERLOAD _CLC_DEF double __clc_sqrt(double x) { uint vcc = x < 0x1p-767; uint exp0 = vcc ? 0x100 : 0; unsigned exp1 = vcc ? 0xffffff80 : 0; double v01 = __clc_ldexp(x, exp0); - double v23 = __clc_builtin_rsq(v01); + double v23 = __builtin_amdgcn_rsq(v01); double v45 = v01 * v23; v23 = v23 * 0.5; diff --git a/libclc/clc/lib/r600/SOURCES b/libclc/clc/lib/r600/SOURCES deleted file mode 100644 index c60ac1e2b043e..0000000000000 --- a/libclc/clc/lib/r600/SOURCES +++ /dev/null @@ -1,4 +0,0 @@ -math/clc_fma.cl -math/clc_native_rsqrt.cl -math/clc_rsqrt.cl -math/clc_sw_fma.cl diff --git a/libclc/clc/lib/r600/math/clc_fma.cl b/libclc/clc/lib/r600/math/clc_fma.cl deleted file mode 100644 index e69ef614e780f..0000000000000 --- a/libclc/clc/lib/r600/math/clc_fma.cl +++ /dev/null @@ -1,14 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -#include <clc/internal/clc.h> -#include <clc/internal/math/clc_sw_fma.h> -#include <clc/math/math.h> - -#define __CLC_BODY <clc_fma.inc> -#include <clc/math/gentype.inc> diff --git a/libclc/clc/lib/r600/math/clc_fma.inc b/libclc/clc/lib/r600/math/clc_fma.inc deleted file mode 100644 index dec1adb66cf89..0000000000000 --- a/libclc/clc/lib/r600/math/clc_fma.inc +++ /dev/null @@ -1,16 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -_CLC_DEF _CLC_OVERLOAD __CLC_GENTYPE __clc_fma(__CLC_GENTYPE a, __CLC_GENTYPE b, - __CLC_GENTYPE c) { -#if __CLC_FPSIZE == 32 - return __clc_sw_fma(a, b, c); -#else - return __builtin_elementwise_fma(a, b, c); -#endif -} diff --git a/libclc/clc/lib/r600/math/clc_native_rsqrt.cl b/libclc/clc/lib/r600/math/clc_native_rsqrt.cl deleted file mode 100644 index cef106e3b4e97..0000000000000 --- a/libclc/clc/lib/r600/math/clc_native_rsqrt.cl +++ /dev/null @@ -1,18 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -#include <clc/internal/clc.h> - -_CLC_OVERLOAD _CLC_DEF float __clc_native_rsqrt(float x) { - return __builtin_r600_recipsqrt_ieeef(x); -} - -#define __CLC_FLOAT_ONLY -#define __CLC_FUNCTION __clc_native_rsqrt -#define __CLC_BODY <clc/shared/unary_def_scalarize.inc> -#include <clc/math/gentype.inc> diff --git a/libclc/clc/lib/r600/math/clc_rsqrt.cl b/libclc/clc/lib/r600/math/clc_rsqrt.cl deleted file mode 100644 index 806ee678e00e4..0000000000000 --- a/libclc/clc/lib/r600/math/clc_rsqrt.cl +++ /dev/null @@ -1,27 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -#include <clc/internal/clc.h> - -_CLC_OVERLOAD _CLC_DEF float __clc_rsqrt(float x) { - return __builtin_r600_recipsqrt_ieeef(x); -} - -#ifdef cl_khr_fp64 - -#pragma OPENCL EXTENSION cl_khr_fp64 : enable - -_CLC_OVERLOAD _CLC_DEF double __clc_rsqrt(double x) { - return __builtin_r600_recipsqrt_ieee(x); -} - -#endif // cl_khr_fp64 - -#define __CLC_FUNCTION __clc_rsqrt -#define __CLC_BODY <clc/shared/unary_def_scalarize.inc> -#include <clc/math/gentype.inc> diff --git a/libclc/clc/lib/r600/math/clc_sw_fma.cl b/libclc/clc/lib/r600/math/clc_sw_fma.cl deleted file mode 100644 index ab5418e569371..0000000000000 --- a/libclc/clc/lib/r600/math/clc_sw_fma.cl +++ /dev/null @@ -1,174 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -#include <clc/clc_as_type.h> -#include <clc/float/definitions.h> -#include <clc/integer/clc_abs.h> -#include <clc/integer/clc_clz.h> -#include <clc/integer/definitions.h> -#include <clc/internal/clc.h> -#include <clc/math/clc_mad.h> -#include <clc/math/math.h> -#include <clc/relational/clc_isinf.h> -#include <clc/relational/clc_isnan.h> -#include <clc/shared/clc_max.h> - -static _CLC_INLINE float __clc_flush_denormal(float x) { - int ix = __clc_as_int(x); - if (((ix & EXPBITS_SP32) == 0) && ((ix & MANTBITS_SP32) != 0)) { - ix &= SIGNBIT_SP32; - x = __clc_as_float(ix); - } - return x; -} - -struct fp { - ulong mantissa; - int exponent; - uint sign; -}; - -_CLC_DEF _CLC_OVERLOAD float __clc_sw_fma(float a, float b, float c) { - /* special cases */ - if (__clc_isnan(a) || __clc_isnan(b) || __clc_isnan(c) || __clc_isinf(a) || - __clc_isinf(b)) { - return __clc_mad(a, b, c); - } - - /* If only c is inf, and both a,b are regular numbers, the result is c*/ - if (__clc_isinf(c)) { - return c; - } - - a = __clc_flush_denormal(a); - b = __clc_flush_denormal(b); - c = __clc_flush_denormal(c); - - if (c == 0) { - return a * b; - } - - struct fp st_a, st_b, st_c; - - st_a.exponent = a == .0f ? 0 : ((__clc_as_uint(a) & 0x7f800000) >> 23) - 127; - st_b.exponent = b == .0f ? 0 : ((__clc_as_uint(b) & 0x7f800000) >> 23) - 127; - st_c.exponent = c == .0f ? 0 : ((__clc_as_uint(c) & 0x7f800000) >> 23) - 127; - - st_a.mantissa = a == .0f ? 0 : (__clc_as_uint(a) & 0x7fffff) | 0x800000; - st_b.mantissa = b == .0f ? 0 : (__clc_as_uint(b) & 0x7fffff) | 0x800000; - st_c.mantissa = c == .0f ? 0 : (__clc_as_uint(c) & 0x7fffff) | 0x800000; - - st_a.sign = __clc_as_uint(a) & 0x80000000; - st_b.sign = __clc_as_uint(b) & 0x80000000; - st_c.sign = __clc_as_uint(c) & 0x80000000; - - // Multiplication. - // Move the product to the highest bits to maximize precision - // mantissa is 24 bits => product is 48 bits, 2bits non-fraction. - // Add one bit for future addition overflow, - // add another bit to detect subtraction underflow - struct fp st_mul; - st_mul.sign = st_a.sign ^ st_b.sign; - st_mul.mantissa = (st_a.mantissa * st_b.mantissa) << 14ul; - st_mul.exponent = st_mul.mantissa ? st_a.exponent + st_b.exponent : 0; - - // FIXME: Detecting a == 0 || b == 0 above crashed GCN isel - if (st_mul.exponent == 0 && st_mul.mantissa == 0) - return c; - -// Mantissa is 23 fractional bits, shift it the same way as product mantissa -#define C_ADJUST 37ul - - // both exponents are bias adjusted - int exp_diff = st_mul.exponent - st_c.exponent; - - st_c.mantissa <<= C_ADJUST; - ulong cutoff_bits = 0; - ulong cutoff_mask = (1ul << __clc_abs(exp_diff)) - 1ul; - if (exp_diff > 0) { - cutoff_bits = - exp_diff >= 64 ? st_c.mantissa : (st_c.mantissa & cutoff_mask); - st_c.mantissa = exp_diff >= 64 ? 0 : (st_c.mantissa >> exp_diff); - } else { - cutoff_bits = - -exp_diff >= 64 ? st_mul.mantissa : (st_mul.mantissa & cutoff_mask); - st_mul.mantissa = -exp_diff >= 64 ? 0 : (st_mul.mantissa >> -exp_diff); - } - - struct fp st_fma; - st_fma.sign = st_mul.sign; - st_fma.exponent = __clc_max(st_mul.exponent, st_c.exponent); - if (st_c.sign == st_mul.sign) { - st_fma.mantissa = st_mul.mantissa + st_c.mantissa; - } else { - // cutoff bits borrow one - st_fma.mantissa = - st_mul.mantissa - st_c.mantissa - - (cutoff_bits && (st_mul.exponent > st_c.exponent) ? 1 : 0); - } - - // underflow: st_c.sign != st_mul.sign, and magnitude switches the sign - if (st_fma.mantissa > LONG_MAX) { - st_fma.mantissa = 0 - st_fma.mantissa; - st_fma.sign = st_mul.sign ^ 0x80000000; - } - - // detect overflow/underflow - int overflow_bits = 3 - __clc_clz(st_fma.mantissa); - - // adjust exponent - st_fma.exponent += overflow_bits; - - // handle underflow - if (overflow_bits < 0) { - st_fma.mantissa <<= -overflow_bits; - overflow_bits = 0; - } - - // rounding - ulong trunc_mask = (1ul << (C_ADJUST + overflow_bits)) - 1; - ulong trunc_bits = (st_fma.mantissa & trunc_mask) | (cutoff_bits != 0); - ulong last_bit = st_fma.mantissa & (1ul << (C_ADJUST + overflow_bits)); - ulong grs_bits = (0x4ul << (C_ADJUST - 3 + overflow_bits)); - - // round to nearest even - if ((trunc_bits > grs_bits) || (trunc_bits == grs_bits && last_bit != 0)) { - st_fma.mantissa += (1ul << (C_ADJUST + overflow_bits)); - } - - // Shift mantissa back to bit 23 - st_fma.mantissa = (st_fma.mantissa >> (C_ADJUST + overflow_bits)); - - // Detect rounding overflow - if (st_fma.mantissa > 0xffffff) { - ++st_fma.exponent; - st_fma.mantissa >>= 1; - } - - if (st_fma.mantissa == 0) { - return .0f; - } - - // Flating point range limit - if (st_fma.exponent > 127) { - return __clc_as_float(__clc_as_uint(INFINITY) | st_fma.sign); - } - - // Flush denormals - if (st_fma.exponent <= -127) { - return __clc_as_float(st_fma.sign); - } - - return __clc_as_float(st_fma.sign | ((st_fma.exponent + 127) << 23) | - ((uint)st_fma.mantissa & 0x7fffff)); -} - -#define __CLC_FLOAT_ONLY -#define __CLC_FUNCTION __clc_sw_fma -#define __CLC_BODY <clc/shared/ternary_def_scalarize.inc> -#include <clc/math/gentype.inc> diff --git a/libclc/opencl/lib/r600/SOURCES b/libclc/opencl/lib/r600/SOURCES deleted file mode 100644 index c4561274d8b2d..0000000000000 --- a/libclc/opencl/lib/r600/SOURCES +++ /dev/null @@ -1,8 +0,0 @@ -synchronization/barrier.cl -workitem/get_global_offset.cl -workitem/get_group_id.cl -workitem/get_global_size.cl -workitem/get_local_id.cl -workitem/get_local_size.cl -workitem/get_num_groups.cl -workitem/get_work_dim.cl diff --git a/libclc/opencl/lib/r600/SOURCES_3.9 b/libclc/opencl/lib/r600/SOURCES_3.9 deleted file mode 100644 index a44a9ce8074f8..0000000000000 --- a/libclc/opencl/lib/r600/SOURCES_3.9 +++ /dev/null @@ -1,15 +0,0 @@ -image/get_image_dim.cl -image/get_image_width.cl -image/get_image_height.cl -image/get_image_depth.cl -image/get_image_channel_data_type.cl -image/get_image_channel_order.cl -image/get_image_attributes_impl.ll -image/read_imagef.cl -image/read_imagei.cl -image/read_imageui.cl -image/read_image_impl.ll -image/write_imagef.cl -image/write_imagei.cl -image/write_imageui.cl -image/write_image_impl.ll diff --git a/libclc/opencl/lib/r600/image/get_image_attributes_impl.ll b/libclc/opencl/lib/r600/image/get_image_attributes_impl.ll deleted file mode 100644 index e7b605df62e43..0000000000000 --- a/libclc/opencl/lib/r600/image/get_image_attributes_impl.ll +++ /dev/null @@ -1,95 +0,0 @@ -;;===----------------------------------------------------------------------===;; -; -; Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -; See https://llvm.org/LICENSE.txt for license information. -; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -; -;;===----------------------------------------------------------------------===;; - -%opencl.image2d_t = type opaque -%opencl.image3d_t = type opaque - -declare i32 @llvm.OpenCL.image.get.resource.id.2d( - %opencl.image2d_t addrspace(1)*) nounwind readnone -declare i32 @llvm.OpenCL.image.get.resource.id.3d( - %opencl.image3d_t addrspace(1)*) nounwind readnone - -declare [3 x i32] @llvm.OpenCL.image.get.size.2d( - %opencl.image2d_t addrspace(1)*) nounwind readnone -declare [3 x i32] @llvm.OpenCL.image.get.size.3d( - %opencl.image3d_t addrspace(1)*) nounwind readnone - -declare [2 x i32] @llvm.OpenCL.image.get.format.2d( - %opencl.image2d_t addrspace(1)*) nounwind readnone -declare [2 x i32] @llvm.OpenCL.image.get.format.3d( - %opencl.image3d_t addrspace(1)*) nounwind readnone - -define i32 @__clc_get_image_width_2d( - %opencl.image2d_t addrspace(1)* nocapture %img) #0 { - %1 = tail call [3 x i32] @llvm.OpenCL.image.get.size.2d( - %opencl.image2d_t addrspace(1)* %img) - %2 = extractvalue [3 x i32] %1, 0 - ret i32 %2 -} -define i32 @__clc_get_image_width_3d( - %opencl.image3d_t addrspace(1)* nocapture %img) #0 { - %1 = tail call [3 x i32] @llvm.OpenCL.image.get.size.3d( - %opencl.image3d_t addrspace(1)* %img) - %2 = extractvalue [3 x i32] %1, 0 - ret i32 %2 -} - -define i32 @__clc_get_image_height_2d( - %opencl.image2d_t addrspace(1)* nocapture %img) #0 { - %1 = tail call [3 x i32] @llvm.OpenCL.image.get.size.2d( - %opencl.image2d_t addrspace(1)* %img) - %2 = extractvalue [3 x i32] %1, 1 - ret i32 %2 -} -define i32 @__clc_get_image_height_3d( - %opencl.image3d_t addrspace(1)* nocapture %img) #0 { - %1 = tail call [3 x i32] @llvm.OpenCL.image.get.size.3d( - %opencl.image3d_t addrspace(1)* %img) - %2 = extractvalue [3 x i32] %1, 1 - ret i32 %2 -} - -define i32 @__clc_get_image_depth_3d( - %opencl.image3d_t addrspace(1)* nocapture %img) #0 { - %1 = tail call [3 x i32] @llvm.OpenCL.image.get.size.3d( - %opencl.image3d_t addrspace(1)* %img) - %2 = extractvalue [3 x i32] %1, 2 - ret i32 %2 -} - -define i32 @__clc_get_image_channel_data_type_2d( - %opencl.image2d_t addrspace(1)* nocapture %img) #0 { - %1 = tail call [2 x i32] @llvm.OpenCL.image.get.format.2d( - %opencl.image2d_t addrspace(1)* %img) - %2 = extractvalue [2 x i32] %1, 0 - ret i32 %2 -} -define i32 @__clc_get_image_channel_data_type_3d( - %opencl.image3d_t addrspace(1)* nocapture %img) #0 { - %1 = tail call [2 x i32] @llvm.OpenCL.image.get.format.3d( - %opencl.image3d_t addrspace(1)* %img) - %2 = extractvalue [2 x i32] %1, 0 - ret i32 %2 -} - -define i32 @__clc_get_image_channel_order_2d( - %opencl.image2d_t addrspace(1)* nocapture %img) #0 { - %1 = tail call [2 x i32] @llvm.OpenCL.image.get.format.2d( - %opencl.image2d_t addrspace(1)* %img) - %2 = extractvalue [2 x i32] %1, 1 - ret i32 %2 -} -define i32 @__clc_get_image_channel_order_3d( - %opencl.image3d_t addrspace(1)* nocapture %img) #0 { - %1 = tail call [2 x i32] @llvm.OpenCL.image.get.format.3d( - %opencl.image3d_t addrspace(1)* %img) - %2 = extractvalue [2 x i32] %1, 1 - ret i32 %2 -} - -attributes #0 = { nounwind readnone alwaysinline } diff --git a/libclc/opencl/lib/r600/image/get_image_channel_data_type.cl b/libclc/opencl/lib/r600/image/get_image_channel_data_type.cl deleted file mode 100644 index 7fed814d3ef15..0000000000000 --- a/libclc/opencl/lib/r600/image/get_image_channel_data_type.cl +++ /dev/null @@ -1,19 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -#include <clc/opencl/opencl-base.h> - -_CLC_DECL int __clc_get_image_channel_data_type_2d(image2d_t); -_CLC_DECL int __clc_get_image_channel_data_type_3d(image3d_t); - -_CLC_OVERLOAD _CLC_DEF int get_image_channel_data_type(image2d_t image) { - return __clc_get_image_channel_data_type_2d(image); -} -_CLC_OVERLOAD _CLC_DEF int get_image_channel_data_type(image3d_t image) { - return __clc_get_image_channel_data_type_3d(image); -} diff --git a/libclc/opencl/lib/r600/image/get_image_channel_order.cl b/libclc/opencl/lib/r600/image/get_image_channel_order.cl deleted file mode 100644 index 1ad4dedb39be4..0000000000000 --- a/libclc/opencl/lib/r600/image/get_image_channel_order.cl +++ /dev/null @@ -1,19 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -#include <clc/opencl/opencl-base.h> - -_CLC_DECL int __clc_get_image_channel_order_2d(image2d_t); -_CLC_DECL int __clc_get_image_channel_order_3d(image3d_t); - -_CLC_OVERLOAD _CLC_DEF int get_image_channel_order(image2d_t image) { - return __clc_get_image_channel_order_2d(image); -} -_CLC_OVERLOAD _CLC_DEF int get_image_channel_order(image3d_t image) { - return __clc_get_image_channel_order_3d(image); -} diff --git a/libclc/opencl/lib/r600/image/get_image_depth.cl b/libclc/opencl/lib/r600/image/get_image_depth.cl deleted file mode 100644 index b7bb8c8b841eb..0000000000000 --- a/libclc/opencl/lib/r600/image/get_image_depth.cl +++ /dev/null @@ -1,15 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/181976 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
