-
Notifications
You must be signed in to change notification settings - Fork 195
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
ptx: Add cp.async.mbarrier.arrive{.noinc}
- Loading branch information
1 parent
5ce5d28
commit 391703c
Showing
4 changed files
with
72 additions
and
0 deletions.
There are no files selected for viewing
10 changes: 10 additions & 0 deletions
10
docs/libcudacxx/ptx/instructions/cp_async_mbarrier_arrive.rst
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
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,10 @@ | ||
.. _libcudacxx-ptx-instructions-cp-async-mbarrier-arrive: | ||
|
||
cp.async.mbarrier.arrive | ||
======================== | ||
|
||
- PTX ISA: | ||
`cp.async.mbarrier.arrive <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-cp-async-mbarrier-arrive>`__ | ||
|
||
.. include:: generated/cp_async_mbarrier_arrive.rst | ||
.. include:: generated/cp_async_mbarrier_arrive_noinc.rst |
38 changes: 38 additions & 0 deletions
38
libcudacxx/include/cuda/__ptx/instructions/cp_async_mbarrier_arrive.h
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
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,38 @@ | ||
// -*- C++ -*- | ||
//===----------------------------------------------------------------------===// | ||
// | ||
// Part of libcu++, the C++ Standard Library for your entire system, | ||
// under the Apache License v2.0 with LLVM Exceptions. | ||
// See https://llvm.org/LICENSE.txt for license information. | ||
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | ||
// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. | ||
// | ||
//===----------------------------------------------------------------------===// | ||
|
||
#ifndef _CUDA_PTX_CP_ASYNC_MBARRIER_ARRIVE_H_ | ||
#define _CUDA_PTX_CP_ASYNC_MBARRIER_ARRIVE_H_ | ||
|
||
#include <cuda/std/detail/__config> | ||
|
||
#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) | ||
# pragma GCC system_header | ||
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) | ||
# pragma clang system_header | ||
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) | ||
# pragma system_header | ||
#endif // no system header | ||
|
||
#include <cuda/__ptx/ptx_dot_variants.h> | ||
#include <cuda/__ptx/ptx_helper_functions.h> | ||
#include <cuda/std/cstdint> | ||
|
||
#include <nv/target> // __CUDA_MINIMUM_ARCH__ and friends | ||
|
||
_LIBCUDACXX_BEGIN_NAMESPACE_CUDA_PTX | ||
|
||
#include <cuda/__ptx/instructions/generated/cp_async_mbarrier_arrive.h> | ||
#include <cuda/__ptx/instructions/generated/cp_async_mbarrier_arrive_noinc.h> | ||
|
||
_LIBCUDACXX_END_NAMESPACE_CUDA_PTX | ||
|
||
#endif // _CUDA_PTX_CP_ASYNC_MBARRIER_ARRIVE_H_ |
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
23 changes: 23 additions & 0 deletions
23
libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.mbarrier.arrive.compile.pass.cpp
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
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,23 @@ | ||
//===----------------------------------------------------------------------===// | ||
// | ||
// Part of libcu++, the C++ Standard Library for your entire system, | ||
// under the Apache License v2.0 with LLVM Exceptions. | ||
// See https://llvm.org/LICENSE.txt for license information. | ||
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | ||
// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. | ||
// | ||
//===----------------------------------------------------------------------===// | ||
// UNSUPPORTED: libcpp-has-no-threads | ||
|
||
// <cuda/ptx> | ||
|
||
#include <cuda/ptx> | ||
#include <cuda/std/utility> | ||
|
||
#include "generated/cp_async_mbarrier_arrive.h" | ||
#include "generated/cp_async_mbarrier_arrive_noinc.h" | ||
|
||
int main(int, char**) | ||
{ | ||
return 0; | ||
} |