Skip to content

Commit 7e6dae3

Browse files
ronliebronlieb
committed
merge main into amd-staging
2 parents 9f4fb8d + 5430d7a commit 7e6dae3

File tree

84 files changed

+6765
-4001
lines changed

Some content is hidden

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

84 files changed

+6765
-4001
lines changed

‎clang/unittests/Format/FormatTestComments.cpp‎

Lines changed: 3210 additions & 3423 deletions
Large diffs are not rendered by default.

‎compiler-rt/test/tsan/Darwin/external.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -70,7 +70,7 @@ int main(int argc, char *argv[]) {
7070
// TEST3: WARNING: ThreadSanitizer: race on MyLibrary::MyObject
7171
// TEST3: {{Modifying|Read-only}} access of MyLibrary::MyObject at
7272
// TEST3: {{ObjectWrite|ObjectRead}}
73-
// TEST3: Previous {{modifying|Read-only}} access of MyLibrary::MyObject at
73+
// TEST3: Previous {{modifying|read-only}} access of MyLibrary::MyObject at
7474
// TEST3: {{ObjectWrite|ObjectRead}}
7575
// TEST3: Location is MyLibrary::MyObject of size 16 at
7676
// TEST3: {{ObjectCreate}}

‎flang-rt/CMakeLists.txt‎

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -331,3 +331,19 @@ if (FLANG_RT_INCLUDE_TESTS)
331331
else ()
332332
add_custom_target(check-flang-rt)
333333
endif()
334+
335+
###################
336+
# Install headers #
337+
###################
338+
339+
if (NOT LLVM_INSTALL_TOOLCHAIN_ONLY)
340+
add_llvm_install_targets(install-flang-rt-headers COMPONENT flang-rt-headers)
341+
342+
install(DIRECTORY include/flang-rt/runtime
343+
DESTINATION "${CMAKE_INSTALL_INCLUDEDIR}/flang-rt"
344+
COMPONENT flang-rt-headers
345+
FILES_MATCHING
346+
PATTERN "*.h"
347+
PATTERN ".git" EXCLUDE
348+
PATTERN "CMakeFiles" EXCLUDE)
349+
endif()

‎flang/lib/Optimizer/Builder/IntrinsicCall.cpp‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -9362,6 +9362,8 @@ static void genTMABulkLoad(fir::FirOpBuilder &builder, mlir::Location loc,
93629362
mlir::Value size = mlir::arith::MulIOp::create(builder, loc, nelem, eleSize);
93639363
auto llvmPtrTy = mlir::LLVM::LLVMPointerType::get(builder.getContext());
93649364
barrier = builder.createConvert(loc, llvmPtrTy, barrier);
9365+
dst = builder.createConvert(loc, llvmPtrTy, dst);
9366+
src = builder.createConvert(loc, llvmPtrTy, src);
93659367
mlir::NVVM::InlinePtxOp::create(
93669368
builder, loc, mlir::TypeRange{}, {dst, src, size, barrier}, {},
93679369
"cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], "

‎flang/test/Lower/CUDA/cuda-device-proc.cuf‎

Lines changed: 19 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -468,7 +468,18 @@ attributes(global) subroutine test_bulk_g2s(a)
468468
end subroutine
469469

470470
! CHECK-LABEL: func.func @_QPtest_bulk_g2s
471-
! CHECK: nvvm.cp.async.bulk.shared.cluster.global %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : <7>, <1>
471+
! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %4 {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_bulk_g2sEbarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>)
472+
! CHECK: %[[DST:.*]]:2 = hlfir.declare %16(%17) {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_bulk_g2sEtmpa"} : (!fir.ref<!fir.array<1024xf64>>, !fir.shape<1>) -> (!fir.ref<!fir.array<1024xf64>>, !fir.ref<!fir.array<1024xf64>>)
473+
! CHECK: %[[COUNT:.*]]:2 = hlfir.declare %19 {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_bulk_g2sEtx_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>)
474+
! CHECK: %[[SRC:.*]] = hlfir.designate %{{.*}} (%{{.*}}) : (!fir.box<!fir.array<?xf64>>, i64) -> !fir.ref<f64>
475+
! CHECK: %[[COUNT_LOAD:.*]] = fir.load %20#0 : !fir.ref<i32>
476+
! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr
477+
! CHECK: %[[BARRIER_3:.*]] = llvm.addrspacecast %[[BARRIER_PTR]] : !llvm.ptr to !llvm.ptr<3>
478+
! CHECK: %[[DST_PTR:.*]] = fir.convert %[[DST]]#0 : (!fir.ref<!fir.array<1024xf64>>) -> !llvm.ptr
479+
! CHECK: %[[DST_7:.*]] = llvm.addrspacecast %[[DST_PTR]] : !llvm.ptr to !llvm.ptr<7>
480+
! CHECK: %[[SRC_PTR:.*]] = fir.convert %[[SRC]] : (!fir.ref<f64>) -> !llvm.ptr
481+
! CHECK: %[[SRC_3:.*]] = llvm.addrspacecast %[[SRC_PTR]] : !llvm.ptr to !llvm.ptr<1>
482+
! CHECK: nvvm.cp.async.bulk.shared.cluster.global %[[DST_7]], %[[SRC_3]], %[[BARRIER_3]], %[[COUNT_LOAD]] : <7>, <1>
472483

473484
attributes(global) subroutine test_bulk_s2g(a)
474485
real(8), device :: a(*)
@@ -533,7 +544,7 @@ end subroutine
533544
! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 8 : i32
534545
! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32
535546
! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr
536-
! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !fir.ref<!fir.array<1024xcomplex<f32>>>, !fir.ref<complex<f32>>, i32, !llvm.ptr)
547+
! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr)
537548
! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32)
538549

539550
attributes(global) subroutine test_tma_bulk_load_c8(a, n)
@@ -552,7 +563,7 @@ end subroutine
552563
! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 16 : i32
553564
! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32
554565
! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr
555-
! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !fir.ref<!fir.array<1024xcomplex<f64>>>, !fir.ref<complex<f64>>, i32, !llvm.ptr)
566+
! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr)
556567
! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32)
557568

558569
attributes(global) subroutine test_tma_bulk_load_i4(a, n)
@@ -571,7 +582,7 @@ end subroutine
571582
! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 4 : i32
572583
! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32
573584
! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr
574-
! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !fir.ref<!fir.array<1024xi32>>, !fir.ref<i32>, i32, !llvm.ptr)
585+
! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr)
575586
! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32)
576587

577588
attributes(global) subroutine test_tma_bulk_load_i8(a, n)
@@ -590,7 +601,7 @@ end subroutine
590601
! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 8 : i32
591602
! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32
592603
! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr
593-
! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !fir.ref<!fir.array<1024xi64>>, !fir.ref<i64>, i32, !llvm.ptr)
604+
! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr)
594605
! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32)
595606

596607
attributes(global) subroutine test_tma_bulk_load_r2(a, n)
@@ -609,7 +620,7 @@ end subroutine
609620
! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 2 : i32
610621
! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32
611622
! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr
612-
! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !fir.ref<!fir.array<1024xf16>>, !fir.ref<f16>, i32, !llvm.ptr)
623+
! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr)
613624
! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32)
614625

615626
attributes(global) subroutine test_tma_bulk_load_r4(a, n)
@@ -628,7 +639,7 @@ end subroutine
628639
! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 4 : i32
629640
! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32
630641
! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr
631-
! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !fir.ref<!fir.array<1024xf32>>, !fir.ref<f32>, i32, !llvm.ptr)
642+
! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr)
632643
! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32)
633644

634645
attributes(global) subroutine test_tma_bulk_load_r8(a, n)
@@ -647,7 +658,7 @@ end subroutine
647658
! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 8 : i32
648659
! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32
649660
! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr
650-
! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !fir.ref<!fir.array<1024xf64>>, !fir.ref<f64>, i32, !llvm.ptr)
661+
! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr)
651662
! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32)
652663

653664
attributes(global) subroutine test_tma_bulk_store_c4(c, n)

‎libcxx/include/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -529,6 +529,7 @@ set(files
529529
__locale_dir/support/freebsd.h
530530
__locale_dir/support/fuchsia.h
531531
__locale_dir/support/linux.h
532+
__locale_dir/support/netbsd.h
532533
__locale_dir/support/no_locale/characters.h
533534
__locale_dir/support/no_locale/strtonum.h
534535
__locale_dir/support/windows.h

‎libcxx/include/__locale_dir/locale_base_api.h‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -115,6 +115,8 @@
115115
# include <__locale_dir/support/apple.h>
116116
# elif defined(__FreeBSD__)
117117
# include <__locale_dir/support/freebsd.h>
118+
# elif defined(__NetBSD__)
119+
# include <__locale_dir/support/netbsd.h>
118120
# elif defined(_LIBCPP_MSVCRT_LIKE)
119121
# include <__locale_dir/support/windows.h>
120122
# elif defined(__Fuchsia__)

‎libcxx/include/__locale_dir/support/bsd_like.h‎

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -24,7 +24,9 @@
2424
# include <wctype.h>
2525
#endif
2626

27-
#include <xlocale.h>
27+
#if __has_include(<xlocale.h>)
28+
# include <xlocale.h>
29+
#endif
2830

2931
#if !defined(_LIBCPP_HAS_NO_PRAGMA_SYSTEM_HEADER)
3032
# pragma GCC system_header

‎libcxx/include/__locale_dir/support/netbsd.h‎

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,20 @@
1+
//===-----------------------------------------------------------------------===//
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 _LIBCPP___LOCALE_DIR_SUPPORT_NETBSD_H
10+
#define _LIBCPP___LOCALE_DIR_SUPPORT_NETBSD_H
11+
12+
#include <__config>
13+
14+
#if !defined(_LIBCPP_HAS_NO_PRAGMA_SYSTEM_HEADER)
15+
# pragma GCC system_header
16+
#endif
17+
18+
#include <__locale_dir/support/bsd_like.h>
19+
20+
#endif // _LIBCPP___LOCALE_DIR_SUPPORT_NETBSD_H

‎libcxx/include/module.modulemap.in‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1587,6 +1587,7 @@ module std [system] {
15871587
textual header "__locale_dir/support/freebsd.h"
15881588
textual header "__locale_dir/support/fuchsia.h"
15891589
textual header "__locale_dir/support/linux.h"
1590+
textual header "__locale_dir/support/netbsd.h"
15901591
textual header "__locale_dir/support/no_locale/characters.h"
15911592
textual header "__locale_dir/support/no_locale/strtonum.h"
15921593
textual header "__locale_dir/support/windows.h"

0 commit comments

Comments
 (0)