https://github.com/arsenm created https://github.com/llvm/llvm-project/pull/181976
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. >From fddef777a60d4b22fb0d766301c983b8bf306842 Mon Sep 17 00:00:00 2001 From: Matt Arsenault <[email protected]> Date: Wed, 18 Feb 2026 10:01:50 +0100 Subject: [PATCH] libclc: Remove r600 support 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. --- libclc/CMakeLists.txt | 9 +- libclc/clc/lib/amdgpu/math/clc_sqrt.cl | 8 +- libclc/clc/lib/r600/SOURCES | 4 - libclc/clc/lib/r600/math/clc_fma.cl | 14 -- libclc/clc/lib/r600/math/clc_fma.inc | 16 -- libclc/clc/lib/r600/math/clc_native_rsqrt.cl | 18 -- libclc/clc/lib/r600/math/clc_rsqrt.cl | 27 --- libclc/clc/lib/r600/math/clc_sw_fma.cl | 174 ------------------ libclc/opencl/lib/r600/SOURCES | 8 - libclc/opencl/lib/r600/SOURCES_3.9 | 15 -- .../r600/image/get_image_attributes_impl.ll | 95 ---------- .../r600/image/get_image_channel_data_type.cl | 19 -- .../lib/r600/image/get_image_channel_order.cl | 19 -- .../opencl/lib/r600/image/get_image_depth.cl | 15 -- libclc/opencl/lib/r600/image/get_image_dim.cl | 17 -- .../opencl/lib/r600/image/get_image_height.cl | 19 -- .../opencl/lib/r600/image/get_image_width.cl | 19 -- .../opencl/lib/r600/image/read_image_impl.ll | 54 ------ libclc/opencl/lib/r600/image/read_imagef.cl | 22 --- libclc/opencl/lib/r600/image/read_imagei.cl | 31 ---- libclc/opencl/lib/r600/image/read_imageui.cl | 31 ---- .../opencl/lib/r600/image/write_image_impl.ll | 60 ------ libclc/opencl/lib/r600/image/write_imagef.cl | 16 -- libclc/opencl/lib/r600/image/write_imagei.cl | 16 -- libclc/opencl/lib/r600/image/write_imageui.cl | 16 -- .../lib/r600/synchronization/barrier.cl | 16 -- .../lib/r600/workitem/get_global_offset.cl | 18 -- .../lib/r600/workitem/get_global_size.cl | 26 --- .../opencl/lib/r600/workitem/get_group_id.cl | 22 --- .../opencl/lib/r600/workitem/get_local_id.cl | 22 --- .../lib/r600/workitem/get_local_size.cl | 26 --- .../lib/r600/workitem/get_num_groups.cl | 26 --- .../opencl/lib/r600/workitem/get_work_dim.cl | 16 -- 33 files changed, 3 insertions(+), 911 deletions(-) delete mode 100644 libclc/clc/lib/r600/SOURCES delete mode 100644 libclc/clc/lib/r600/math/clc_fma.cl delete mode 100644 libclc/clc/lib/r600/math/clc_fma.inc delete mode 100644 libclc/clc/lib/r600/math/clc_native_rsqrt.cl delete mode 100644 libclc/clc/lib/r600/math/clc_rsqrt.cl delete mode 100644 libclc/clc/lib/r600/math/clc_sw_fma.cl delete mode 100644 libclc/opencl/lib/r600/SOURCES delete mode 100644 libclc/opencl/lib/r600/SOURCES_3.9 delete mode 100644 libclc/opencl/lib/r600/image/get_image_attributes_impl.ll delete mode 100644 libclc/opencl/lib/r600/image/get_image_channel_data_type.cl delete mode 100644 libclc/opencl/lib/r600/image/get_image_channel_order.cl delete mode 100644 libclc/opencl/lib/r600/image/get_image_depth.cl delete mode 100644 libclc/opencl/lib/r600/image/get_image_dim.cl delete mode 100644 libclc/opencl/lib/r600/image/get_image_height.cl delete mode 100644 libclc/opencl/lib/r600/image/get_image_width.cl delete mode 100644 libclc/opencl/lib/r600/image/read_image_impl.ll delete mode 100644 libclc/opencl/lib/r600/image/read_imagef.cl delete mode 100644 libclc/opencl/lib/r600/image/read_imagei.cl delete mode 100644 libclc/opencl/lib/r600/image/read_imageui.cl delete mode 100644 libclc/opencl/lib/r600/image/write_image_impl.ll delete mode 100644 libclc/opencl/lib/r600/image/write_imagef.cl delete mode 100644 libclc/opencl/lib/r600/image/write_imagei.cl delete mode 100644 libclc/opencl/lib/r600/image/write_imageui.cl delete mode 100644 libclc/opencl/lib/r600/synchronization/barrier.cl delete mode 100644 libclc/opencl/lib/r600/workitem/get_global_offset.cl delete mode 100644 libclc/opencl/lib/r600/workitem/get_global_size.cl delete mode 100644 libclc/opencl/lib/r600/workitem/get_group_id.cl delete mode 100644 libclc/opencl/lib/r600/workitem/get_local_id.cl delete mode 100644 libclc/opencl/lib/r600/workitem/get_local_size.cl delete mode 100644 libclc/opencl/lib/r600/workitem/get_num_groups.cl delete mode 100644 libclc/opencl/lib/r600/workitem/get_work_dim.cl 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-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -#include <clc/opencl/opencl-base.h> - -_CLC_DECL int __clc_get_image_depth_3d(image3d_t); - -_CLC_OVERLOAD _CLC_DEF int get_image_depth(image3d_t image) { - return __clc_get_image_depth_3d(image); -} diff --git a/libclc/opencl/lib/r600/image/get_image_dim.cl b/libclc/opencl/lib/r600/image/get_image_dim.cl deleted file mode 100644 index 91986156c1735..0000000000000 --- a/libclc/opencl/lib/r600/image/get_image_dim.cl +++ /dev/null @@ -1,17 +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_OVERLOAD _CLC_DEF int2 get_image_dim(image2d_t image) { - return (int2)(get_image_width(image), get_image_height(image)); -} -_CLC_OVERLOAD _CLC_DEF int4 get_image_dim(image3d_t image) { - return (int4)(get_image_width(image), get_image_height(image), - get_image_depth(image), 0); -} diff --git a/libclc/opencl/lib/r600/image/get_image_height.cl b/libclc/opencl/lib/r600/image/get_image_height.cl deleted file mode 100644 index cfdfa8972d31d..0000000000000 --- a/libclc/opencl/lib/r600/image/get_image_height.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_height_2d(image2d_t); -_CLC_DECL int __clc_get_image_height_3d(image3d_t); - -_CLC_OVERLOAD _CLC_DEF int get_image_height(image2d_t image) { - return __clc_get_image_height_2d(image); -} -_CLC_OVERLOAD _CLC_DEF int get_image_height(image3d_t image) { - return __clc_get_image_height_3d(image); -} diff --git a/libclc/opencl/lib/r600/image/get_image_width.cl b/libclc/opencl/lib/r600/image/get_image_width.cl deleted file mode 100644 index eb7bd73ea2f56..0000000000000 --- a/libclc/opencl/lib/r600/image/get_image_width.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_width_2d(image2d_t); -_CLC_DECL int __clc_get_image_width_3d(image3d_t); - -_CLC_OVERLOAD _CLC_DEF int get_image_width(image2d_t image) { - return __clc_get_image_width_2d(image); -} -_CLC_OVERLOAD _CLC_DEF int get_image_width(image3d_t image) { - return __clc_get_image_width_3d(image); -} diff --git a/libclc/opencl/lib/r600/image/read_image_impl.ll b/libclc/opencl/lib/r600/image/read_image_impl.ll deleted file mode 100644 index 00b80b63f4f7d..0000000000000 --- a/libclc/opencl/lib/r600/image/read_image_impl.ll +++ /dev/null @@ -1,54 +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 - -declare <4 x float> @llvm.R600.tex(<4 x float>, i32, i32, i32, i32, i32, i32, - i32, i32, i32) readnone -declare i32 @llvm.OpenCL.image.get.resource.id.2d( - %opencl.image2d_t addrspace(1)*) nounwind readnone -declare i32 @llvm.OpenCL.sampler.get.resource.id(i32) readnone - -define <4 x float> @__clc_v4f_from_v2f(<2 x float> %v) alwaysinline { - %e0 = extractelement <2 x float> %v, i32 0 - %e1 = extractelement <2 x float> %v, i32 1 - %res.0 = insertelement <4 x float> poison, float %e0, i32 0 - %res.1 = insertelement <4 x float> %res.0, float %e1, i32 1 - %res.2 = insertelement <4 x float> %res.1, float 0.0, i32 2 - %res.3 = insertelement <4 x float> %res.2, float 0.0, i32 3 - ret <4 x float> %res.3 -} - -define <4 x float> @__clc_read_imagef_tex( - %opencl.image2d_t addrspace(1)* nocapture %img, - i32 %sampler, <2 x float> %coord) alwaysinline { -entry: - %coord_v4 = call <4 x float> @__clc_v4f_from_v2f(<2 x float> %coord) - %smp_id = call i32 @llvm.OpenCL.sampler.get.resource.id(i32 %sampler) - %img_id = call i32 @llvm.OpenCL.image.get.resource.id.2d( - %opencl.image2d_t addrspace(1)* %img) - %tex_id = add i32 %img_id, 2 ; First 2 IDs are reserved. - - %coord_norm = and i32 %sampler, 1 - %is_norm = icmp eq i32 %coord_norm, 1 - br i1 %is_norm, label %NormCoord, label %UnnormCoord -NormCoord: - %data.norm = call <4 x float> @llvm.R600.tex( - <4 x float> %coord_v4, - i32 0, i32 0, i32 0, ; Offset. - i32 2, i32 %smp_id, - i32 1, i32 1, i32 1, i32 1) ; Normalized coords. - ret <4 x float> %data.norm -UnnormCoord: - %data.unnorm = call <4 x float> @llvm.R600.tex( - <4 x float> %coord_v4, - i32 0, i32 0, i32 0, ; Offset. - i32 %tex_id, i32 %smp_id, - i32 0, i32 0, i32 0, i32 0) ; Unnormalized coords. - ret <4 x float> %data.unnorm -} diff --git a/libclc/opencl/lib/r600/image/read_imagef.cl b/libclc/opencl/lib/r600/image/read_imagef.cl deleted file mode 100644 index 263972f7c22fc..0000000000000 --- a/libclc/opencl/lib/r600/image/read_imagef.cl +++ /dev/null @@ -1,22 +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 float4 __clc_read_imagef_tex(image2d_t, sampler_t, float2); - -_CLC_OVERLOAD _CLC_DEF float4 read_imagef(image2d_t image, sampler_t sampler, - int2 coord) { - float2 coord_float = (float2)(coord.x, coord.y); - return __clc_read_imagef_tex(image, sampler, coord_float); -} - -_CLC_OVERLOAD _CLC_DEF float4 read_imagef(image2d_t image, sampler_t sampler, - float2 coord) { - return __clc_read_imagef_tex(image, sampler, coord); -} diff --git a/libclc/opencl/lib/r600/image/read_imagei.cl b/libclc/opencl/lib/r600/image/read_imagei.cl deleted file mode 100644 index 24ec6b1a6661a..0000000000000 --- a/libclc/opencl/lib/r600/image/read_imagei.cl +++ /dev/null @@ -1,31 +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 float4 __clc_read_imagef_tex(image2d_t, sampler_t, float2); - -int4 __clc_reinterpret_v4f_to_v4i(float4 v) { - union { - int4 v4i; - float4 v4f; - } res = {.v4f = v}; - return res.v4i; -} - -_CLC_OVERLOAD _CLC_DEF int4 read_imagei(image2d_t image, sampler_t sampler, - int2 coord) { - float2 coord_float = (float2)(coord.x, coord.y); - return __clc_reinterpret_v4f_to_v4i( - __clc_read_imagef_tex(image, sampler, coord_float)); -} -_CLC_OVERLOAD _CLC_DEF int4 read_imagei(image2d_t image, sampler_t sampler, - float2 coord) { - return __clc_reinterpret_v4f_to_v4i( - __clc_read_imagef_tex(image, sampler, coord)); -} diff --git a/libclc/opencl/lib/r600/image/read_imageui.cl b/libclc/opencl/lib/r600/image/read_imageui.cl deleted file mode 100644 index 30bbe766d671c..0000000000000 --- a/libclc/opencl/lib/r600/image/read_imageui.cl +++ /dev/null @@ -1,31 +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 float4 __clc_read_imagef_tex(image2d_t, sampler_t, float2); - -uint4 __clc_reinterpret_v4f_to_v4ui(float4 v) { - union { - uint4 v4ui; - float4 v4f; - } res = {.v4f = v}; - return res.v4ui; -} - -_CLC_OVERLOAD _CLC_DEF uint4 read_imageui(image2d_t image, sampler_t sampler, - int2 coord) { - float2 coord_float = (float2)(coord.x, coord.y); - return __clc_reinterpret_v4f_to_v4ui( - __clc_read_imagef_tex(image, sampler, coord_float)); -} -_CLC_OVERLOAD _CLC_DEF uint4 read_imageui(image2d_t image, sampler_t sampler, - float2 coord) { - return __clc_reinterpret_v4f_to_v4ui( - __clc_read_imagef_tex(image, sampler, coord)); -} diff --git a/libclc/opencl/lib/r600/image/write_image_impl.ll b/libclc/opencl/lib/r600/image/write_image_impl.ll deleted file mode 100644 index 0759f0e2f9449..0000000000000 --- a/libclc/opencl/lib/r600/image/write_image_impl.ll +++ /dev/null @@ -1,60 +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 void @llvm.r600.rat.store.typed(<4 x i32> %color, <4 x i32> %coord, i32 %rat_id) - -define void @__clc_write_imageui_2d( - %opencl.image2d_t addrspace(1)* nocapture %img, - <2 x i32> %coord, <4 x i32> %color) #0 { - - ; Coordinate int2 -> int4. - %e0 = extractelement <2 x i32> %coord, i32 0 - %e1 = extractelement <2 x i32> %coord, i32 1 - %coord.0 = insertelement <4 x i32> poison, i32 %e0, i32 0 - %coord.1 = insertelement <4 x i32> %coord.0, i32 %e1, i32 1 - %coord.2 = insertelement <4 x i32> %coord.1, i32 0, i32 2 - %coord.3 = insertelement <4 x i32> %coord.2, i32 0, i32 3 - - ; Get RAT ID. - %img_id = call i32 @llvm.OpenCL.image.get.resource.id.2d( - %opencl.image2d_t addrspace(1)* %img) - %rat_id = add i32 %img_id, 1 - - ; Call store intrinsic. - call void @llvm.r600.rat.store.typed(<4 x i32> %color, <4 x i32> %coord.3, i32 %rat_id) - ret void -} - -define void @__clc_write_imagei_2d( - %opencl.image2d_t addrspace(1)* nocapture %img, - <2 x i32> %coord, <4 x i32> %color) #0 { - call void @__clc_write_imageui_2d( - %opencl.image2d_t addrspace(1)* nocapture %img, - <2 x i32> %coord, <4 x i32> %color) - ret void -} - -define void @__clc_write_imagef_2d( - %opencl.image2d_t addrspace(1)* nocapture %img, - <2 x i32> %coord, <4 x float> %color) #0 { - %color.i32 = bitcast <4 x float> %color to <4 x i32> - call void @__clc_write_imageui_2d( - %opencl.image2d_t addrspace(1)* nocapture %img, - <2 x i32> %coord, <4 x i32> %color.i32) - ret void -} - -attributes #0 = { alwaysinline } diff --git a/libclc/opencl/lib/r600/image/write_imagef.cl b/libclc/opencl/lib/r600/image/write_imagef.cl deleted file mode 100644 index 85d9a0bce86d9..0000000000000 --- a/libclc/opencl/lib/r600/image/write_imagef.cl +++ /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 -// -//===----------------------------------------------------------------------===// - -#include <clc/opencl/opencl-base.h> - -_CLC_DECL void __clc_write_imagef_2d(image2d_t image, int2 coord, float4 color); - -_CLC_OVERLOAD _CLC_DEF void write_imagef(image2d_t image, int2 coord, - float4 color) { - __clc_write_imagef_2d(image, coord, color); -} diff --git a/libclc/opencl/lib/r600/image/write_imagei.cl b/libclc/opencl/lib/r600/image/write_imagei.cl deleted file mode 100644 index 73bfb94644a7e..0000000000000 --- a/libclc/opencl/lib/r600/image/write_imagei.cl +++ /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 -// -//===----------------------------------------------------------------------===// - -#include <clc/opencl/opencl-base.h> - -_CLC_DECL void __clc_write_imagei_2d(image2d_t image, int2 coord, int4 color); - -_CLC_OVERLOAD _CLC_DEF void write_imagei(image2d_t image, int2 coord, - int4 color) { - __clc_write_imagei_2d(image, coord, color); -} diff --git a/libclc/opencl/lib/r600/image/write_imageui.cl b/libclc/opencl/lib/r600/image/write_imageui.cl deleted file mode 100644 index c2f477d39b039..0000000000000 --- a/libclc/opencl/lib/r600/image/write_imageui.cl +++ /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 -// -//===----------------------------------------------------------------------===// - -#include <clc/opencl/opencl-base.h> - -_CLC_DECL void __clc_write_imageui_2d(image2d_t image, int2 coord, uint4 color); - -_CLC_OVERLOAD _CLC_DEF void write_imageui(image2d_t image, int2 coord, - uint4 color) { - __clc_write_imageui_2d(image, coord, color); -} diff --git a/libclc/opencl/lib/r600/synchronization/barrier.cl b/libclc/opencl/lib/r600/synchronization/barrier.cl deleted file mode 100644 index a6a7dcfc6be2e..0000000000000 --- a/libclc/opencl/lib/r600/synchronization/barrier.cl +++ /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 -// -//===----------------------------------------------------------------------===// - -#include <clc/opencl/opencl-base.h> - -_CLC_DEF void __clc_r600_barrier(void) __asm("llvm.r600.group.barrier"); - -_CLC_DEF _CLC_OVERLOAD void barrier(uint flags) { - // We should call mem_fence here, but that is not implemented for r600 yet - __clc_r600_barrier(); -} diff --git a/libclc/opencl/lib/r600/workitem/get_global_offset.cl b/libclc/opencl/lib/r600/workitem/get_global_offset.cl deleted file mode 100644 index 477e8f405ec07..0000000000000 --- a/libclc/opencl/lib/r600/workitem/get_global_offset.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/opencl/opencl-base.h> - -_CLC_DEF _CLC_OVERLOAD uint get_global_offset(uint dim) { - __attribute__((address_space(7))) uint *ptr = - (__attribute__((address_space(7))) - uint *)__builtin_r600_implicitarg_ptr(); - if (dim < 3) - return ptr[dim + 1]; - return 0; -} diff --git a/libclc/opencl/lib/r600/workitem/get_global_size.cl b/libclc/opencl/lib/r600/workitem/get_global_size.cl deleted file mode 100644 index 3ea3881d2f3c6..0000000000000 --- a/libclc/opencl/lib/r600/workitem/get_global_size.cl +++ /dev/null @@ -1,26 +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> - -uint __clc_r600_get_global_size_x(void) __asm("llvm.r600.read.global.size.x"); -uint __clc_r600_get_global_size_y(void) __asm("llvm.r600.read.global.size.y"); -uint __clc_r600_get_global_size_z(void) __asm("llvm.r600.read.global.size.z"); - -_CLC_DEF _CLC_OVERLOAD size_t get_global_size(uint dim) { - switch (dim) { - case 0: - return __clc_r600_get_global_size_x(); - case 1: - return __clc_r600_get_global_size_y(); - case 2: - return __clc_r600_get_global_size_z(); - default: - return 1; - } -} diff --git a/libclc/opencl/lib/r600/workitem/get_group_id.cl b/libclc/opencl/lib/r600/workitem/get_group_id.cl deleted file mode 100644 index bf426cc2bb4dd..0000000000000 --- a/libclc/opencl/lib/r600/workitem/get_group_id.cl +++ /dev/null @@ -1,22 +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_DEF _CLC_OVERLOAD uint get_group_id(uint dim) { - switch (dim) { - case 0: - return __builtin_r600_read_tgid_x(); - case 1: - return __builtin_r600_read_tgid_y(); - case 2: - return __builtin_r600_read_tgid_z(); - default: - return 1; - } -} diff --git a/libclc/opencl/lib/r600/workitem/get_local_id.cl b/libclc/opencl/lib/r600/workitem/get_local_id.cl deleted file mode 100644 index 4915e07e7f181..0000000000000 --- a/libclc/opencl/lib/r600/workitem/get_local_id.cl +++ /dev/null @@ -1,22 +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_DEF _CLC_OVERLOAD uint get_local_id(uint dim) { - switch (dim) { - case 0: - return __builtin_r600_read_tidig_x(); - case 1: - return __builtin_r600_read_tidig_y(); - case 2: - return __builtin_r600_read_tidig_z(); - default: - return 1; - } -} diff --git a/libclc/opencl/lib/r600/workitem/get_local_size.cl b/libclc/opencl/lib/r600/workitem/get_local_size.cl deleted file mode 100644 index 877d9c359342c..0000000000000 --- a/libclc/opencl/lib/r600/workitem/get_local_size.cl +++ /dev/null @@ -1,26 +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> - -uint __clc_r600_get_local_size_x(void) __asm("llvm.r600.read.local.size.x"); -uint __clc_r600_get_local_size_y(void) __asm("llvm.r600.read.local.size.y"); -uint __clc_r600_get_local_size_z(void) __asm("llvm.r600.read.local.size.z"); - -_CLC_DEF _CLC_OVERLOAD size_t get_local_size(uint dim) { - switch (dim) { - case 0: - return __clc_r600_get_local_size_x(); - case 1: - return __clc_r600_get_local_size_y(); - case 2: - return __clc_r600_get_local_size_z(); - default: - return 1; - } -} diff --git a/libclc/opencl/lib/r600/workitem/get_num_groups.cl b/libclc/opencl/lib/r600/workitem/get_num_groups.cl deleted file mode 100644 index d04a69f9e0aa4..0000000000000 --- a/libclc/opencl/lib/r600/workitem/get_num_groups.cl +++ /dev/null @@ -1,26 +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> - -uint __clc_r600_get_num_groups_x(void) __asm("llvm.r600.read.ngroups.x"); -uint __clc_r600_get_num_groups_y(void) __asm("llvm.r600.read.ngroups.y"); -uint __clc_r600_get_num_groups_z(void) __asm("llvm.r600.read.ngroups.z"); - -_CLC_DEF _CLC_OVERLOAD size_t get_num_groups(uint dim) { - switch (dim) { - case 0: - return __clc_r600_get_num_groups_x(); - case 1: - return __clc_r600_get_num_groups_y(); - case 2: - return __clc_r600_get_num_groups_z(); - default: - return 1; - } -} diff --git a/libclc/opencl/lib/r600/workitem/get_work_dim.cl b/libclc/opencl/lib/r600/workitem/get_work_dim.cl deleted file mode 100644 index 1b743ae925a37..0000000000000 --- a/libclc/opencl/lib/r600/workitem/get_work_dim.cl +++ /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 -// -//===----------------------------------------------------------------------===// - -#include <clc/opencl/opencl-base.h> - -_CLC_DEF _CLC_OVERLOAD uint get_work_dim(void) { - __attribute__((address_space(7))) uint *ptr = - (__attribute__((address_space(7))) - uint *)__builtin_r600_implicitarg_ptr(); - return ptr[0]; -} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
