llvm encode decode #87187
Closed
llvm encode decode #87187
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
|
Thank you for submitting a Pull Request (PR) to the LLVM Project! This PR will be automatically labeled and the relevant teams will be If you wish to, you can add reviewers by using the "Reviewers" section on this page. If this is not working for you, it is probably because you do not have write If you have received no comments on your PR for a week, you can request a review If you have further questions, they may be answered by the LLVM GitHub User Guide. You can also ask questions in a comment on this PR, on the LLVM Discord or on the forums. |
llvmbot
added
clang
|
@llvm/pr-subscribers-libcxx @llvm/pr-subscribers-lld-elf Author: Jiang zixian (jiang-zixian) Changesthe 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:
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]
|
|
@llvm/pr-subscribers-lld Author: Jiang zixian (jiang-zixian) Changesthe 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:
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]
|
|
@llvm/pr-subscribers-clang-driver Author: Jiang zixian (jiang-zixian) Changesthe 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:
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]
|
|
@llvm/pr-subscribers-libcxx Author: Jiang zixian (jiang-zixian) Changesthe 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:
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]
|
ldionne
removed
the
libc++
llvmbot
added
the
libc++
|
Please send this PR to your downstream fork https://github.com/x-codingman/llvm-project. |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
the decoding and encoding are complete