Script 'mail_helper' called by obssrc Hello community, here is the log from the commit of package libclc for openSUSE:Factory checked in at 2021-04-18 21:45:18 ++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ Comparing /work/SRC/openSUSE:Factory/libclc (Old) and /work/SRC/openSUSE:Factory/.libclc.new.12324 (New) ++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++
Package is "libclc" Sun Apr 18 21:45:18 2021 rev:15 rq:886350 version:0.2.0+llvm12.0.0 Changes: -------- --- /work/SRC/openSUSE:Factory/libclc/libclc.changes 2021-01-15 19:43:45.537793032 +0100 +++ /work/SRC/openSUSE:Factory/.libclc.new.12324/libclc.changes 2021-04-18 21:45:28.244758402 +0200 @@ -1,0 +2,5 @@ +Wed Apr 14 23:57:28 UTC 2021 - Aaron Puchert <[email protected]> + +- Update to version 0.2.0+llvm12.0.0. + +------------------------------------------------------------------- Old: ---- libclc-11.0.1.src.tar.xz New: ---- libclc-12.0.0.src.tar.xz ++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ Other differences: ------------------ ++++++ libclc.spec ++++++ --- /var/tmp/diff_new_pack.lrwwRN/_old 2021-04-18 21:45:28.760759275 +0200 +++ /var/tmp/diff_new_pack.lrwwRN/_new 2021-04-18 21:45:28.764759281 +0200 @@ -16,13 +16,13 @@ # -%define _libclc_llvm_ver 11.0.1 +%define _libclc_llvm_ver 12.0.0 Name: libclc Version: 0.2.0+llvm%{_libclc_llvm_ver} Release: 0 Summary: OpenCL C programming language library -License: (BSD-3-Clause OR MIT) AND Apache-2.0 WITH LLVM-exception +License: Apache-2.0 WITH LLVM-exception AND (BSD-3-Clause OR MIT) Group: Development/Libraries/C and C++ URL: https://libclc.llvm.org/ Source0: https://github.com/llvm/llvm-project/releases/download/llvmorg-%{_libclc_llvm_ver}/%{name}-%{_libclc_llvm_ver}.src.tar.xz @@ -44,11 +44,12 @@ # The libraries are bitcode files, so LTO is neither supported nor does it help. %define _lto_cflags %{nil} -sed -i "s|python|python3|g" CMakeLists.txt +# TODO: For building all targets, we need llvm-spirv. %cmake \ -DCMAKE_C_COMPILER=clang \ -DCMAKE_CXX_COMPILER=clang++ \ - -DENABLE_RUNTIME_SUBNORMAL:BOOL=ON + -DENABLE_RUNTIME_SUBNORMAL:BOOL=ON \ + -DLIBCLC_TARGETS_TO_BUILD="amdgcn--;amdgcn--amdhsa;amdgcn-mesa-mesa3d;r600--;nvptx--;nvptx64--;nvptx--nvidiacl;nvptx64--nvidiacl" %cmake_build %install ++++++ libclc-11.0.1.src.tar.xz -> libclc-12.0.0.src.tar.xz ++++++ diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/CMakeLists.txt new/libclc-12.0.0.src/CMakeLists.txt --- old/libclc-11.0.1.src/CMakeLists.txt 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/CMakeLists.txt 2021-04-06 18:38:18.000000000 +0200 @@ -1,4 +1,4 @@ -cmake_minimum_required( VERSION 3.9.2 ) +cmake_minimum_required(VERSION 3.13.4) project( libclc VERSION 0.2.0 LANGUAGES CXX ) include( GNUInstallDirs ) @@ -10,7 +10,9 @@ generic/lib/SOURCES; ptx/lib/SOURCES; ptx-nvidiacl/lib/SOURCES; - r600/lib/SOURCES + r600/lib/SOURCES; + spirv/lib/SOURCES; + spirv64/lib/SOURCES ) # List of all targets @@ -22,6 +24,8 @@ nvptx64-- nvptx--nvidiacl nvptx64--nvidiacl + spirv-mesa3d- + spirv64-mesa3d- ) set( LIBCLC_MIN_LLVM "3.9.0" ) @@ -53,8 +57,6 @@ set( LIBCLC_TARGETS_TO_BUILD ${LIBCLC_TARGETS_ALL} ) endif() -list( SORT LIBCLC_TARGETS_TO_BUILD ) - execute_process( COMMAND ${LLVM_CONFIG} "--system-libs" OUTPUT_VARIABLE LLVM_SYSTEM_LIBS OUTPUT_STRIP_TRAILING_WHITESPACE ) @@ -93,17 +95,27 @@ find_program( LLVM_AS llvm-as PATHS ${LLVM_BINDIR} NO_DEFAULT_PATH ) find_program( LLVM_LINK llvm-link PATHS ${LLVM_BINDIR} NO_DEFAULT_PATH ) find_program( LLVM_OPT opt PATHS ${LLVM_BINDIR} NO_DEFAULT_PATH ) +find_program( LLVM_SPIRV llvm-spirv PATHS ${LLVM_BINDIR} NO_DEFAULT_PATH ) # Print toolchain message( "clang: ${LLVM_CLANG}" ) message( "llvm-as: ${LLVM_AS}" ) message( "llvm-link: ${LLVM_LINK}" ) message( "opt: ${LLVM_OPT}" ) +message( "llvm-spirv: ${LLVM_SPIRV}" ) message( "" ) if( NOT LLVM_CLANG OR NOT LLVM_OPT OR NOT LLVM_AS OR NOT LLVM_LINK ) message( FATAL_ERROR "toolchain incomplete!" ) endif() +list( SORT LIBCLC_TARGETS_TO_BUILD ) + +if( "spirv-mesa3d-" IN_LIST LIBCLC_TARGETS_TO_BUILD OR "spirv64-mesa3d-" IN_LIST LIBCLC_TARGETS_TO_BUILD ) + if( NOT LLVM_SPIRV ) + message( FATAL_ERROR "SPIR-V targets requested, but spirv-tools is not installed" ) + endif() +endif() + set( CMAKE_MODULE_PATH ${CMAKE_SOURCE_DIR}/cmake ) set( CMAKE_CLC_COMPILER ${LLVM_CLANG} ) set( CMAKE_CLC_ARCHIVE ${LLVM_LINK} ) @@ -137,6 +149,8 @@ set( nvptx64--_devices none ) set( nvptx--nvidiacl_devices none ) set( nvptx64--nvidiacl_devices none ) +set( spirv-mesa3d-_devices none ) +set( spirv64-mesa3d-_devices none ) # Setup aliases set( cedar_aliases palm sumo sumo2 redwood juniper ) @@ -170,11 +184,11 @@ DESTINATION ${CMAKE_INSTALL_DATADIR}/clc ) endif() -find_program( PYTHON python ) +find_package( Python3 REQUIRED COMPONENTS Interpreter ) file( TO_CMAKE_PATH ${CMAKE_SOURCE_DIR}/generic/lib/gen_convert.py script_loc ) add_custom_command( OUTPUT convert.cl - COMMAND ${PYTHON} ${script_loc} > convert.cl + COMMAND ${Python3_EXECUTABLE} ${script_loc} > convert.cl DEPENDS ${script_loc} ) add_custom_target( "generate_convert.cl" DEPENDS convert.cl ) @@ -187,9 +201,14 @@ list( GET TRIPLE 1 VENDOR ) list( GET TRIPLE 2 OS ) - set( dirs generic ) + set( dirs ) + + if ( NOT ${ARCH} STREQUAL spirv AND NOT ${ARCH} STREQUAL spirv64 ) + LIST( APPEND dirs generic ) + endif() + if( ${ARCH} STREQUAL r600 OR ${ARCH} STREQUAL amdgcn ) - set( dirs ${dirs} amdgpu ) + list( APPEND dirs amdgpu ) endif() #nvptx is special @@ -215,10 +234,15 @@ # Add the generated convert.cl here to prevent adding # the one listed in SOURCES - set( rel_files convert.cl ) - set( objects convert.cl ) - if( NOT ENABLE_RUNTIME_SUBNORMAL ) - list( APPEND rel_files generic/lib/subnormal_use_default.ll ) + if( NOT ${ARCH} STREQUAL "spirv" AND NOT ${ARCH} STREQUAL "spirv64" ) + set( rel_files convert.cl ) + set( objects convert.cl ) + if( NOT ENABLE_RUNTIME_SUBNORMAL ) + list( APPEND rel_files generic/lib/subnormal_use_default.ll ) + endif() + else() + set( rel_files ) + set( objects ) endif() foreach( l ${source_list} ) @@ -242,7 +266,7 @@ foreach( d ${${t}_devices} ) # Some targets don't have a specific GPU to target - if( ${d} STREQUAL "none" ) + if( ${d} STREQUAL "none" OR ${ARCH} STREQUAL "spirv" OR ${ARCH} STREQUAL "spirv64" ) set( mcpu ) set( arch_suffix "${t}" ) else() @@ -251,6 +275,20 @@ endif() message( " DEVICE: ${d} ( ${${d}_aliases} )" ) + if ( ${ARCH} STREQUAL "spirv" OR ${ARCH} STREQUAL "spirv64" ) + if( ${ARCH} STREQUAL "spirv" ) + set( t "spir--" ) + else() + set( t "spir64--" ) + endif() + set( build_flags -O0 -finline-hint-functions ) + set( opt_flags ) + set( spvflags --spirv-max-version=1.1 ) + else() + set( build_flags ) + set( opt_flags -O3 ) + endif() + add_library( builtins.link.${arch_suffix} STATIC ${rel_files} ) # Make sure we depend on the pseudo target to prevent # multiple invocations @@ -261,8 +299,11 @@ "generic/include" ) target_compile_definitions( builtins.link.${arch_suffix} PRIVATE "__CLC_INTERNAL" ) + string( TOUPPER "-DCLC_${ARCH}" CLC_TARGET_DEFINE ) + target_compile_definitions( builtins.link.${arch_suffix} PRIVATE + ${CLC_TARGET_DEFINE} ) target_compile_options( builtins.link.${arch_suffix} PRIVATE -target - ${t} ${mcpu} -fno-builtin -nostdlib ) + ${t} ${mcpu} -fno-builtin -nostdlib ${build_flags} ) set_target_properties( builtins.link.${arch_suffix} PROPERTIES LINKER_LANGUAGE CLC ) @@ -270,42 +311,56 @@ # Add opt target add_custom_command( OUTPUT "builtins.opt.${obj_suffix}" - COMMAND ${LLVM_OPT} -O3 -o + COMMAND ${LLVM_OPT} ${opt_flags} -o "builtins.opt.${obj_suffix}" "builtins.link.${obj_suffix}" DEPENDS "builtins.link.${arch_suffix}" ) add_custom_target( "opt.${obj_suffix}" ALL DEPENDS "builtins.opt.${obj_suffix}" ) - # Add prepare target - add_custom_command( OUTPUT "${obj_suffix}" - COMMAND prepare_builtins -o - "${obj_suffix}" - "builtins.opt.${obj_suffix}" - DEPENDS "opt.${obj_suffix}" - "builtins.opt.${obj_suffix}" - prepare_builtins ) - add_custom_target( "prepare-${obj_suffix}" ALL - DEPENDS "${obj_suffix}" ) - install( FILES ${CMAKE_CURRENT_BINARY_DIR}/${obj_suffix} DESTINATION ${CMAKE_INSTALL_DATADIR}/clc ) - # nvptx-- targets don't include workitem builtins - if( NOT ${t} MATCHES ".*ptx.*--$" ) - add_test( NAME external-calls-${obj_suffix} - COMMAND ./check_external_calls.sh ${CMAKE_CURRENT_BINARY_DIR}/${obj_suffix} - WORKING_DIRECTORY ${CMAKE_SOURCE_DIR} ) - set_tests_properties( external-calls-${obj_suffix} - PROPERTIES ENVIRONMENT "LLVM_CONFIG=${LLVM_CONFIG}" ) - endif() + if( ${ARCH} STREQUAL "spirv" OR ${ARCH} STREQUAL "spirv64" ) + set( spv_suffix ${arch_suffix}.spv ) + add_custom_command( OUTPUT "${spv_suffix}" + COMMAND ${LLVM_SPIRV} ${spvflags} + -o "${spv_suffix}" + "builtins.link.${obj_suffix}" + DEPENDS "builtins.link.${arch_suffix}" ) + add_custom_target( "prepare-${spv_suffix}" ALL + DEPENDS "${spv_suffix}" ) + install( FILES ${CMAKE_CURRENT_BINARY_DIR}/${spv_suffix} + DESTINATION ${CMAKE_INSTALL_DATADIR}/clc ) + else() + # Add prepare target + add_custom_command( OUTPUT "${obj_suffix}" + COMMAND prepare_builtins -o + "${obj_suffix}" + "builtins.opt.${obj_suffix}" + DEPENDS "opt.${obj_suffix}" + "builtins.opt.${obj_suffix}" + prepare_builtins ) + add_custom_target( "prepare-${obj_suffix}" ALL + DEPENDS "${obj_suffix}" ) + + # nvptx-- targets don't include workitem builtins + if( NOT ${t} MATCHES ".*ptx.*--$" ) + add_test( NAME external-calls-${obj_suffix} + COMMAND ./check_external_calls.sh ${CMAKE_CURRENT_BINARY_DIR}/${obj_suffix} + WORKING_DIRECTORY ${CMAKE_SOURCE_DIR} ) + set_tests_properties( external-calls-${obj_suffix} + PROPERTIES ENVIRONMENT "LLVM_CONFIG=${LLVM_CONFIG}" ) + endif() - foreach( a ${${d}_aliases} ) - set( alias_suffix "${a}-${t}.bc" ) - add_custom_target( ${alias_suffix} ALL - COMMAND ${CMAKE_COMMAND} -E - create_symlink ${obj_suffix} - ${alias_suffix} - DEPENDS "prepare-${obj_suffix}" ) - install( FILES ${CMAKE_CURRENT_BINARY_DIR}/${alias_suffix} DESTINATION ${CMAKE_INSTALL_DATADIR}/clc ) - endforeach( a ) + install( FILES ${CMAKE_CURRENT_BINARY_DIR}/${obj_suffix} DESTINATION ${CMAKE_INSTALL_DATADIR}/clc ) + foreach( a ${${d}_aliases} ) + set( alias_suffix "${a}-${t}.bc" ) + add_custom_target( ${alias_suffix} ALL + COMMAND ${CMAKE_COMMAND} -E + create_symlink ${obj_suffix} + ${alias_suffix} + DEPENDS "prepare-${obj_suffix}" ) + install( FILES ${CMAKE_CURRENT_BINARY_DIR}/${alias_suffix} DESTINATION ${CMAKE_INSTALL_DATADIR}/clc ) + endforeach( a ) + endif() endforeach( d ) endforeach( t ) diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/amdgcn/lib/mem_fence/fence.cl new/libclc-12.0.0.src/amdgcn/lib/mem_fence/fence.cl --- old/libclc-11.0.1.src/amdgcn/lib/mem_fence/fence.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/amdgcn/lib/mem_fence/fence.cl 2021-04-06 18:38:18.000000000 +0200 @@ -17,24 +17,21 @@ _CLC_DEF void __clc_amdgcn_s_waitcnt(unsigned) __asm("llvm.amdgcn.s.waitcnt"); #endif -_CLC_DEF void mem_fence(cl_mem_fence_flags flags) -{ - if (flags & CLK_GLOBAL_MEM_FENCE) { - // scalar loads are counted with LGKM but we don't know whether - // the compiler turned any loads to scalar - __waitcnt(0); - } else if (flags & CLK_LOCAL_MEM_FENCE) - __waitcnt(0xff); // LGKM is [12:8] +_CLC_DEF _CLC_OVERLOAD void mem_fence(cl_mem_fence_flags flags) { + if (flags & CLK_GLOBAL_MEM_FENCE) { + // scalar loads are counted with LGKM but we don't know whether + // the compiler turned any loads to scalar + __waitcnt(0); + } else if (flags & CLK_LOCAL_MEM_FENCE) + __waitcnt(0xff); // LGKM is [12:8] } #undef __waitcnt // We don't have separate mechanism for read and write fences -_CLC_DEF void read_mem_fence(cl_mem_fence_flags flags) -{ - mem_fence(flags); +_CLC_DEF _CLC_OVERLOAD void read_mem_fence(cl_mem_fence_flags flags) { + mem_fence(flags); } -_CLC_DEF void write_mem_fence(cl_mem_fence_flags flags) -{ - mem_fence(flags); +_CLC_DEF _CLC_OVERLOAD void write_mem_fence(cl_mem_fence_flags flags) { + mem_fence(flags); } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/amdgcn/lib/synchronization/barrier.cl new/libclc-12.0.0.src/amdgcn/lib/synchronization/barrier.cl --- old/libclc-11.0.1.src/amdgcn/lib/synchronization/barrier.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/amdgcn/lib/synchronization/barrier.cl 2021-04-06 18:38:18.000000000 +0200 @@ -1,7 +1,6 @@ #include <clc/clc.h> -_CLC_DEF void barrier(cl_mem_fence_flags flags) -{ - mem_fence(flags); - __builtin_amdgcn_s_barrier(); +_CLC_DEF _CLC_OVERLOAD void barrier(cl_mem_fence_flags flags) { + mem_fence(flags); + __builtin_amdgcn_s_barrier(); } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/amdgcn/lib/workitem/get_global_offset.cl new/libclc-12.0.0.src/amdgcn/lib/workitem/get_global_offset.cl --- old/libclc-11.0.1.src/amdgcn/lib/workitem/get_global_offset.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/amdgcn/lib/workitem/get_global_offset.cl 2021-04-06 18:38:18.000000000 +0200 @@ -8,11 +8,9 @@ #define CONST_AS __attribute__((address_space(2))) #endif -_CLC_DEF size_t get_global_offset(uint dim) -{ - CONST_AS uint * ptr = - (CONST_AS uint *) __builtin_amdgcn_implicitarg_ptr(); - if (dim < 3) - return ptr[dim + 1]; - return 0; +_CLC_DEF _CLC_OVERLOAD size_t get_global_offset(uint dim) { + CONST_AS uint *ptr = (CONST_AS uint *)__builtin_amdgcn_implicitarg_ptr(); + if (dim < 3) + return ptr[dim + 1]; + return 0; } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/amdgcn/lib/workitem/get_global_size.cl new/libclc-12.0.0.src/amdgcn/lib/workitem/get_global_size.cl --- old/libclc-11.0.1.src/amdgcn/lib/workitem/get_global_size.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/amdgcn/lib/workitem/get_global_size.cl 2021-04-06 18:38:18.000000000 +0200 @@ -4,12 +4,15 @@ uint __clc_amdgcn_get_global_size_y(void) __asm("llvm.r600.read.global.size.y"); uint __clc_amdgcn_get_global_size_z(void) __asm("llvm.r600.read.global.size.z"); -_CLC_DEF size_t get_global_size(uint dim) -{ - switch (dim) { - case 0: return __clc_amdgcn_get_global_size_x(); - case 1: return __clc_amdgcn_get_global_size_y(); - case 2: return __clc_amdgcn_get_global_size_z(); - default: return 1; - } +_CLC_DEF _CLC_OVERLOAD size_t get_global_size(uint dim) { + switch (dim) { + case 0: + return __clc_amdgcn_get_global_size_x(); + case 1: + return __clc_amdgcn_get_global_size_y(); + case 2: + return __clc_amdgcn_get_global_size_z(); + default: + return 1; + } } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/amdgcn/lib/workitem/get_group_id.cl new/libclc-12.0.0.src/amdgcn/lib/workitem/get_group_id.cl --- old/libclc-11.0.1.src/amdgcn/lib/workitem/get_group_id.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/amdgcn/lib/workitem/get_group_id.cl 2021-04-06 18:38:18.000000000 +0200 @@ -1,11 +1,14 @@ #include <clc/clc.h> -_CLC_DEF size_t get_group_id(uint dim) -{ - switch(dim) { - case 0: return __builtin_amdgcn_workgroup_id_x(); - case 1: return __builtin_amdgcn_workgroup_id_y(); - case 2: return __builtin_amdgcn_workgroup_id_z(); - default: return 1; - } +_CLC_DEF _CLC_OVERLOAD size_t get_group_id(uint dim) { + switch (dim) { + case 0: + return __builtin_amdgcn_workgroup_id_x(); + case 1: + return __builtin_amdgcn_workgroup_id_y(); + case 2: + return __builtin_amdgcn_workgroup_id_z(); + default: + return 1; + } } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/amdgcn/lib/workitem/get_local_id.cl new/libclc-12.0.0.src/amdgcn/lib/workitem/get_local_id.cl --- old/libclc-11.0.1.src/amdgcn/lib/workitem/get_local_id.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/amdgcn/lib/workitem/get_local_id.cl 2021-04-06 18:38:18.000000000 +0200 @@ -1,11 +1,14 @@ #include <clc/clc.h> -_CLC_DEF size_t get_local_id(uint dim) -{ - switch(dim) { - case 0: return __builtin_amdgcn_workitem_id_x(); - case 1: return __builtin_amdgcn_workitem_id_y(); - case 2: return __builtin_amdgcn_workitem_id_z(); - default: return 1; - } +_CLC_DEF _CLC_OVERLOAD size_t get_local_id(uint dim) { + switch (dim) { + case 0: + return __builtin_amdgcn_workitem_id_x(); + case 1: + return __builtin_amdgcn_workitem_id_y(); + case 2: + return __builtin_amdgcn_workitem_id_z(); + default: + return 1; + } } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/amdgcn/lib/workitem/get_local_size.cl new/libclc-12.0.0.src/amdgcn/lib/workitem/get_local_size.cl --- old/libclc-11.0.1.src/amdgcn/lib/workitem/get_local_size.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/amdgcn/lib/workitem/get_local_size.cl 2021-04-06 18:38:18.000000000 +0200 @@ -4,12 +4,15 @@ uint __clc_amdgcn_get_local_size_y(void) __asm("llvm.r600.read.local.size.y"); uint __clc_amdgcn_get_local_size_z(void) __asm("llvm.r600.read.local.size.z"); -_CLC_DEF size_t get_local_size(uint dim) -{ - switch (dim) { - case 0: return __clc_amdgcn_get_local_size_x(); - case 1: return __clc_amdgcn_get_local_size_y(); - case 2: return __clc_amdgcn_get_local_size_z(); - default: return 1; - } +_CLC_DEF _CLC_OVERLOAD size_t get_local_size(uint dim) { + switch (dim) { + case 0: + return __clc_amdgcn_get_local_size_x(); + case 1: + return __clc_amdgcn_get_local_size_y(); + case 2: + return __clc_amdgcn_get_local_size_z(); + default: + return 1; + } } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/amdgcn/lib/workitem/get_num_groups.cl new/libclc-12.0.0.src/amdgcn/lib/workitem/get_num_groups.cl --- old/libclc-11.0.1.src/amdgcn/lib/workitem/get_num_groups.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/amdgcn/lib/workitem/get_num_groups.cl 2021-04-06 18:38:18.000000000 +0200 @@ -4,12 +4,15 @@ uint __clc_amdgcn_get_num_groups_y(void) __asm("llvm.r600.read.ngroups.y"); uint __clc_amdgcn_get_num_groups_z(void) __asm("llvm.r600.read.ngroups.z"); -_CLC_DEF size_t get_num_groups(uint dim) -{ - switch (dim) { - case 0: return __clc_amdgcn_get_num_groups_x(); - case 1: return __clc_amdgcn_get_num_groups_y(); - case 2: return __clc_amdgcn_get_num_groups_z(); - default: return 1; - } +_CLC_DEF _CLC_OVERLOAD size_t get_num_groups(uint dim) { + switch (dim) { + case 0: + return __clc_amdgcn_get_num_groups_x(); + case 1: + return __clc_amdgcn_get_num_groups_y(); + case 2: + return __clc_amdgcn_get_num_groups_z(); + default: + return 1; + } } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/amdgcn/lib/workitem/get_work_dim.cl new/libclc-12.0.0.src/amdgcn/lib/workitem/get_work_dim.cl --- old/libclc-11.0.1.src/amdgcn/lib/workitem/get_work_dim.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/amdgcn/lib/workitem/get_work_dim.cl 2021-04-06 18:38:18.000000000 +0200 @@ -8,9 +8,7 @@ #define CONST_AS __attribute__((address_space(2))) #endif -_CLC_DEF uint get_work_dim(void) -{ - CONST_AS uint * ptr = - (CONST_AS uint *) __builtin_amdgcn_implicitarg_ptr(); - return ptr[0]; +_CLC_DEF _CLC_OVERLOAD uint get_work_dim(void) { + CONST_AS uint *ptr = (CONST_AS uint *)__builtin_amdgcn_implicitarg_ptr(); + return ptr[0]; } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/amdgcn-amdhsa/lib/workitem/get_global_size.cl new/libclc-12.0.0.src/amdgcn-amdhsa/lib/workitem/get_global_size.cl --- old/libclc-11.0.1.src/amdgcn-amdhsa/lib/workitem/get_global_size.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/amdgcn-amdhsa/lib/workitem/get_global_size.cl 2021-04-06 18:38:18.000000000 +0200 @@ -15,10 +15,9 @@ CONST_AS uchar * __clc_amdgcn_dispatch_ptr(void) __asm("llvm.amdgcn.dispatch.ptr"); #endif -_CLC_DEF size_t get_global_size(uint dim) -{ - CONST_AS uint * ptr = (CONST_AS uint *) __dispatch_ptr(); - if (dim < 3) - return ptr[3 + dim]; - return 1; +_CLC_DEF _CLC_OVERLOAD size_t get_global_size(uint dim) { + CONST_AS uint *ptr = (CONST_AS uint *)__dispatch_ptr(); + if (dim < 3) + return ptr[3 + dim]; + return 1; } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/amdgcn-amdhsa/lib/workitem/get_local_size.cl new/libclc-12.0.0.src/amdgcn-amdhsa/lib/workitem/get_local_size.cl --- old/libclc-11.0.1.src/amdgcn-amdhsa/lib/workitem/get_local_size.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/amdgcn-amdhsa/lib/workitem/get_local_size.cl 2021-04-06 18:38:18.000000000 +0200 @@ -15,16 +15,15 @@ CONST_AS char * __clc_amdgcn_dispatch_ptr(void) __asm("llvm.amdgcn.dispatch.ptr"); #endif -_CLC_DEF size_t get_local_size(uint dim) -{ - CONST_AS uint * ptr = (CONST_AS uint *) __dispatch_ptr(); - switch (dim) { - case 0: - return ptr[1] & 0xffffu; - case 1: - return ptr[1] >> 16; - case 2: - return ptr[2] & 0xffffu; - } - return 1; +_CLC_DEF _CLC_OVERLOAD size_t get_local_size(uint dim) { + CONST_AS uint *ptr = (CONST_AS uint *)__dispatch_ptr(); + switch (dim) { + case 0: + return ptr[1] & 0xffffu; + case 1: + return ptr[1] >> 16; + case 2: + return ptr[2] & 0xffffu; + } + return 1; } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/amdgcn-amdhsa/lib/workitem/get_num_groups.cl new/libclc-12.0.0.src/amdgcn-amdhsa/lib/workitem/get_num_groups.cl --- old/libclc-11.0.1.src/amdgcn-amdhsa/lib/workitem/get_num_groups.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/amdgcn-amdhsa/lib/workitem/get_num_groups.cl 2021-04-06 18:38:18.000000000 +0200 @@ -1,7 +1,7 @@ #include <clc/clc.h> -_CLC_DEF size_t get_num_groups(uint dim) { +_CLC_DEF _CLC_OVERLOAD size_t get_num_groups(uint dim) { size_t global_size = get_global_size(dim); size_t local_size = get_local_size(dim); size_t num_groups = global_size / local_size; diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/generic/include/clc/async/wait_group_events.h new/libclc-12.0.0.src/generic/include/clc/async/wait_group_events.h --- old/libclc-11.0.1.src/generic/include/clc/async/wait_group_events.h 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/generic/include/clc/async/wait_group_events.h 2021-04-06 18:38:18.000000000 +0200 @@ -1 +1,2 @@ -void wait_group_events(int num_events, event_t *event_list); +_CLC_DECL _CLC_OVERLOAD void wait_group_events(int num_events, + event_t *event_list); diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/generic/include/clc/clcfunc.h new/libclc-12.0.0.src/generic/include/clc/clcfunc.h --- old/libclc-11.0.1.src/generic/include/clc/clcfunc.h 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/generic/include/clc/clcfunc.h 2021-04-06 18:38:18.000000000 +0200 @@ -1,4 +1,10 @@ #define _CLC_OVERLOAD __attribute__((overloadable)) #define _CLC_DECL -#define _CLC_DEF __attribute__((always_inline)) #define _CLC_INLINE __attribute__((always_inline)) inline + +/* avoid inlines for SPIR-V since we'll optimise later in the chain */ +#if defined(CLC_SPIRV) || defined(CLC_SPIRV64) +#define _CLC_DEF +#else +#define _CLC_DEF __attribute__((always_inline)) +#endif diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/generic/include/clc/explicit_fence/explicit_memory_fence.h new/libclc-12.0.0.src/generic/include/clc/explicit_fence/explicit_memory_fence.h --- old/libclc-11.0.1.src/generic/include/clc/explicit_fence/explicit_memory_fence.h 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/generic/include/clc/explicit_fence/explicit_memory_fence.h 2021-04-06 18:38:18.000000000 +0200 @@ -1,3 +1,3 @@ -_CLC_DECL void mem_fence(cl_mem_fence_flags flags); -_CLC_DECL void read_mem_fence(cl_mem_fence_flags flags); -_CLC_DECL void write_mem_fence(cl_mem_fence_flags flags); +_CLC_DECL _CLC_OVERLOAD void mem_fence(cl_mem_fence_flags flags); +_CLC_DECL _CLC_OVERLOAD void read_mem_fence(cl_mem_fence_flags flags); +_CLC_DECL _CLC_OVERLOAD void write_mem_fence(cl_mem_fence_flags flags); diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/generic/include/clc/float/definitions.h new/libclc-12.0.0.src/generic/include/clc/float/definitions.h --- old/libclc-11.0.1.src/generic/include/clc/float/definitions.h 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/generic/include/clc/float/definitions.h 2021-04-06 18:38:18.000000000 +0200 @@ -15,7 +15,7 @@ #define FLT_EPSILON 0x1.0p-23f #define FP_ILOGB0 (-2147483647 - 1) -#define FP_ILOGBNAN (-2147483647 - 1) +#define FP_ILOGBNAN 2147483647 #define M_E_F 0x1.5bf0a8p+1f #define M_LOG2E_F 0x1.715476p+0f diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/generic/include/clc/synchronization/barrier.h new/libclc-12.0.0.src/generic/include/clc/synchronization/barrier.h --- old/libclc-11.0.1.src/generic/include/clc/synchronization/barrier.h 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/generic/include/clc/synchronization/barrier.h 2021-04-06 18:38:18.000000000 +0200 @@ -1 +1 @@ -_CLC_DECL void barrier(cl_mem_fence_flags flags); +_CLC_DECL _CLC_OVERLOAD void barrier(cl_mem_fence_flags flags); diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/generic/include/clc/workitem/get_global_id.h new/libclc-12.0.0.src/generic/include/clc/workitem/get_global_id.h --- old/libclc-11.0.1.src/generic/include/clc/workitem/get_global_id.h 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/generic/include/clc/workitem/get_global_id.h 2021-04-06 18:38:18.000000000 +0200 @@ -1 +1 @@ -_CLC_DECL size_t get_global_id(uint dim); +_CLC_DECL _CLC_OVERLOAD size_t get_global_id(uint dim); diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/generic/include/clc/workitem/get_global_offset.h new/libclc-12.0.0.src/generic/include/clc/workitem/get_global_offset.h --- old/libclc-11.0.1.src/generic/include/clc/workitem/get_global_offset.h 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/generic/include/clc/workitem/get_global_offset.h 2021-04-06 18:38:18.000000000 +0200 @@ -1 +1 @@ -_CLC_DECL size_t get_global_offset(uint dim); +_CLC_DECL _CLC_OVERLOAD size_t get_global_offset(uint dim); diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/generic/include/clc/workitem/get_global_size.h new/libclc-12.0.0.src/generic/include/clc/workitem/get_global_size.h --- old/libclc-11.0.1.src/generic/include/clc/workitem/get_global_size.h 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/generic/include/clc/workitem/get_global_size.h 2021-04-06 18:38:18.000000000 +0200 @@ -1 +1 @@ -_CLC_DECL size_t get_global_size(uint dim); +_CLC_DECL _CLC_OVERLOAD size_t get_global_size(uint dim); diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/generic/include/clc/workitem/get_group_id.h new/libclc-12.0.0.src/generic/include/clc/workitem/get_group_id.h --- old/libclc-11.0.1.src/generic/include/clc/workitem/get_group_id.h 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/generic/include/clc/workitem/get_group_id.h 2021-04-06 18:38:18.000000000 +0200 @@ -1 +1 @@ -_CLC_DECL size_t get_group_id(uint dim); +_CLC_DECL _CLC_OVERLOAD size_t get_group_id(uint dim); diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/generic/include/clc/workitem/get_local_id.h new/libclc-12.0.0.src/generic/include/clc/workitem/get_local_id.h --- old/libclc-11.0.1.src/generic/include/clc/workitem/get_local_id.h 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/generic/include/clc/workitem/get_local_id.h 2021-04-06 18:38:18.000000000 +0200 @@ -1 +1 @@ -_CLC_DECL size_t get_local_id(uint dim); +_CLC_DECL _CLC_OVERLOAD size_t get_local_id(uint dim); diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/generic/include/clc/workitem/get_local_size.h new/libclc-12.0.0.src/generic/include/clc/workitem/get_local_size.h --- old/libclc-11.0.1.src/generic/include/clc/workitem/get_local_size.h 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/generic/include/clc/workitem/get_local_size.h 2021-04-06 18:38:18.000000000 +0200 @@ -1 +1 @@ -_CLC_DECL size_t get_local_size(uint dim); +_CLC_DECL _CLC_OVERLOAD size_t get_local_size(uint dim); diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/generic/include/clc/workitem/get_num_groups.h new/libclc-12.0.0.src/generic/include/clc/workitem/get_num_groups.h --- old/libclc-11.0.1.src/generic/include/clc/workitem/get_num_groups.h 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/generic/include/clc/workitem/get_num_groups.h 2021-04-06 18:38:18.000000000 +0200 @@ -1 +1 @@ -_CLC_DECL size_t get_num_groups(uint dim); +_CLC_DECL _CLC_OVERLOAD size_t get_num_groups(uint dim); diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/generic/include/clc/workitem/get_work_dim.h new/libclc-12.0.0.src/generic/include/clc/workitem/get_work_dim.h --- old/libclc-11.0.1.src/generic/include/clc/workitem/get_work_dim.h 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/generic/include/clc/workitem/get_work_dim.h 2021-04-06 18:38:18.000000000 +0200 @@ -1 +1 @@ -_CLC_DECL uint get_work_dim(void); +_CLC_DECL _CLC_OVERLOAD uint get_work_dim(void); diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/generic/lib/async/wait_group_events.cl new/libclc-12.0.0.src/generic/lib/async/wait_group_events.cl --- old/libclc-11.0.1.src/generic/lib/async/wait_group_events.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/generic/lib/async/wait_group_events.cl 2021-04-06 18:38:18.000000000 +0200 @@ -1,5 +1,6 @@ #include <clc/clc.h> -_CLC_DEF void wait_group_events(int num_events, event_t *event_list) { +_CLC_DEF _CLC_OVERLOAD void wait_group_events(int num_events, + event_t *event_list) { barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE); } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/generic/lib/common/smoothstep.cl new/libclc-12.0.0.src/generic/lib/common/smoothstep.cl --- old/libclc-11.0.1.src/generic/lib/common/smoothstep.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/generic/lib/common/smoothstep.cl 2021-04-06 18:38:18.000000000 +0200 @@ -46,10 +46,12 @@ _CLC_TERNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, smoothstep, double, double, double); +#if !defined(CLC_SPIRV) && !defined(CLC_SPIRV64) SMOOTH_STEP_DEF(float, double, SMOOTH_STEP_IMPL_D); SMOOTH_STEP_DEF(double, float, SMOOTH_STEP_IMPL_D); _CLC_V_S_S_V_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, smoothstep, float, float, double); _CLC_V_S_S_V_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, smoothstep, double, double, float); +#endif #endif diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/generic/lib/common/step.cl new/libclc-12.0.0.src/generic/lib/common/step.cl --- old/libclc-11.0.1.src/generic/lib/common/step.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/generic/lib/common/step.cl 2021-04-06 18:38:18.000000000 +0200 @@ -45,10 +45,12 @@ _CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, step, double, double); _CLC_V_S_V_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, step, double, double); +#if !defined(CLC_SPIRV) && !defined(CLC_SPIRV64) STEP_DEF(float, double); STEP_DEF(double, float); _CLC_V_S_V_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, step, float, double); _CLC_V_S_V_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, step, double, float); +#endif #endif diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/generic/lib/gen_convert.py new/libclc-12.0.0.src/generic/lib/gen_convert.py --- old/libclc-11.0.1.src/generic/lib/gen_convert.py 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/generic/lib/gen_convert.py 2021-04-06 18:38:18.000000000 +0200 @@ -1,5 +1,3 @@ -#!/usr/bin/env python3 - # OpenCL built-in library: type conversion functions # # Copyright (c) 2013 Victor Oliveira <[email protected]> diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/generic/lib/math/ilogb.cl new/libclc-12.0.0.src/generic/lib/math/ilogb.cl --- old/libclc-11.0.1.src/generic/lib/math/ilogb.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/generic/lib/math/ilogb.cl 2021-04-06 18:38:18.000000000 +0200 @@ -31,7 +31,15 @@ int rs = -118 - (int) clz(ux & MANTBITS_SP32); int r = (int) (ax >> EXPSHIFTBITS_SP32) - EXPBIAS_SP32; r = ax < 0x00800000U ? rs : r; - r = ax > EXPBITS_SP32 | ax == 0 ? 0x80000000 : r; + r = ax == 0 ? FP_ILOGB0 : r; + + // We could merge those 2 tests and have: + // + // r = ax >= EXPBITS_SP32 ? 0x7fffffff : r + // + // since FP_ILOGBNAN is set to INT_MAX, but it's clearer this way and + // FP_ILOGBNAN can change without requiring changes to ilogb() code. + r = ax > EXPBITS_SP32 ? FP_ILOGBNAN : r; r = ax == EXPBITS_SP32 ? 0x7fffffff : r; return r; } @@ -47,7 +55,15 @@ int r = (int) (ax >> EXPSHIFTBITS_DP64) - EXPBIAS_DP64; int rs = -1011 - (int) clz(ax & MANTBITS_DP64); r = ax < 0x0010000000000000UL ? rs : r; - r = ax > 0x7ff0000000000000UL | ax == 0UL ? 0x80000000 : r; + r = ax == 0UL ? FP_ILOGB0 : r; + + // We could merge those 2 tests and have: + // + // r = ax >= 0x7ff0000000000000UL ? 0x7fffffff : r + // + // since FP_ILOGBNAN is set to INT_MAX, but it's clearer this way and + // FP_ILOGBNAN can change without requiring changes to ilogb() code. + r = ax > 0x7ff0000000000000UL ? FP_ILOGBNAN : r; r = ax == 0x7ff0000000000000UL ? 0x7fffffff : r; return r; } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/generic/lib/math/math.h new/libclc-12.0.0.src/generic/lib/math/math.h --- old/libclc-11.0.1.src/generic/lib/math/math.h 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/generic/lib/math/math.h 2021-04-06 18:38:18.000000000 +0200 @@ -40,6 +40,9 @@ #if (defined __AMDGCN__ || defined __R600__) && !defined __HAS_FMAF__ #define HAVE_HW_FMA32() (0) +#elif defined CLC_SPIRV || defined CLC_SPIRV64 +bool __attribute__((noinline)) __clc_runtime_has_hw_fma32(void); +#define HAVE_HW_FMA32() __clc_runtime_has_hw_fma32() #else #define HAVE_HW_FMA32() (1) #endif diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/generic/lib/workitem/get_global_id.cl new/libclc-12.0.0.src/generic/lib/workitem/get_global_id.cl --- old/libclc-11.0.1.src/generic/lib/workitem/get_global_id.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/generic/lib/workitem/get_global_id.cl 2021-04-06 18:38:18.000000000 +0200 @@ -1,5 +1,5 @@ #include <clc/clc.h> -_CLC_DEF size_t get_global_id(uint dim) { +_CLC_DEF _CLC_OVERLOAD size_t get_global_id(uint dim) { return get_group_id(dim) * get_local_size(dim) + get_local_id(dim) + get_global_offset(dim); } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/generic/lib/workitem/get_global_size.cl new/libclc-12.0.0.src/generic/lib/workitem/get_global_size.cl --- old/libclc-11.0.1.src/generic/lib/workitem/get_global_size.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/generic/lib/workitem/get_global_size.cl 2021-04-06 18:38:18.000000000 +0200 @@ -1,5 +1,5 @@ #include <clc/clc.h> -_CLC_DEF size_t get_global_size(uint dim) { +_CLC_DEF _CLC_OVERLOAD size_t get_global_size(uint dim) { return get_num_groups(dim)*get_local_size(dim); } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/ptx-nvidiacl/lib/mem_fence/fence.cl new/libclc-12.0.0.src/ptx-nvidiacl/lib/mem_fence/fence.cl --- old/libclc-11.0.1.src/ptx-nvidiacl/lib/mem_fence/fence.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/ptx-nvidiacl/lib/mem_fence/fence.cl 2021-04-06 18:38:18.000000000 +0200 @@ -1,15 +1,15 @@ #include <clc/clc.h> -_CLC_DEF void mem_fence(cl_mem_fence_flags flags) { - if (flags & (CLK_GLOBAL_MEM_FENCE | CLK_LOCAL_MEM_FENCE)) - __nvvm_membar_cta(); +_CLC_DEF _CLC_OVERLOAD void mem_fence(cl_mem_fence_flags flags) { + if (flags & (CLK_GLOBAL_MEM_FENCE | CLK_LOCAL_MEM_FENCE)) + __nvvm_membar_cta(); } // We do not have separate mechanism for read and write fences. -_CLC_DEF void read_mem_fence(cl_mem_fence_flags flags) { +_CLC_DEF _CLC_OVERLOAD void read_mem_fence(cl_mem_fence_flags flags) { mem_fence(flags); } -_CLC_DEF void write_mem_fence(cl_mem_fence_flags flags) { +_CLC_DEF _CLC_OVERLOAD void write_mem_fence(cl_mem_fence_flags flags) { mem_fence(flags); } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/ptx-nvidiacl/lib/synchronization/barrier.cl new/libclc-12.0.0.src/ptx-nvidiacl/lib/synchronization/barrier.cl --- old/libclc-11.0.1.src/ptx-nvidiacl/lib/synchronization/barrier.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/ptx-nvidiacl/lib/synchronization/barrier.cl 2021-04-06 18:38:18.000000000 +0200 @@ -1,6 +1,5 @@ #include <clc/clc.h> -_CLC_DEF void barrier(cl_mem_fence_flags flags) { +_CLC_DEF _CLC_OVERLOAD void barrier(cl_mem_fence_flags flags) { __syncthreads(); } - diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/ptx-nvidiacl/lib/workitem/get_global_id.cl new/libclc-12.0.0.src/ptx-nvidiacl/lib/workitem/get_global_id.cl --- old/libclc-11.0.1.src/ptx-nvidiacl/lib/workitem/get_global_id.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/ptx-nvidiacl/lib/workitem/get_global_id.cl 2021-04-06 18:38:18.000000000 +0200 @@ -1,5 +1,5 @@ #include <clc/clc.h> -_CLC_DEF size_t get_global_id(uint dim) { +_CLC_DEF _CLC_OVERLOAD size_t get_global_id(uint dim) { return get_group_id(dim) * get_local_size(dim) + get_local_id(dim); } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/ptx-nvidiacl/lib/workitem/get_group_id.cl new/libclc-12.0.0.src/ptx-nvidiacl/lib/workitem/get_group_id.cl --- old/libclc-11.0.1.src/ptx-nvidiacl/lib/workitem/get_group_id.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/ptx-nvidiacl/lib/workitem/get_group_id.cl 2021-04-06 18:38:18.000000000 +0200 @@ -1,6 +1,6 @@ #include <clc/clc.h> -_CLC_DEF size_t get_group_id(uint dim) { +_CLC_DEF _CLC_OVERLOAD size_t get_group_id(uint dim) { switch (dim) { case 0: return __nvvm_read_ptx_sreg_ctaid_x(); case 1: return __nvvm_read_ptx_sreg_ctaid_y(); diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/ptx-nvidiacl/lib/workitem/get_local_id.cl new/libclc-12.0.0.src/ptx-nvidiacl/lib/workitem/get_local_id.cl --- old/libclc-11.0.1.src/ptx-nvidiacl/lib/workitem/get_local_id.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/ptx-nvidiacl/lib/workitem/get_local_id.cl 2021-04-06 18:38:18.000000000 +0200 @@ -1,6 +1,6 @@ #include <clc/clc.h> -_CLC_DEF size_t get_local_id(uint dim) { +_CLC_DEF _CLC_OVERLOAD size_t get_local_id(uint dim) { switch (dim) { case 0: return __nvvm_read_ptx_sreg_tid_x(); case 1: return __nvvm_read_ptx_sreg_tid_y(); diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/ptx-nvidiacl/lib/workitem/get_local_size.cl new/libclc-12.0.0.src/ptx-nvidiacl/lib/workitem/get_local_size.cl --- old/libclc-11.0.1.src/ptx-nvidiacl/lib/workitem/get_local_size.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/ptx-nvidiacl/lib/workitem/get_local_size.cl 2021-04-06 18:38:18.000000000 +0200 @@ -1,6 +1,6 @@ #include <clc/clc.h> -_CLC_DEF size_t get_local_size(uint dim) { +_CLC_DEF _CLC_OVERLOAD size_t get_local_size(uint dim) { switch (dim) { case 0: return __nvvm_read_ptx_sreg_ntid_x(); case 1: return __nvvm_read_ptx_sreg_ntid_y(); diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/ptx-nvidiacl/lib/workitem/get_num_groups.cl new/libclc-12.0.0.src/ptx-nvidiacl/lib/workitem/get_num_groups.cl --- old/libclc-11.0.1.src/ptx-nvidiacl/lib/workitem/get_num_groups.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/ptx-nvidiacl/lib/workitem/get_num_groups.cl 2021-04-06 18:38:18.000000000 +0200 @@ -1,6 +1,6 @@ #include <clc/clc.h> -_CLC_DEF size_t get_num_groups(uint dim) { +_CLC_DEF _CLC_OVERLOAD size_t get_num_groups(uint dim) { switch (dim) { case 0: return __nvvm_read_ptx_sreg_nctaid_x(); case 1: return __nvvm_read_ptx_sreg_nctaid_y(); diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/r600/lib/synchronization/barrier.cl new/libclc-12.0.0.src/r600/lib/synchronization/barrier.cl --- old/libclc-11.0.1.src/r600/lib/synchronization/barrier.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/r600/lib/synchronization/barrier.cl 2021-04-06 18:38:18.000000000 +0200 @@ -2,8 +2,7 @@ _CLC_DEF void __clc_r600_barrier(void) __asm("llvm.r600.group.barrier"); -_CLC_DEF void barrier(uint flags) -{ +_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 -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/r600/lib/workitem/get_global_offset.cl new/libclc-12.0.0.src/r600/lib/workitem/get_global_offset.cl --- old/libclc-11.0.1.src/r600/lib/workitem/get_global_offset.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/r600/lib/workitem/get_global_offset.cl 2021-04-06 18:38:18.000000000 +0200 @@ -1,11 +1,10 @@ #include <clc/clc.h> -_CLC_DEF 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; +_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 -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/r600/lib/workitem/get_global_size.cl new/libclc-12.0.0.src/r600/lib/workitem/get_global_size.cl --- old/libclc-11.0.1.src/r600/lib/workitem/get_global_size.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/r600/lib/workitem/get_global_size.cl 2021-04-06 18:38:18.000000000 +0200 @@ -4,12 +4,15 @@ 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 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; - } +_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 -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/r600/lib/workitem/get_group_id.cl new/libclc-12.0.0.src/r600/lib/workitem/get_group_id.cl --- old/libclc-11.0.1.src/r600/lib/workitem/get_group_id.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/r600/lib/workitem/get_group_id.cl 2021-04-06 18:38:18.000000000 +0200 @@ -1,11 +1,14 @@ #include <clc/clc.h> -_CLC_DEF 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; - } +_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 -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/r600/lib/workitem/get_local_id.cl new/libclc-12.0.0.src/r600/lib/workitem/get_local_id.cl --- old/libclc-11.0.1.src/r600/lib/workitem/get_local_id.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/r600/lib/workitem/get_local_id.cl 2021-04-06 18:38:18.000000000 +0200 @@ -1,11 +1,14 @@ #include <clc/clc.h> -_CLC_DEF 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; - } +_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 -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/r600/lib/workitem/get_local_size.cl new/libclc-12.0.0.src/r600/lib/workitem/get_local_size.cl --- old/libclc-11.0.1.src/r600/lib/workitem/get_local_size.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/r600/lib/workitem/get_local_size.cl 2021-04-06 18:38:18.000000000 +0200 @@ -4,12 +4,15 @@ 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 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; - } +_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 -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/r600/lib/workitem/get_num_groups.cl new/libclc-12.0.0.src/r600/lib/workitem/get_num_groups.cl --- old/libclc-11.0.1.src/r600/lib/workitem/get_num_groups.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/r600/lib/workitem/get_num_groups.cl 2021-04-06 18:38:18.000000000 +0200 @@ -4,12 +4,15 @@ 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 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; - } +_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 -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/r600/lib/workitem/get_work_dim.cl new/libclc-12.0.0.src/r600/lib/workitem/get_work_dim.cl --- old/libclc-11.0.1.src/r600/lib/workitem/get_work_dim.cl 2020-12-18 20:57:38.000000000 +0100 +++ new/libclc-12.0.0.src/r600/lib/workitem/get_work_dim.cl 2021-04-06 18:38:18.000000000 +0200 @@ -1,9 +1,8 @@ #include <clc/clc.h> -_CLC_DEF uint get_work_dim(void) -{ - __attribute__((address_space(7))) uint * ptr = - (__attribute__((address_space(7))) uint *) - __builtin_r600_implicitarg_ptr(); - return ptr[0]; +_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]; } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/spirv/lib/SOURCES new/libclc-12.0.0.src/spirv/lib/SOURCES --- old/libclc-11.0.1.src/spirv/lib/SOURCES 1970-01-01 01:00:00.000000000 +0100 +++ new/libclc-12.0.0.src/spirv/lib/SOURCES 2021-04-06 18:38:18.000000000 +0200 @@ -0,0 +1,90 @@ +subnormal_config.cl +../../generic/lib/async/async_work_group_strided_copy.cl +../../generic/lib/async/wait_group_events.cl +../../generic/lib/common/degrees.cl +../../generic/lib/common/mix.cl +../../generic/lib/common/radians.cl +../../generic/lib/common/sign.cl +../../generic/lib/common/smoothstep.cl +../../generic/lib/common/step.cl +../../generic/lib/geometric/cross.cl +../../generic/lib/geometric/distance.cl +../../generic/lib/geometric/dot.cl +../../generic/lib/geometric/fast_distance.cl +../../generic/lib/geometric/fast_length.cl +../../generic/lib/geometric/fast_normalize.cl +../../generic/lib/geometric/length.cl +../../generic/lib/geometric/normalize.cl +../../generic/lib/integer/rotate.cl +../../generic/lib/integer/mad_sat.cl +../../generic/lib/math/acos.cl +../../generic/lib/math/acosh.cl +../../generic/lib/math/acospi.cl +../../generic/lib/math/asin.cl +../../generic/lib/math/asinh.cl +../../generic/lib/math/asinpi.cl +../../generic/lib/math/atan.cl +../../generic/lib/math/atan2.cl +../../generic/lib/math/atan2pi.cl +../../generic/lib/math/atanh.cl +../../generic/lib/math/atanpi.cl +../../generic/lib/math/cbrt.cl +../../generic/lib/math/cos.cl +../../generic/lib/math/cosh.cl +../../generic/lib/math/cospi.cl +../../generic/lib/math/ep_log.cl +../../generic/lib/math/erf.cl +../../generic/lib/math/erfc.cl +../../generic/lib/math/exp.cl +../../generic/lib/math/exp_helper.cl +../../generic/lib/math/expm1.cl +../../generic/lib/math/exp2.cl +../../generic/lib/math/clc_exp10.cl +../../generic/lib/math/exp10.cl +../../generic/lib/math/clc_fma.cl +math/fma.cl +../../generic/lib/math/clc_fmod.cl +../../generic/lib/math/fmod.cl +../../generic/lib/math/fract.cl +../../generic/lib/math/frexp.cl +../../generic/lib/math/half_rsqrt.cl +../../generic/lib/math/half_sqrt.cl +../../generic/lib/math/clc_hypot.cl +../../generic/lib/math/hypot.cl +../../generic/lib/math/ilogb.cl +../../generic/lib/math/clc_ldexp.cl +../../generic/lib/math/ldexp.cl +../../generic/lib/math/lgamma.cl +../../generic/lib/math/lgamma_r.cl +../../generic/lib/math/log.cl +../../generic/lib/math/log10.cl +../../generic/lib/math/log1p.cl +../../generic/lib/math/log2.cl +../../generic/lib/math/logb.cl +../../generic/lib/math/modf.cl +../../generic/lib/math/tables.cl +../../generic/lib/math/clc_pow.cl +../../generic/lib/math/pow.cl +../../generic/lib/math/clc_pown.cl +../../generic/lib/math/pown.cl +../../generic/lib/math/clc_powr.cl +../../generic/lib/math/powr.cl +../../generic/lib/math/clc_remainder.cl +../../generic/lib/math/remainder.cl +../../generic/lib/math/clc_remquo.cl +../../generic/lib/math/remquo.cl +../../generic/lib/math/clc_rootn.cl +../../generic/lib/math/rootn.cl +../../generic/lib/math/sin.cl +../../generic/lib/math/sincos.cl +../../generic/lib/math/sincos_helpers.cl +../../generic/lib/math/sinh.cl +../../generic/lib/math/sinpi.cl +../../generic/lib/math/clc_tan.cl +../../generic/lib/math/tan.cl +../../generic/lib/math/tanh.cl +../../generic/lib/math/clc_tanpi.cl +../../generic/lib/math/tanpi.cl +../../generic/lib/math/tgamma.cl +../../generic/lib/shared/vload.cl +../../generic/lib/shared/vstore.cl diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/spirv/lib/math/fma.cl new/libclc-12.0.0.src/spirv/lib/math/fma.cl --- old/libclc-11.0.1.src/spirv/lib/math/fma.cl 1970-01-01 01:00:00.000000000 +0100 +++ new/libclc-12.0.0.src/spirv/lib/math/fma.cl 2021-04-06 18:38:18.000000000 +0200 @@ -0,0 +1,11 @@ +#include <clc/clc.h> +#include <math/clc_fma.h> + +#define __CLC_BODY <fma.inc> +#define __FLOAT_ONLY +#include <clc/math/gentype.inc> + +bool __clc_runtime_has_hw_fma32() +{ + return false; +} diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/spirv/lib/math/fma.inc new/libclc-12.0.0.src/spirv/lib/math/fma.inc --- old/libclc-11.0.1.src/spirv/lib/math/fma.inc 1970-01-01 01:00:00.000000000 +0100 +++ new/libclc-12.0.0.src/spirv/lib/math/fma.inc 2021-04-06 18:38:18.000000000 +0200 @@ -0,0 +1,3 @@ +_CLC_DEF _CLC_OVERLOAD __CLC_GENTYPE fma(__CLC_GENTYPE a, __CLC_GENTYPE b, __CLC_GENTYPE c) { + return __clc_sw_fma(a, b, c); +} diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/spirv/lib/subnormal_config.cl new/libclc-12.0.0.src/spirv/lib/subnormal_config.cl --- old/libclc-11.0.1.src/spirv/lib/subnormal_config.cl 1970-01-01 01:00:00.000000000 +0100 +++ new/libclc-12.0.0.src/spirv/lib/subnormal_config.cl 2021-04-06 18:38:18.000000000 +0200 @@ -0,0 +1,31 @@ +/* + * Copyright (c) 2015 Advanced Micro Devices, Inc. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include <clc/clc.h> + +#include "config.h" + +_CLC_DEF bool __clc_fp16_subnormals_supported() { return false; } + +_CLC_DEF bool __clc_fp32_subnormals_supported() { return false; } + +_CLC_DEF bool __clc_fp64_subnormals_supported() { return false; } diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/spirv64/lib/SOURCES new/libclc-12.0.0.src/spirv64/lib/SOURCES --- old/libclc-11.0.1.src/spirv64/lib/SOURCES 1970-01-01 01:00:00.000000000 +0100 +++ new/libclc-12.0.0.src/spirv64/lib/SOURCES 2021-04-06 18:38:18.000000000 +0200 @@ -0,0 +1,90 @@ +subnormal_config.cl +../../generic/lib/async/async_work_group_strided_copy.cl +../../generic/lib/async/wait_group_events.cl +../../generic/lib/common/degrees.cl +../../generic/lib/common/mix.cl +../../generic/lib/common/radians.cl +../../generic/lib/common/sign.cl +../../generic/lib/common/smoothstep.cl +../../generic/lib/common/step.cl +../../generic/lib/geometric/cross.cl +../../generic/lib/geometric/distance.cl +../../generic/lib/geometric/dot.cl +../../generic/lib/geometric/fast_distance.cl +../../generic/lib/geometric/fast_length.cl +../../generic/lib/geometric/fast_normalize.cl +../../generic/lib/geometric/length.cl +../../generic/lib/geometric/normalize.cl +../../generic/lib/integer/rotate.cl +../../generic/lib/integer/mad_sat.cl +../../generic/lib/math/acos.cl +../../generic/lib/math/acosh.cl +../../generic/lib/math/acospi.cl +../../generic/lib/math/asin.cl +../../generic/lib/math/asinh.cl +../../generic/lib/math/asinpi.cl +../../generic/lib/math/atan.cl +../../generic/lib/math/atan2.cl +../../generic/lib/math/atan2pi.cl +../../generic/lib/math/atanh.cl +../../generic/lib/math/atanpi.cl +../../generic/lib/math/cbrt.cl +../../generic/lib/math/cos.cl +../../generic/lib/math/cosh.cl +../../generic/lib/math/cospi.cl +../../generic/lib/math/ep_log.cl +../../generic/lib/math/erf.cl +../../generic/lib/math/erfc.cl +../../generic/lib/math/exp.cl +../../generic/lib/math/exp_helper.cl +../../generic/lib/math/expm1.cl +../../generic/lib/math/exp2.cl +../../generic/lib/math/clc_exp10.cl +../../generic/lib/math/exp10.cl +../../generic/lib/math/clc_fma.cl +math/fma.cl +../../generic/lib/math/clc_fmod.cl +../../generic/lib/math/fmod.cl +../../generic/lib/math/fract.cl +../../generic/lib/math/frexp.cl +../../generic/lib/math/half_rsqrt.cl +../../generic/lib/math/half_sqrt.cl +../../generic/lib/math/clc_hypot.cl +../../generic/lib/math/hypot.cl +../../generic/lib/math/ilogb.cl +../../generic/lib/math/clc_ldexp.cl +../../generic/lib/math/ldexp.cl +../../generic/lib/math/lgamma.cl +../../generic/lib/math/lgamma_r.cl +../../generic/lib/math/log.cl +../../generic/lib/math/log10.cl +../../generic/lib/math/log1p.cl +../../generic/lib/math/log2.cl +../../generic/lib/math/logb.cl +../../generic/lib/math/modf.cl +../../generic/lib/math/tables.cl +../../generic/lib/math/clc_pow.cl +../../generic/lib/math/pow.cl +../../generic/lib/math/clc_pown.cl +../../generic/lib/math/pown.cl +../../generic/lib/math/clc_powr.cl +../../generic/lib/math/powr.cl +../../generic/lib/math/clc_remainder.cl +../../generic/lib/math/remainder.cl +../../generic/lib/math/clc_remquo.cl +../../generic/lib/math/remquo.cl +../../generic/lib/math/clc_rootn.cl +../../generic/lib/math/rootn.cl +../../generic/lib/math/sin.cl +../../generic/lib/math/sincos.cl +../../generic/lib/math/sincos_helpers.cl +../../generic/lib/math/sinh.cl +../../generic/lib/math/sinpi.cl +../../generic/lib/math/clc_tan.cl +../../generic/lib/math/tan.cl +../../generic/lib/math/tanh.cl +../../generic/lib/math/clc_tanpi.cl +../../generic/lib/math/tanpi.cl +../../generic/lib/math/tgamma.cl +../../generic/lib/shared/vload.cl +../../generic/lib/shared/vstore.cl diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/spirv64/lib/math/fma.cl new/libclc-12.0.0.src/spirv64/lib/math/fma.cl --- old/libclc-11.0.1.src/spirv64/lib/math/fma.cl 1970-01-01 01:00:00.000000000 +0100 +++ new/libclc-12.0.0.src/spirv64/lib/math/fma.cl 2021-04-06 18:38:18.000000000 +0200 @@ -0,0 +1,11 @@ +#include <clc/clc.h> +#include <math/clc_fma.h> + +#define __CLC_BODY <fma.inc> +#define __FLOAT_ONLY +#include <clc/math/gentype.inc> + +bool __clc_runtime_has_hw_fma32() +{ + return false; +} diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/spirv64/lib/math/fma.inc new/libclc-12.0.0.src/spirv64/lib/math/fma.inc --- old/libclc-11.0.1.src/spirv64/lib/math/fma.inc 1970-01-01 01:00:00.000000000 +0100 +++ new/libclc-12.0.0.src/spirv64/lib/math/fma.inc 2021-04-06 18:38:18.000000000 +0200 @@ -0,0 +1,3 @@ +_CLC_DEF _CLC_OVERLOAD __CLC_GENTYPE fma(__CLC_GENTYPE a, __CLC_GENTYPE b, __CLC_GENTYPE c) { + return __clc_sw_fma(a, b, c); +} diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/libclc-11.0.1.src/spirv64/lib/subnormal_config.cl new/libclc-12.0.0.src/spirv64/lib/subnormal_config.cl --- old/libclc-11.0.1.src/spirv64/lib/subnormal_config.cl 1970-01-01 01:00:00.000000000 +0100 +++ new/libclc-12.0.0.src/spirv64/lib/subnormal_config.cl 2021-04-06 18:38:18.000000000 +0200 @@ -0,0 +1,31 @@ +/* + * Copyright (c) 2015 Advanced Micro Devices, Inc. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include <clc/clc.h> + +#include "config.h" + +_CLC_DEF bool __clc_fp16_subnormals_supported() { return false; } + +_CLC_DEF bool __clc_fp32_subnormals_supported() { return false; } + +_CLC_DEF bool __clc_fp64_subnormals_supported() { return false; }
