https://github.com/arsenm created https://github.com/llvm/llvm-project/pull/181701
These versions are old, and it hasn't made any sense to check against the compiler version since the monorepo. >From d065b73c5384443eab78198540fe7aba9a24b86a Mon Sep 17 00:00:00 2001 From: Matt Arsenault <[email protected]> Date: Mon, 16 Feb 2026 16:04:58 +0100 Subject: [PATCH] libclc: Remove old clang version checks These versions are old, and it hasn't made any sense to check against the compiler version since the monorepo. --- libclc/CMakeLists.txt | 6 ------ .../amdgcn/workitem/clc_get_global_offset.cl | 10 +--------- .../lib/amdgcn/workitem/clc_get_work_dim.cl | 10 +--------- .../amdgcn-amdhsa/workitem/get_global_size.cl | 18 +----------------- .../amdgcn-amdhsa/workitem/get_local_size.cl | 18 +----------------- 5 files changed, 4 insertions(+), 58 deletions(-) diff --git a/libclc/CMakeLists.txt b/libclc/CMakeLists.txt index e12a971452132..6430c4bf9c2be 100644 --- a/libclc/CMakeLists.txt +++ b/libclc/CMakeLists.txt @@ -36,8 +36,6 @@ set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS clc/lib/spirv/SOURCES; ) -set( LIBCLC_MIN_LLVM 3.9.0 ) - # A runtimes cross-build should only use the requested target. set( LIBCLC_DEFAULT_TARGET "all" ) if( LLVM_RUNTIMES_BUILD AND LLVM_DEFAULT_TARGET_TRIPLE MATCHES "^nvptx|^amdgcn" ) @@ -67,10 +65,6 @@ if( LIBCLC_STANDALONE_BUILD OR CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DI message( STATUS "libclc LLVM version: ${LLVM_PACKAGE_VERSION}" ) - if( LLVM_PACKAGE_VERSION VERSION_LESS LIBCLC_MIN_LLVM ) - message( FATAL_ERROR "libclc needs at least LLVM ${LIBCLC_MIN_LLVM}" ) - endif() - # Import required tools if( NOT EXISTS ${LIBCLC_CUSTOM_LLVM_TOOLS_BINARY_DIR} ) foreach( tool IN ITEMS clang llvm-as llvm-link opt ) diff --git a/libclc/clc/lib/amdgcn/workitem/clc_get_global_offset.cl b/libclc/clc/lib/amdgcn/workitem/clc_get_global_offset.cl index 24b04fb1679d8..cd7839a43dd69 100644 --- a/libclc/clc/lib/amdgcn/workitem/clc_get_global_offset.cl +++ b/libclc/clc/lib/amdgcn/workitem/clc_get_global_offset.cl @@ -8,16 +8,8 @@ #include <clc/workitem/clc_get_global_offset.h> -#if __clang_major__ >= 8 -#define CONST_AS __constant -#elif __clang_major__ >= 7 -#define CONST_AS __attribute__((address_space(4))) -#else -#define CONST_AS __attribute__((address_space(2))) -#endif - _CLC_DEF _CLC_OVERLOAD size_t __clc_get_global_offset(uint dim) { - CONST_AS uint *ptr = (CONST_AS uint *)__builtin_amdgcn_implicitarg_ptr(); + __constant uint *ptr = (__constant uint *)__builtin_amdgcn_implicitarg_ptr(); if (dim < 3) return ptr[dim + 1]; return 0; diff --git a/libclc/clc/lib/amdgcn/workitem/clc_get_work_dim.cl b/libclc/clc/lib/amdgcn/workitem/clc_get_work_dim.cl index 93cab09b4109f..0f079b948f9cb 100644 --- a/libclc/clc/lib/amdgcn/workitem/clc_get_work_dim.cl +++ b/libclc/clc/lib/amdgcn/workitem/clc_get_work_dim.cl @@ -8,15 +8,7 @@ #include <clc/workitem/clc_get_work_dim.h> -#if __clang_major__ >= 8 -#define CONST_AS __constant -#elif __clang_major__ >= 7 -#define CONST_AS __attribute__((address_space(4))) -#else -#define CONST_AS __attribute__((address_space(2))) -#endif - _CLC_OVERLOAD _CLC_DEF uint __clc_get_work_dim() { - CONST_AS uint *ptr = (CONST_AS uint *)__builtin_amdgcn_implicitarg_ptr(); + __constant uint *ptr = (__constant uint *)__builtin_amdgcn_implicitarg_ptr(); return ptr[0]; } diff --git a/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_global_size.cl b/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_global_size.cl index b450530c7dde5..c10cdd2d02efb 100644 --- a/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_global_size.cl +++ b/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_global_size.cl @@ -8,24 +8,8 @@ #include <clc/opencl/opencl-base.h> -#if __clang_major__ >= 8 -#define CONST_AS __constant -#elif __clang_major__ >= 7 -#define CONST_AS __attribute__((address_space(4))) -#else -#define CONST_AS __attribute__((address_space(2))) -#endif - -#if __clang_major__ >= 6 -#define __dispatch_ptr __builtin_amdgcn_dispatch_ptr -#else -#define __dispatch_ptr __clc_amdgcn_dispatch_ptr -CONST_AS uchar * -__clc_amdgcn_dispatch_ptr(void) __asm("llvm.amdgcn.dispatch.ptr"); -#endif - _CLC_DEF _CLC_OVERLOAD size_t get_global_size(uint dim) { - CONST_AS uint *ptr = (CONST_AS uint *)__dispatch_ptr(); + __constant uint *ptr = (__constant uint *)__builtin_amdgcn_dispatch_ptr(); if (dim < 3) return ptr[3 + dim]; return 1; diff --git a/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_local_size.cl b/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_local_size.cl index 6f92d9759a220..a95c58ca18534 100644 --- a/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_local_size.cl +++ b/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_local_size.cl @@ -8,24 +8,8 @@ #include <clc/opencl/opencl-base.h> -#if __clang_major__ >= 8 -#define CONST_AS __constant -#elif __clang_major__ >= 7 -#define CONST_AS __attribute__((address_space(4))) -#else -#define CONST_AS __attribute__((address_space(2))) -#endif - -#if __clang_major__ >= 6 -#define __dispatch_ptr __builtin_amdgcn_dispatch_ptr -#else -#define __dispatch_ptr __clc_amdgcn_dispatch_ptr -CONST_AS char * -__clc_amdgcn_dispatch_ptr(void) __asm("llvm.amdgcn.dispatch.ptr"); -#endif - _CLC_DEF _CLC_OVERLOAD size_t get_local_size(uint dim) { - CONST_AS uint *ptr = (CONST_AS uint *)__dispatch_ptr(); + __constant uint *ptr = (__constant uint *)__builtin_amdgcn_dispatch_ptr(); switch (dim) { case 0: return ptr[1] & 0xffffu; _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
