[libclc] ac545c7 - libclc: Remove old clang version checks (#181701)

[via cfe-commits]

Author: Matt Arsenault
Date: 2026-02-17T00:25:52+01:00
New Revision: ac545c7ea8c1afce005e7bdaa34e86d899509d9c

URL: https://github.com/llvm/llvm-project/commit/ac545c7ea8c1afce005e7bdaa34e86d899509d9c
DIFF: https://github.com/llvm/llvm-project/commit/ac545c7ea8c1afce005e7bdaa34e86d899509d9c.diff

LOG: libclc: Remove old clang version checks (#181701)

Added: 
    

Modified: 
    libclc/CMakeLists.txt
    libclc/clc/lib/amdgcn/workitem/clc_get_global_offset.cl
    libclc/clc/lib/amdgcn/workitem/clc_get_work_dim.cl
    libclc/opencl/lib/amdgcn-amdhsa/workitem/get_global_size.cl
    libclc/opencl/lib/amdgcn-amdhsa/workitem/get_local_size.cl

Removed: 
    


################################################################################
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;