llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang-driver Author: Jiang zixian (jiang-zixian) <details> <summary>Changes</summary> the decoding and encoding are complete --- Patch is 430.17 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/87187.diff 101 Files Affected: - (removed) clang/test/Driver/Inputs/CUDA-symlinks/usr/bin/ptxas (-1) - (added) clang/test/Driver/Inputs/CUDA-symlinks/usr/bin/ptxas () - (removed) clang/test/Driver/Inputs/basic_cross_linux_tree/usr/bin/i386-unknown-linux-gnu-ld (-1) - (added) clang/test/Driver/Inputs/basic_cross_linux_tree/usr/bin/i386-unknown-linux-gnu-ld (+1) - (removed) clang/test/Driver/Inputs/basic_cross_linux_tree/usr/bin/x86_64-unknown-linux-gnu-ld (-1) - (added) clang/test/Driver/Inputs/basic_cross_linux_tree/usr/bin/x86_64-unknown-linux-gnu-ld (+1) - (removed) clang/test/Driver/Inputs/basic_cross_linux_tree/usr/i386-unknown-linux-gnu/bin/ld (-1) - (added) clang/test/Driver/Inputs/basic_cross_linux_tree/usr/i386-unknown-linux-gnu/bin/ld (+1) - (removed) clang/test/Driver/Inputs/basic_cross_linux_tree/usr/x86_64-unknown-linux-gnu/bin/ld (-1) - (added) clang/test/Driver/Inputs/basic_cross_linux_tree/usr/x86_64-unknown-linux-gnu/bin/ld (+1) - (removed) clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/bin/as (-1) - (added) clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/bin/as (+1) - (removed) clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/bin/ld (-1) - (added) clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/bin/ld (+1) - (removed) clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/i386-unknown-linux/bin/as (-1) - (added) clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/i386-unknown-linux/bin/as (+1) - (removed) clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/i386-unknown-linux/bin/ld (-1) - (added) clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/i386-unknown-linux/bin/ld (+1) - (removed) clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/bin/as (-1) - (added) clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/bin/as (+1) - (removed) clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/bin/ld (-1) - (added) clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/bin/ld (+1) - (removed) clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/x86_64-unknown-linux/bin/as (-1) - (added) clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/x86_64-unknown-linux/bin/as (+1) - (removed) clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/x86_64-unknown-linux/bin/ld (-1) - (added) clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/x86_64-unknown-linux/bin/ld (+1) - (removed) libclc/amdgcn-mesa3d (-1) - (added) libclc/amdgcn-mesa3d/lib/SOURCES (+3) - (added) libclc/amdgcn-mesa3d/lib/workitem/get_global_size.cl (+23) - (added) libclc/amdgcn-mesa3d/lib/workitem/get_local_size.cl (+29) - (added) libclc/amdgcn-mesa3d/lib/workitem/get_num_groups.cl (+12) - (removed) libclc/clspv64 (-1) - (added) libclc/clspv64/lib/SOURCES (+48) - (added) libclc/clspv64/lib/math/fma.cl (+256) - (added) libclc/clspv64/lib/math/nextafter.cl (+5) - (added) libclc/clspv64/lib/math/nextafter.inc (+3) - (added) libclc/clspv64/lib/subnormal_config.cl (+31) - (removed) libcxx/test/std/pstl (-1) - (added) libcxx/test/std/pstl/algorithms/alg.merge/inplace_merge.pass.cpp (+159) - (added) libcxx/test/std/pstl/algorithms/alg.merge/merge.pass.cpp (+113) - (added) libcxx/test/std/pstl/algorithms/alg.modifying.operations/alg.copy/copy_if.pass.cpp (+147) - (added) libcxx/test/std/pstl/algorithms/alg.modifying.operations/alg.partitions/is_partitioned.pass.cpp (+101) - (added) libcxx/test/std/pstl/algorithms/alg.modifying.operations/alg.partitions/partition.pass.cpp (+178) - (added) libcxx/test/std/pstl/algorithms/alg.modifying.operations/alg.partitions/partition_copy.pass.cpp (+116) - (added) libcxx/test/std/pstl/algorithms/alg.modifying.operations/alg.reverse/reverse.pass.cpp (+104) - (added) libcxx/test/std/pstl/algorithms/alg.modifying.operations/alg.reverse/reverse_copy.pass.cpp (+130) - (added) libcxx/test/std/pstl/algorithms/alg.modifying.operations/copy_move.pass.cpp (+197) - (added) libcxx/test/std/pstl/algorithms/alg.modifying.operations/fill.pass.cpp (+100) - (added) libcxx/test/std/pstl/algorithms/alg.modifying.operations/generate.pass.cpp (+104) - (added) libcxx/test/std/pstl/algorithms/alg.modifying.operations/remove.pass.cpp (+161) - (added) libcxx/test/std/pstl/algorithms/alg.modifying.operations/remove_copy.pass.cpp (+91) - (added) libcxx/test/std/pstl/algorithms/alg.modifying.operations/replace.pass.cpp (+160) - (added) libcxx/test/std/pstl/algorithms/alg.modifying.operations/replace_copy.pass.cpp (+105) - (added) libcxx/test/std/pstl/algorithms/alg.modifying.operations/rotate.pass.cpp (+176) - (added) libcxx/test/std/pstl/algorithms/alg.modifying.operations/rotate_copy.pass.cpp (+146) - (added) libcxx/test/std/pstl/algorithms/alg.modifying.operations/swap_ranges.pass.cpp (+133) - (added) libcxx/test/std/pstl/algorithms/alg.modifying.operations/transform_binary.pass.cpp (+122) - (added) libcxx/test/std/pstl/algorithms/alg.modifying.operations/transform_unary.pass.cpp (+91) - (added) libcxx/test/std/pstl/algorithms/alg.modifying.operations/unique.pass.cpp (+163) - (added) libcxx/test/std/pstl/algorithms/alg.modifying.operations/unique_copy_equal.pass.cpp (+135) - (added) libcxx/test/std/pstl/algorithms/alg.nonmodifying/adjacent_find.pass.cpp (+114) - (added) libcxx/test/std/pstl/algorithms/alg.nonmodifying/all_of.pass.cpp (+117) - (added) libcxx/test/std/pstl/algorithms/alg.nonmodifying/any_of.pass.cpp (+103) - (added) libcxx/test/std/pstl/algorithms/alg.nonmodifying/count.pass.cpp (+108) - (added) libcxx/test/std/pstl/algorithms/alg.nonmodifying/equal.pass.cpp (+168) - (added) libcxx/test/std/pstl/algorithms/alg.nonmodifying/find.pass.cpp (+96) - (added) libcxx/test/std/pstl/algorithms/alg.nonmodifying/find_end.pass.cpp (+123) - (added) libcxx/test/std/pstl/algorithms/alg.nonmodifying/find_first_of.pass.cpp (+112) - (added) libcxx/test/std/pstl/algorithms/alg.nonmodifying/find_if.pass.cpp (+109) - (added) libcxx/test/std/pstl/algorithms/alg.nonmodifying/for_each.pass.cpp (+102) - (added) libcxx/test/std/pstl/algorithms/alg.nonmodifying/mismatch.pass.cpp (+132) - (added) libcxx/test/std/pstl/algorithms/alg.nonmodifying/none_of.pass.cpp (+101) - (added) libcxx/test/std/pstl/algorithms/alg.nonmodifying/nth_element.pass.cpp (+175) - (added) libcxx/test/std/pstl/algorithms/alg.nonmodifying/search_n.pass.cpp (+109) - (added) libcxx/test/std/pstl/algorithms/alg.sorting/alg.heap.operations/is_heap.pass.cpp (+146) - (added) libcxx/test/std/pstl/algorithms/alg.sorting/alg.lex.comparison/lexicographical_compare.pass.cpp (+175) - (added) libcxx/test/std/pstl/algorithms/alg.sorting/alg.min.max/minmax_element.pass.cpp (+192) - (added) libcxx/test/std/pstl/algorithms/alg.sorting/alg.set.operations/includes.pass.cpp (+106) - (added) libcxx/test/std/pstl/algorithms/alg.sorting/alg.set.operations/set.pass.cpp (+280) - (added) libcxx/test/std/pstl/algorithms/alg.sorting/is_sorted.pass.cpp (+100) - (added) libcxx/test/std/pstl/algorithms/alg.sorting/partial_sort.pass.cpp (+149) - (added) libcxx/test/std/pstl/algorithms/alg.sorting/partial_sort_copy.pass.cpp (+196) - (added) libcxx/test/std/pstl/algorithms/alg.sorting/sort.pass.cpp (+247) - (added) libcxx/test/std/pstl/lit.local.cfg (+2) - (added) libcxx/test/std/pstl/numerics/numeric.ops/adjacent_difference.pass.cpp (+170) - (added) libcxx/test/std/pstl/numerics/numeric.ops/reduce.pass.cpp (+114) - (added) libcxx/test/std/pstl/numerics/numeric.ops/scan.fail.cpp (+28) - (added) libcxx/test/std/pstl/numerics/numeric.ops/scan.pass.cpp (+201) - (added) libcxx/test/std/pstl/numerics/numeric.ops/transform_reduce.pass.cpp (+129) - (added) libcxx/test/std/pstl/numerics/numeric.ops/transform_scan.pass.cpp (+177) - (added) libcxx/test/std/pstl/utilities/memory/specialized.algorithms/uninitialized_construct.pass.cpp (+123) - (added) libcxx/test/std/pstl/utilities/memory/specialized.algorithms/uninitialized_copy_move.pass.cpp (+143) - (added) libcxx/test/std/pstl/utilities/memory/specialized.algorithms/uninitialized_fill_destroy.pass.cpp (+93) - (modified) lld/ELF/Arch/ARM.cpp (+12) - (modified) llvm/include/llvm/CodeGen/MachineBasicBlock.h (+14) - (added) llvm/lib/Target/ARM/ARMEncodeDecode.cpp (+628) - (added) llvm/lib/Target/ARM/ARMEncodeDecode.h (+38) - (modified) llvm/lib/Target/ARM/ARMInstrThumb2.td (+21) - (added) llvm/lib/Target/ARM/ARMRandezvousCLR.cpp (+335) - (added) llvm/lib/Target/ARM/ARMRandezvousCLR.h (+54) - (added) llvm/lib/Target/ARM/ARMRandezvousInstrumentor.cpp (+824) - (added) llvm/lib/Target/ARM/ARMRandezvousInstrumentor.h (+178) - (added) llvm/lib/Target/ARM/ARMRandezvousOptions.cpp (+246) - (added) llvm/lib/Target/ARM/ARMRandezvousOptions.h (+67) - (added) llvm/lib/Target/ARM/ARMRandezvousShadowStack.cpp (+861) - (added) llvm/lib/Target/ARM/ARMRandezvousShadowStack.h (+40) - (modified) llvm/lib/Target/ARM/ARMTargetMachine.cpp (+17) - (modified) llvm/lib/Target/ARM/CMakeLists.txt (+13) - (modified) llvm/lib/Target/ARM/MCTargetDesc/ARMAsmBackend.cpp (+30) - (modified) llvm/lib/Target/ARM/MCTargetDesc/ARMELFObjectWriter.cpp (+12) - (modified) llvm/lib/Target/ARM/MCTargetDesc/ARMFixupKinds.h (+2-1) - (modified) llvm/lib/Target/ARM/MCTargetDesc/ARMMCCodeEmitter.cpp (+4) - (modified) llvm/lib/Transforms/Hello/Hello.cpp (+1) - (removed) openmp/tools/analyzer/llvm-openmp-analyzer++ (-1) - (added) openmp/tools/analyzer/llvm-openmp-analyzer++ (+45) ``````````diff diff --git a/clang/test/Driver/Inputs/CUDA-symlinks/usr/bin/ptxas b/clang/test/Driver/Inputs/CUDA-symlinks/usr/bin/ptxas deleted file mode 120000 index 59eefd95a9023c..00000000000000 --- a/clang/test/Driver/Inputs/CUDA-symlinks/usr/bin/ptxas +++ /dev/null @@ -1 +0,0 @@ -../../opt/cuda/bin/ptxas \ No newline at end of file diff --git a/clang/test/Driver/Inputs/CUDA-symlinks/usr/bin/ptxas b/clang/test/Driver/Inputs/CUDA-symlinks/usr/bin/ptxas new file mode 100755 index 00000000000000..e69de29bb2d1d6 diff --git a/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/bin/i386-unknown-linux-gnu-ld b/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/bin/i386-unknown-linux-gnu-ld deleted file mode 120000 index 7e0a9cfe2ddbd6..00000000000000 --- a/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/bin/i386-unknown-linux-gnu-ld +++ /dev/null @@ -1 +0,0 @@ -i386-unknown-linux-gnu-ld.gold \ No newline at end of file diff --git a/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/bin/i386-unknown-linux-gnu-ld b/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/bin/i386-unknown-linux-gnu-ld new file mode 100755 index 00000000000000..b23e55619b2ff0 --- /dev/null +++ b/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/bin/i386-unknown-linux-gnu-ld @@ -0,0 +1 @@ +#!/bin/true diff --git a/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/bin/x86_64-unknown-linux-gnu-ld b/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/bin/x86_64-unknown-linux-gnu-ld deleted file mode 120000 index ce36ac093b6176..00000000000000 --- a/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/bin/x86_64-unknown-linux-gnu-ld +++ /dev/null @@ -1 +0,0 @@ -x86_64-unknown-linux-gnu-ld.gold \ No newline at end of file diff --git a/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/bin/x86_64-unknown-linux-gnu-ld b/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/bin/x86_64-unknown-linux-gnu-ld new file mode 100755 index 00000000000000..b23e55619b2ff0 --- /dev/null +++ b/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/bin/x86_64-unknown-linux-gnu-ld @@ -0,0 +1 @@ +#!/bin/true diff --git a/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/i386-unknown-linux-gnu/bin/ld b/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/i386-unknown-linux-gnu/bin/ld deleted file mode 120000 index 6cd03701cdda78..00000000000000 --- a/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/i386-unknown-linux-gnu/bin/ld +++ /dev/null @@ -1 +0,0 @@ -ld.gold \ No newline at end of file diff --git a/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/i386-unknown-linux-gnu/bin/ld b/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/i386-unknown-linux-gnu/bin/ld new file mode 100755 index 00000000000000..b23e55619b2ff0 --- /dev/null +++ b/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/i386-unknown-linux-gnu/bin/ld @@ -0,0 +1 @@ +#!/bin/true diff --git a/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/x86_64-unknown-linux-gnu/bin/ld b/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/x86_64-unknown-linux-gnu/bin/ld deleted file mode 120000 index 6cd03701cdda78..00000000000000 --- a/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/x86_64-unknown-linux-gnu/bin/ld +++ /dev/null @@ -1 +0,0 @@ -ld.gold \ No newline at end of file diff --git a/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/x86_64-unknown-linux-gnu/bin/ld b/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/x86_64-unknown-linux-gnu/bin/ld new file mode 100755 index 00000000000000..b23e55619b2ff0 --- /dev/null +++ b/clang/test/Driver/Inputs/basic_cross_linux_tree/usr/x86_64-unknown-linux-gnu/bin/ld @@ -0,0 +1 @@ +#!/bin/true diff --git a/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/bin/as b/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/bin/as deleted file mode 120000 index 0065315cfd1de8..00000000000000 --- a/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/bin/as +++ /dev/null @@ -1 +0,0 @@ -i386-unknown-linux-gnu-as \ No newline at end of file diff --git a/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/bin/as b/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/bin/as new file mode 100755 index 00000000000000..b23e55619b2ff0 --- /dev/null +++ b/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/bin/as @@ -0,0 +1 @@ +#!/bin/true diff --git a/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/bin/ld b/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/bin/ld deleted file mode 120000 index 9e5574285c70e4..00000000000000 --- a/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/bin/ld +++ /dev/null @@ -1 +0,0 @@ -i386-unknown-linux-gnu-ld \ No newline at end of file diff --git a/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/bin/ld b/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/bin/ld new file mode 100755 index 00000000000000..b23e55619b2ff0 --- /dev/null +++ b/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/bin/ld @@ -0,0 +1 @@ +#!/bin/true diff --git a/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/i386-unknown-linux/bin/as b/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/i386-unknown-linux/bin/as deleted file mode 120000 index 2aa12fdef91620..00000000000000 --- a/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/i386-unknown-linux/bin/as +++ /dev/null @@ -1 +0,0 @@ -../../bin/i386-unknown-linux-gnu-as \ No newline at end of file diff --git a/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/i386-unknown-linux/bin/as b/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/i386-unknown-linux/bin/as new file mode 100755 index 00000000000000..b23e55619b2ff0 --- /dev/null +++ b/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/i386-unknown-linux/bin/as @@ -0,0 +1 @@ +#!/bin/true diff --git a/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/i386-unknown-linux/bin/ld b/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/i386-unknown-linux/bin/ld deleted file mode 120000 index 5aeaff619662a8..00000000000000 --- a/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/i386-unknown-linux/bin/ld +++ /dev/null @@ -1 +0,0 @@ -../../bin/i386-unknown-linux-gnu-ld \ No newline at end of file diff --git a/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/i386-unknown-linux/bin/ld b/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/i386-unknown-linux/bin/ld new file mode 100755 index 00000000000000..b23e55619b2ff0 --- /dev/null +++ b/clang/test/Driver/Inputs/multilib_32bit_linux_tree/usr/i386-unknown-linux/bin/ld @@ -0,0 +1 @@ +#!/bin/true diff --git a/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/bin/as b/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/bin/as deleted file mode 120000 index 477cbc9635fcbf..00000000000000 --- a/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/bin/as +++ /dev/null @@ -1 +0,0 @@ -x86_64-unknown-linux-gnu-as \ No newline at end of file diff --git a/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/bin/as b/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/bin/as new file mode 100755 index 00000000000000..b23e55619b2ff0 --- /dev/null +++ b/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/bin/as @@ -0,0 +1 @@ +#!/bin/true diff --git a/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/bin/ld b/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/bin/ld deleted file mode 120000 index 5343caf34d8f34..00000000000000 --- a/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/bin/ld +++ /dev/null @@ -1 +0,0 @@ -x86_64-unknown-linux-gnu-ld \ No newline at end of file diff --git a/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/bin/ld b/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/bin/ld new file mode 100755 index 00000000000000..b23e55619b2ff0 --- /dev/null +++ b/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/bin/ld @@ -0,0 +1 @@ +#!/bin/true diff --git a/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/x86_64-unknown-linux/bin/as b/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/x86_64-unknown-linux/bin/as deleted file mode 120000 index 84a9113f2671f4..00000000000000 --- a/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/x86_64-unknown-linux/bin/as +++ /dev/null @@ -1 +0,0 @@ -../../bin/x86_64-unknown-linux-gnu-as \ No newline at end of file diff --git a/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/x86_64-unknown-linux/bin/as b/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/x86_64-unknown-linux/bin/as new file mode 100755 index 00000000000000..b23e55619b2ff0 --- /dev/null +++ b/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/x86_64-unknown-linux/bin/as @@ -0,0 +1 @@ +#!/bin/true diff --git a/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/x86_64-unknown-linux/bin/ld b/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/x86_64-unknown-linux/bin/ld deleted file mode 120000 index c417e3afaa4945..00000000000000 --- a/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/x86_64-unknown-linux/bin/ld +++ /dev/null @@ -1 +0,0 @@ -../../bin/x86_64-unknown-linux-gnu-ld \ No newline at end of file diff --git a/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/x86_64-unknown-linux/bin/ld b/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/x86_64-unknown-linux/bin/ld new file mode 100755 index 00000000000000..b23e55619b2ff0 --- /dev/null +++ b/clang/test/Driver/Inputs/multilib_64bit_linux_tree/usr/x86_64-unknown-linux/bin/ld @@ -0,0 +1 @@ +#!/bin/true diff --git a/libclc/amdgcn-mesa3d b/libclc/amdgcn-mesa3d deleted file mode 120000 index 400782833efe6c..00000000000000 --- a/libclc/amdgcn-mesa3d +++ /dev/null @@ -1 +0,0 @@ -amdgcn-amdhsa \ No newline at end of file diff --git a/libclc/amdgcn-mesa3d/lib/SOURCES b/libclc/amdgcn-mesa3d/lib/SOURCES new file mode 100644 index 00000000000000..8224b7721b2ca5 --- /dev/null +++ b/libclc/amdgcn-mesa3d/lib/SOURCES @@ -0,0 +1,3 @@ +workitem/get_global_size.cl +workitem/get_local_size.cl +workitem/get_num_groups.cl diff --git a/libclc/amdgcn-mesa3d/lib/workitem/get_global_size.cl b/libclc/amdgcn-mesa3d/lib/workitem/get_global_size.cl new file mode 100644 index 00000000000000..62bd2ba283523f --- /dev/null +++ b/libclc/amdgcn-mesa3d/lib/workitem/get_global_size.cl @@ -0,0 +1,23 @@ +#include <clc/clc.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(); + if (dim < 3) + return ptr[3 + dim]; + return 1; +} diff --git a/libclc/amdgcn-mesa3d/lib/workitem/get_local_size.cl b/libclc/amdgcn-mesa3d/lib/workitem/get_local_size.cl new file mode 100644 index 00000000000000..9f09fd5a16ec66 --- /dev/null +++ b/libclc/amdgcn-mesa3d/lib/workitem/get_local_size.cl @@ -0,0 +1,29 @@ +#include <clc/clc.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(); + switch (dim) { + case 0: + return ptr[1] & 0xffffu; + case 1: + return ptr[1] >> 16; + case 2: + return ptr[2] & 0xffffu; + } + return 1; +} diff --git a/libclc/amdgcn-mesa3d/lib/workitem/get_num_groups.cl b/libclc/amdgcn-mesa3d/lib/workitem/get_num_groups.cl new file mode 100644 index 00000000000000..35dc2218852114 --- /dev/null +++ b/libclc/amdgcn-mesa3d/lib/workitem/get_num_groups.cl @@ -0,0 +1,12 @@ + +#include <clc/clc.h> + +_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; + if (global_size % local_size != 0) { + num_groups++; + } + return num_groups; +} diff --git a/libclc/clspv64 b/libclc/clspv64 deleted file mode 120000 index ea01ba94bc6368..00000000000000 --- a/libclc/clspv64 +++ /dev/null @@ -1 +0,0 @@ -clspv \ No newline at end of file diff --git a/libclc/clspv64/lib/SOURCES b/libclc/clspv64/lib/SOURCES new file mode 100644 index 00000000000000..0466345cee0271 --- /dev/null +++ b/libclc/clspv64/lib/SOURCES @@ -0,0 +1,48 @@ +subnormal_config.cl +../../generic/lib/geometric/distance.cl +../../generic/lib/geometric/length.cl +math/fma.cl +math/nextafter.cl +../../generic/lib/math/acosh.cl +../../generic/lib/math/asinh.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/clc_fmod.cl +../../generic/lib/math/clc_hypot.cl +../../generic/lib/math/clc_ldexp.cl +../../generic/lib/math/clc_nextafter.cl +../../generic/lib/math/clc_remainder.cl +../../generic/lib/math/clc_remquo.cl +../../generic/lib/math/clc_rootn.cl +../../generic/lib/math/clc_sqrt.cl +../../generic/lib/math/clc_tan.cl +../../generic/lib/math/erf.cl +../../generic/lib/math/erfc.cl +../../generic/lib/math/fmod.cl +../../generic/lib/math/fract.cl +../../generic/lib/math/frexp.cl +../../generic/lib/math/half_divide.cl +../../generic/lib/math/half_recip.cl +../../generic/lib/math/half_sqrt.cl +../../generic/lib/math/hypot.cl +../../generic/lib/math/ilogb.cl +../../generic/lib/math/ldexp.cl +../../generic/lib/math/lgamma.cl +../../generic/lib/math/lgamma_r.cl +../../generic/lib/math/logb.cl +../../generic/lib/math/maxmag.cl +../../generic/lib/math/minmag.cl +../../generic/lib/math/modf.cl +../../generic/lib/math/nan.cl +../../generic/lib/math/remainder.cl +../../generic/lib/math/remquo.cl +../../generic/lib/math/rootn.cl +../../generic/lib/math/rsqrt.cl +../../generic/lib/math/sqrt.cl +../../generic/lib/math/tables.cl +../../generic/lib/math/tanh.cl +../../generic/lib/math/tgamma.cl diff --git a/libclc/clspv64/lib/math/fma.cl b/libclc/clspv64/lib/math/fma.cl new file mode 100644 index 00000000000000..fdc8b8b296876c --- /dev/null +++ b/libclc/clspv64/lib/math/fma.cl @@ -0,0 +1,256 @@ +/* + * Copyright (c) 2014 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. + */ + +// This version is derived from the generic fma software implementation +// (__clc_sw_fma), but avoids the use of ulong in favor of uint2. The logic has +// been updated as appropriate. + +#include <clc/clc.h> +#include "../../../generic/lib/clcmacro.h" +#include "../../../generic/lib/math/math.h" + +struct fp { + uint2 mantissa; + int exponent; + uint sign; +}; + +_CLC_DEF _CLC_OVERLOAD float fma(float a, float b, float c) { + /* special cases */ + if (isnan(a) || isnan(b) || isnan(c) || isinf(a) || isinf(b)) { + return mad(a, b, c); + } + + /* If only c is inf, and both a,b are regular numbers, the result is c*/ + if (isinf(c)) { + return c; + } + + a = __clc_flush_denormal_if_not_supported(a); + b = __clc_flush_denormal_if_not_supported(b); + c = __clc_flush_denormal_if_not_supported(c); + + if (a == 0.0f || b == 0.0f) { + return c; + } + + if (c == 0) { + return a * b; + } + + struct fp st_a, st_b, st_c; + + st_a.exponent = a == .0f ? 0 : ((as_uint(a) & 0x7f800000) >> 23) - 127; + st_b.exponent = b == .0f ? 0 : ((as_uint(b) & 0x7f800000) >> 23) - 127; + st_c.exponent = c == .0f ? 0 : ((as_uint(c) & 0x7f800000) >> 23) - 127; + + st_a.mantissa.lo = a == .0f ? 0 : (as_uint(a) & 0x7fffff) | 0x800000; + st_b.mantissa.lo = b == .0f ? 0 : (as_uint(b) & 0x7fffff) | 0x800000; + st_c.mantissa.lo = c == .0f ? 0 : (as_uint(c) & 0x7fffff) | 0x800000; + st_a.mantissa.hi = 0; + st_b.mantissa.hi = 0; + st_c.mantissa.hi = 0; + + st_a.sign = as_uint(a) & 0x80000000; + st_b.sign = as_uint(b) & 0x80000000; + st_c.sign = 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.hi = mul_hi(st_a.mantissa.lo, st_b.mantissa.lo); + st_mul.mantissa.lo = st_a.mantissa.lo * st_b.mantissa.lo; + uint upper_14bits = (st_mul.mantissa.lo >> 18) & 0x3fff; + st_mul.mantissa.lo <<= 14; + st_mul.mantissa.hi <<= 14; + st_mul.mantissa.hi |= upper_14bits; + st_mul.exponent = (st_mul.mantissa.lo != 0 || st_mul.mantissa.hi != 0) + ? st_a.exponent + st_b.exponent + : 0; + +// 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; + + uint abs_exp_diff = abs(exp_diff); + st_c.mantissa.hi = (st_c.mantissa.lo << 5); + st_c.mantissa.lo = 0; + uint2 cutoff_bits = (uint2)(0, 0); + uint2 cutoff_mask = (uint2)(0, 0); + if (abs_exp_diff < 32) { + cutoff_mask.lo = (1u << abs(exp_diff)) - 1u; + } else if (abs_exp_diff < 64) { + cutoff_mask.lo = 0xffffffff; + uint remaining = abs_exp_diff - 32; + cutoff_mask.hi = (1u << remaining) - 1u; + } else { + cutoff_mask = (uint2)(0, 0); + } + uint2 tmp = (exp_diff > 0) ? st_c.mantissa : st_mul.mantissa; + if (abs_exp_diff > 0) { + cutoff_bits = abs_exp_diff >= 64 ? tmp : (tmp & cutoff_mask); + if (abs_exp_diff < 32) { + // shift some of the hi bits into the shifted lo bits. + uint shift_mask = (1u << abs_exp_diff) - 1; + uint upper_saved_bits = tmp.hi & shift_mask; + upper_saved_bits = upper_saved_bits << (32 - abs_exp_diff); + tmp.hi >>= abs_exp_diff; + tmp.lo >>= abs_exp_diff; + tmp.lo |= upper_saved_bits; + } else if (abs_exp_diff < 64) { + tmp.lo = (tmp.hi >> (abs_exp_diff - 32)); + tmp.hi = 0; + } else { + tmp = (uint2)(0, 0); + } + } + if (exp_diff > 0) + st_c.mantissa = tmp; + else + st_mul.mantissa = tmp; + + struct fp st_fma; + st_fma.sign = st_mul.sign; + st_fma.exponent = max(st_mul.exponent, st_c.exponent); + st_fma.mantissa = (uint2)(0, 0); + if (st_c.sign == st_mul.sign) { + uint carry = (hadd(st_mul.mantissa.lo, st_c.mantissa.lo) >> 31) & 0x1; + st_fma.mantissa = st_mul.mantissa + st_c.mantissa; + st_fma.mantissa.hi += carry; + } else { + // cutoff bits borrow one + uint cutoff_borrow = ((cutoff_bits.lo != 0 || cutoff_bits.hi != 0) && + (st_mul.exponent > st_c.exponent)) + ? 1 + : 0; + uint borrow = 0; + if (st_c.mantissa.lo > st_mul.mantissa.lo) { + borrow = 1; + } else if (st_c.mantissa.lo == UINT_MAX && cutoff_borrow == 1) { + borrow = 1; + } else if ((st_c.mantissa.lo + cutoff_borrow) > st_mul.mantissa.lo) { + borrow = 1; + } + + st_fma.mantissa.lo = st_mul.mantissa.lo - st... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/87187 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits