Skip to content

Commit e2cc94e

Browse files
author
z1_cciauto
authored
merge main into amd-staging (llvm#3400)
2 parents 214a7e7 + 7c59f8e commit e2cc94e

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

47 files changed

+1960
-631
lines changed

‎clang/include/clang/Basic/BuiltinsAMDGPU.def

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -726,6 +726,12 @@ TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_fp8_f32_e5m3, "ifiiIi", "nc", "fp8e5m3-in
726726
TARGET_BUILTIN(__builtin_amdgcn_sat_pk4_i4_i8, "UsUi", "nc", "gfx1250-insts")
727727
TARGET_BUILTIN(__builtin_amdgcn_sat_pk4_u4_u8, "UsUi", "nc", "gfx1250-insts")
728728

729+
TARGET_BUILTIN(__builtin_amdgcn_permlane_bcast, "iiii", "nc", "gfx1250-insts,wavefrontsize32")
730+
TARGET_BUILTIN(__builtin_amdgcn_permlane_up, "iiii", "nc", "gfx1250-insts,wavefrontsize32")
731+
TARGET_BUILTIN(__builtin_amdgcn_permlane_down, "iiii", "nc", "gfx1250-insts,wavefrontsize32")
732+
TARGET_BUILTIN(__builtin_amdgcn_permlane_xor, "iiii", "nc", "gfx1250-insts,wavefrontsize32")
733+
TARGET_BUILTIN(__builtin_amdgcn_permlane_idx_gen, "iii", "nc", "gfx1250-insts,wavefrontsize32")
734+
729735
// GFX1250 WMMA builtins
730736
TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x4_f32, "V8fIbV2fIbV2fIsV8fIbIb", "nc", "gfx1250-insts,wavefrontsize32")
731737
TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x32_bf16, "V8fIbV16yIbV16yIsV8fIbIb", "nc", "gfx1250-insts,wavefrontsize32")

‎clang/lib/CodeGen/CGBuiltin.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -973,6 +973,9 @@ class StructFieldAccess
973973
AddrOfSeen = false;
974974
return Visit(E->getSubExpr());
975975
}
976+
const Expr *VisitBinaryOperator(const clang::BinaryOperator *Op) {
977+
return Op->isCommaOp() ? Visit(Op->getRHS()) : nullptr;
978+
}
976979
};
977980

978981
} // end anonymous namespace

‎clang/lib/Headers/__clang_cuda_runtime_wrapper.h

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -385,7 +385,12 @@ __host__ __device__ void __nv_tex_surf_handler(const char *name, T *ptr,
385385
#endif // CUDA_VERSION
386386
#endif // __cplusplus >= 201103L && CUDA_VERSION >= 9000
387387
#include "surface_indirect_functions.h"
388+
#if CUDA_VERSION < 13000
389+
// Direct texture fetch functions had been deprecated since CUDA-11.
390+
// The file in CUDA-12 only carried unused texture types, and is no longer
391+
// needed.
388392
#include "texture_fetch_functions.h"
393+
#endif // CUDA_VERSION < 13000
389394
#include "texture_indirect_functions.h"
390395

391396
// Restore state of __CUDA_ARCH__ and __THROW we had on entry.

‎clang/test/CodeGen/attr-counted-by.c

Lines changed: 36 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2445,3 +2445,39 @@ struct {
24452445
size_t test36() {
24462446
return __builtin_dynamic_object_size(&x.dev_addr[4], 1);
24472447
}
2448+
2449+
// SANITIZE-WITH-ATTR-LABEL: define dso_local range(i64 -8589934592, 8589934589) i64 @test37(
2450+
// SANITIZE-WITH-ATTR-SAME: ptr noundef [[PTR:%.*]]) local_unnamed_addr #[[ATTR0]] {
2451+
// SANITIZE-WITH-ATTR-NEXT: entry:
2452+
// SANITIZE-WITH-ATTR-NEXT: [[COUNTED_BY_GEP:%.*]] = getelementptr inbounds nuw i8, ptr [[PTR]], i64 8
2453+
// SANITIZE-WITH-ATTR-NEXT: [[COUNTED_BY_LOAD:%.*]] = load i32, ptr [[COUNTED_BY_GEP]], align 4
2454+
// SANITIZE-WITH-ATTR-NEXT: [[COUNT:%.*]] = sext i32 [[COUNTED_BY_LOAD]] to i64
2455+
// SANITIZE-WITH-ATTR-NEXT: [[FLEXIBLE_ARRAY_MEMBER_SIZE:%.*]] = shl nsw i64 [[COUNT]], 2
2456+
// SANITIZE-WITH-ATTR-NEXT: [[TMP0:%.*]] = icmp sgt i32 [[COUNTED_BY_LOAD]], -1
2457+
// SANITIZE-WITH-ATTR-NEXT: [[TMP1:%.*]] = select i1 [[TMP0]], i64 [[FLEXIBLE_ARRAY_MEMBER_SIZE]], i64 0
2458+
// SANITIZE-WITH-ATTR-NEXT: ret i64 [[TMP1]]
2459+
//
2460+
// NO-SANITIZE-WITH-ATTR-LABEL: define dso_local range(i64 -8589934592, 8589934589) i64 @test37(
2461+
// NO-SANITIZE-WITH-ATTR-SAME: ptr noundef readonly captures(none) [[PTR:%.*]]) local_unnamed_addr #[[ATTR2]] {
2462+
// NO-SANITIZE-WITH-ATTR-NEXT: entry:
2463+
// NO-SANITIZE-WITH-ATTR-NEXT: [[COUNTED_BY_GEP:%.*]] = getelementptr inbounds nuw i8, ptr [[PTR]], i64 8
2464+
// NO-SANITIZE-WITH-ATTR-NEXT: [[COUNTED_BY_LOAD:%.*]] = load i32, ptr [[COUNTED_BY_GEP]], align 4
2465+
// NO-SANITIZE-WITH-ATTR-NEXT: [[COUNT:%.*]] = sext i32 [[COUNTED_BY_LOAD]] to i64
2466+
// NO-SANITIZE-WITH-ATTR-NEXT: [[FLEXIBLE_ARRAY_MEMBER_SIZE:%.*]] = shl nsw i64 [[COUNT]], 2
2467+
// NO-SANITIZE-WITH-ATTR-NEXT: [[TMP0:%.*]] = icmp sgt i32 [[COUNTED_BY_LOAD]], -1
2468+
// NO-SANITIZE-WITH-ATTR-NEXT: [[TMP1:%.*]] = select i1 [[TMP0]], i64 [[FLEXIBLE_ARRAY_MEMBER_SIZE]], i64 0
2469+
// NO-SANITIZE-WITH-ATTR-NEXT: ret i64 [[TMP1]]
2470+
//
2471+
// SANITIZE-WITHOUT-ATTR-LABEL: define dso_local i64 @test37(
2472+
// SANITIZE-WITHOUT-ATTR-SAME: ptr noundef [[PTR:%.*]]) local_unnamed_addr #[[ATTR0]] {
2473+
// SANITIZE-WITHOUT-ATTR-NEXT: entry:
2474+
// SANITIZE-WITHOUT-ATTR-NEXT: ret i64 -1
2475+
//
2476+
// NO-SANITIZE-WITHOUT-ATTR-LABEL: define dso_local i64 @test37(
2477+
// NO-SANITIZE-WITHOUT-ATTR-SAME: ptr noundef readnone [[PTR:%.*]]) local_unnamed_addr #[[ATTR1]] {
2478+
// NO-SANITIZE-WITHOUT-ATTR-NEXT: entry:
2479+
// NO-SANITIZE-WITHOUT-ATTR-NEXT: ret i64 -1
2480+
//
2481+
size_t test37(struct annotated *ptr) {
2482+
return __builtin_dynamic_object_size((1, 2, (4, 5, (7, 8, 9, (10, ptr->array)))), 1);
2483+
}

‎clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl

Lines changed: 126 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -744,6 +744,132 @@ void test_permlane16_swap(global uint2* out, uint old, uint src) {
744744
*out = __builtin_amdgcn_permlane16_swap(old, src, false, true);
745745
}
746746

747+
// CHECK-LABEL: @test_permlane_bcast(
748+
// CHECK-NEXT: entry:
749+
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
750+
// CHECK-NEXT: [[SRC0_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
751+
// CHECK-NEXT: [[SRC1_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
752+
// CHECK-NEXT: [[SRC2_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
753+
// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
754+
// CHECK-NEXT: [[SRC0_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC0_ADDR]] to ptr
755+
// CHECK-NEXT: [[SRC1_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC1_ADDR]] to ptr
756+
// CHECK-NEXT: [[SRC2_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC2_ADDR]] to ptr
757+
// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]], align 8
758+
// CHECK-NEXT: store i32 [[SRC0:%.*]], ptr [[SRC0_ADDR_ASCAST]], align 4
759+
// CHECK-NEXT: store i32 [[SRC1:%.*]], ptr [[SRC1_ADDR_ASCAST]], align 4
760+
// CHECK-NEXT: store i32 [[SRC2:%.*]], ptr [[SRC2_ADDR_ASCAST]], align 4
761+
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[SRC0_ADDR_ASCAST]], align 4
762+
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[SRC1_ADDR_ASCAST]], align 4
763+
// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[SRC2_ADDR_ASCAST]], align 4
764+
// CHECK-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.permlane.bcast(i32 [[TMP0]], i32 [[TMP1]], i32 [[TMP2]])
765+
// CHECK-NEXT: [[TMP4:%.*]] = load ptr addrspace(1), ptr [[OUT_ADDR_ASCAST]], align 8
766+
// CHECK-NEXT: store i32 [[TMP3]], ptr addrspace(1) [[TMP4]], align 4
767+
// CHECK-NEXT: ret void
768+
//
769+
void test_permlane_bcast(global uint* out, uint src0, uint src1, uint src2) {
770+
*out = __builtin_amdgcn_permlane_bcast(src0, src1, src2);
771+
}
772+
773+
// CHECK-LABEL: @test_permlane_down(
774+
// CHECK-NEXT: entry:
775+
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
776+
// CHECK-NEXT: [[SRC0_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
777+
// CHECK-NEXT: [[SRC1_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
778+
// CHECK-NEXT: [[SRC2_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
779+
// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
780+
// CHECK-NEXT: [[SRC0_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC0_ADDR]] to ptr
781+
// CHECK-NEXT: [[SRC1_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC1_ADDR]] to ptr
782+
// CHECK-NEXT: [[SRC2_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC2_ADDR]] to ptr
783+
// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]], align 8
784+
// CHECK-NEXT: store i32 [[SRC0:%.*]], ptr [[SRC0_ADDR_ASCAST]], align 4
785+
// CHECK-NEXT: store i32 [[SRC1:%.*]], ptr [[SRC1_ADDR_ASCAST]], align 4
786+
// CHECK-NEXT: store i32 [[SRC2:%.*]], ptr [[SRC2_ADDR_ASCAST]], align 4
787+
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[SRC0_ADDR_ASCAST]], align 4
788+
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[SRC1_ADDR_ASCAST]], align 4
789+
// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[SRC2_ADDR_ASCAST]], align 4
790+
// CHECK-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.permlane.down(i32 [[TMP0]], i32 [[TMP1]], i32 [[TMP2]])
791+
// CHECK-NEXT: [[TMP4:%.*]] = load ptr addrspace(1), ptr [[OUT_ADDR_ASCAST]], align 8
792+
// CHECK-NEXT: store i32 [[TMP3]], ptr addrspace(1) [[TMP4]], align 4
793+
// CHECK-NEXT: ret void
794+
//
795+
void test_permlane_down(global uint* out, uint src0, uint src1, uint src2) {
796+
*out = __builtin_amdgcn_permlane_down(src0, src1, src2);
797+
}
798+
799+
// CHECK-LABEL: @test_permlane_up(
800+
// CHECK-NEXT: entry:
801+
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
802+
// CHECK-NEXT: [[SRC0_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
803+
// CHECK-NEXT: [[SRC1_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
804+
// CHECK-NEXT: [[SRC2_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
805+
// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
806+
// CHECK-NEXT: [[SRC0_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC0_ADDR]] to ptr
807+
// CHECK-NEXT: [[SRC1_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC1_ADDR]] to ptr
808+
// CHECK-NEXT: [[SRC2_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC2_ADDR]] to ptr
809+
// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]], align 8
810+
// CHECK-NEXT: store i32 [[SRC0:%.*]], ptr [[SRC0_ADDR_ASCAST]], align 4
811+
// CHECK-NEXT: store i32 [[SRC1:%.*]], ptr [[SRC1_ADDR_ASCAST]], align 4
812+
// CHECK-NEXT: store i32 [[SRC2:%.*]], ptr [[SRC2_ADDR_ASCAST]], align 4
813+
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[SRC0_ADDR_ASCAST]], align 4
814+
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[SRC1_ADDR_ASCAST]], align 4
815+
// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[SRC2_ADDR_ASCAST]], align 4
816+
// CHECK-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.permlane.up(i32 [[TMP0]], i32 [[TMP1]], i32 [[TMP2]])
817+
// CHECK-NEXT: [[TMP4:%.*]] = load ptr addrspace(1), ptr [[OUT_ADDR_ASCAST]], align 8
818+
// CHECK-NEXT: store i32 [[TMP3]], ptr addrspace(1) [[TMP4]], align 4
819+
// CHECK-NEXT: ret void
820+
//
821+
void test_permlane_up(global uint* out, uint src0, uint src1, uint src2) {
822+
*out = __builtin_amdgcn_permlane_up(src0, src1, src2);
823+
}
824+
825+
// CHECK-LABEL: @test_permlane_xor(
826+
// CHECK-NEXT: entry:
827+
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
828+
// CHECK-NEXT: [[SRC0_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
829+
// CHECK-NEXT: [[SRC1_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
830+
// CHECK-NEXT: [[SRC2_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
831+
// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
832+
// CHECK-NEXT: [[SRC0_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC0_ADDR]] to ptr
833+
// CHECK-NEXT: [[SRC1_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC1_ADDR]] to ptr
834+
// CHECK-NEXT: [[SRC2_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC2_ADDR]] to ptr
835+
// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]], align 8
836+
// CHECK-NEXT: store i32 [[SRC0:%.*]], ptr [[SRC0_ADDR_ASCAST]], align 4
837+
// CHECK-NEXT: store i32 [[SRC1:%.*]], ptr [[SRC1_ADDR_ASCAST]], align 4
838+
// CHECK-NEXT: store i32 [[SRC2:%.*]], ptr [[SRC2_ADDR_ASCAST]], align 4
839+
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[SRC0_ADDR_ASCAST]], align 4
840+
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[SRC1_ADDR_ASCAST]], align 4
841+
// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[SRC2_ADDR_ASCAST]], align 4
842+
// CHECK-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.permlane.xor(i32 [[TMP0]], i32 [[TMP1]], i32 [[TMP2]])
843+
// CHECK-NEXT: [[TMP4:%.*]] = load ptr addrspace(1), ptr [[OUT_ADDR_ASCAST]], align 8
844+
// CHECK-NEXT: store i32 [[TMP3]], ptr addrspace(1) [[TMP4]], align 4
845+
// CHECK-NEXT: ret void
846+
//
847+
void test_permlane_xor(global uint* out, uint src0, uint src1, uint src2) {
848+
*out = __builtin_amdgcn_permlane_xor(src0, src1, src2);
849+
}
850+
851+
// CHECK-LABEL: @test_permlane_idx_gen(
852+
// CHECK-NEXT: entry:
853+
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
854+
// CHECK-NEXT: [[SRC0_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
855+
// CHECK-NEXT: [[SRC1_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
856+
// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
857+
// CHECK-NEXT: [[SRC0_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC0_ADDR]] to ptr
858+
// CHECK-NEXT: [[SRC1_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC1_ADDR]] to ptr
859+
// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]], align 8
860+
// CHECK-NEXT: store i32 [[SRC0:%.*]], ptr [[SRC0_ADDR_ASCAST]], align 4
861+
// CHECK-NEXT: store i32 [[SRC1:%.*]], ptr [[SRC1_ADDR_ASCAST]], align 4
862+
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[SRC0_ADDR_ASCAST]], align 4
863+
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[SRC1_ADDR_ASCAST]], align 4
864+
// CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.permlane.idx.gen(i32 [[TMP0]], i32 [[TMP1]])
865+
// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(1), ptr [[OUT_ADDR_ASCAST]], align 8
866+
// CHECK-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[TMP3]], align 4
867+
// CHECK-NEXT: ret void
868+
//
869+
void test_permlane_idx_gen(global uint* out, uint src0, uint src1) {
870+
*out = __builtin_amdgcn_permlane_idx_gen(src0, src1);
871+
}
872+
747873
// CHECK-LABEL: @test_prefetch(
748874
// CHECK-NEXT: entry:
749875
// CHECK-NEXT: [[FPTR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)

‎libc/shared/math.h

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -25,8 +25,10 @@
2525
#include "math/atan.h"
2626
#include "math/atan2.h"
2727
#include "math/atan2f.h"
28+
#include "math/atan2f128.h"
2829
#include "math/atanf.h"
2930
#include "math/atanf16.h"
31+
#include "math/atanhf.h"
3032
#include "math/erff.h"
3133
#include "math/exp.h"
3234
#include "math/exp10.h"

‎libc/shared/math/atan2f128.h

Lines changed: 29 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,29 @@
1+
//===-- Shared atan2f128 function -------------------------------*- C++ -*-===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#ifndef LLVM_LIBC_SHARED_MATH_ATAN2F128_H
10+
#define LLVM_LIBC_SHARED_MATH_ATAN2F128_H
11+
12+
#include "include/llvm-libc-types/float128.h"
13+
14+
#ifdef LIBC_TYPES_HAS_FLOAT128
15+
16+
#include "shared/libc_common.h"
17+
#include "src/__support/math/atan2f128.h"
18+
19+
namespace LIBC_NAMESPACE_DECL {
20+
namespace shared {
21+
22+
using math::atan2f128;
23+
24+
} // namespace shared
25+
} // namespace LIBC_NAMESPACE_DECL
26+
27+
#endif // LIBC_TYPES_HAS_FLOAT128
28+
29+
#endif // LLVM_LIBC_SHARED_MATH_ATAN2F128_H

‎libc/shared/math/atanhf.h

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,23 @@
1+
//===-- Shared atanhf function ----------------------------------*- C++ -*-===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#ifndef LLVM_LIBC_SHARED_MATH_ATANHF_H
10+
#define LLVM_LIBC_SHARED_MATH_ATANHF_H
11+
12+
#include "shared/libc_common.h"
13+
#include "src/__support/math/atanhf.h"
14+
15+
namespace LIBC_NAMESPACE_DECL {
16+
namespace shared {
17+
18+
using math::atanhf;
19+
20+
} // namespace shared
21+
} // namespace LIBC_NAMESPACE_DECL
22+
23+
#endif // LLVM_LIBC_SHARED_MATH_ATANHF_H

‎libc/src/__support/math/CMakeLists.txt

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -230,6 +230,21 @@ add_header_library(
230230
libc.src.__support.macros.optimization
231231
)
232232

233+
add_header_library(
234+
atan2f128
235+
HDRS
236+
atan2f128.h
237+
DEPENDS
238+
.atan_utils
239+
libc.src.__support.integer_literals
240+
libc.src.__support.uint128
241+
libc.src.__support.FPUtil.dyadic_float
242+
libc.src.__support.FPUtil.fp_bits
243+
libc.src.__support.FPUtil.multiply_add
244+
libc.src.__support.FPUtil.nearest_integer
245+
libc.src.__support.macros.optimization
246+
)
247+
233248
add_header_library(
234249
atanf
235250
HDRS
@@ -260,6 +275,17 @@ add_header_library(
260275
libc.src.__support.macros.optimization
261276
)
262277

278+
add_header_library(
279+
atanhf
280+
HDRS
281+
atanhf.h
282+
DEPENDS
283+
.acoshf_utils
284+
libc.src.__support.FPUtil.fp_bits
285+
libc.src.__support.FPUtil.fenv_impl
286+
libc.src.__support.macros.optimization
287+
)
288+
263289
add_header_library(
264290
asinf
265291
HDRS

0 commit comments

Comments
 (0)