shylie/llvm-project
1
0
Fork 0
Code Issues Pull Requests Actions 6 Packages Projects Releases Wiki Activity

libclc: Remove old clang version checks (#181701)

Browse Source
This commit is contained in:
Matt Arsenault 2026-02-17 00:25:52 +01:00 committed by GitHub
parent 4a4ab28f94
commit ac545c7ea8
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
5 changed files with 4 additions and 58 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
libclc/CMakeLists.txt
View File

@ -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 )

10
libclc/clc/lib/amdgcn/workitem/clc_get_global_offset.cl
View File

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

10
libclc/clc/lib/amdgcn/workitem/clc_get_work_dim.cl
View File

@ -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];
}

18
libclc/opencl/lib/amdgcn-amdhsa/workitem/get_global_size.cl
View File

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

18
libclc/opencl/lib/amdgcn-amdhsa/workitem/get_local_size.cl
View File

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