Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

ptx: Test alternative to NV_IF_TARGET [Do not merge] #3158

Draft
wants to merge 2 commits into
base: main
Choose a base branch
from
Draft

ptx: Test alternative to NV_IF_TARGET [Do not merge] #3158

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
1045021
ptx: Test alternative to NV_IF_TARGET
ahendriksen Dec 13, 2024
99da44e
Include enable_if.h
ahendriksen Dec 13, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions libcudacxx/include/cuda/__ptx/instructions/barrier_cluster.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,6 +25,7 @@
#include <cuda/__ptx/ptx_dot_variants.h>
#include <cuda/__ptx/ptx_helper_functions.h>
#include <cuda/std/cstdint>
#include <cuda/std/__type_traits/enable_if.h>

#include <nv/target> // __CUDA_MINIMUM_ARCH__ and friends

Expand Down
1 change: 1 addition & 0 deletions libcudacxx/include/cuda/__ptx/instructions/cp_async_bulk.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,6 +25,7 @@
#include <cuda/__ptx/ptx_dot_variants.h>
#include <cuda/__ptx/ptx_helper_functions.h>
#include <cuda/std/cstdint>
#include <cuda/std/__type_traits/enable_if.h>

#include <nv/target> // __CUDA_MINIMUM_ARCH__ and friends

Expand Down
109 changes: 61 additions & 48 deletions libcudacxx/include/cuda/__ptx/instructions/generated/barrier_cluster.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,15 +14,17 @@ extern "C" _CCCL_DEVICE void __cuda_ptx_barrier_cluster_arrive_is_not_supported_
template <typename = void>
_CCCL_DEVICE static inline void barrier_cluster_arrive()
{
NV_IF_ELSE_TARGET(
NV_PROVIDES_SM_90,
(asm volatile("barrier.cluster.arrive;"
:
:
: "memory");),
(
// Unsupported architectures will have a linker error with a semi-decent error message
__cuda_ptx_barrier_cluster_arrive_is_not_supported_before_SM_90__();));
#if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 900
asm volatile (
"barrier.cluster.arrive;"
:
:
: "memory"
);
#else
// Unsupported architectures will have a linker error with a semi-decent error message
__cuda_ptx_barrier_cluster_arrive_is_not_supported_before_SM_90__();
#endif
}
#endif // __cccl_ptx_isa >= 780

Expand All @@ -37,15 +39,17 @@ extern "C" _CCCL_DEVICE void __cuda_ptx_barrier_cluster_wait_is_not_supported_be
template <typename = void>
_CCCL_DEVICE static inline void barrier_cluster_wait()
{
NV_IF_ELSE_TARGET(
NV_PROVIDES_SM_90,
(asm volatile("barrier.cluster.wait;"
:
:
: "memory");),
(
// Unsupported architectures will have a linker error with a semi-decent error message
__cuda_ptx_barrier_cluster_wait_is_not_supported_before_SM_90__();));
#if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 900
asm volatile (
"barrier.cluster.wait;"
:
:
: "memory"
);
#else
// Unsupported architectures will have a linker error with a semi-decent error message
__cuda_ptx_barrier_cluster_wait_is_not_supported_before_SM_90__();
#endif
}
#endif // __cccl_ptx_isa >= 780

Expand All @@ -60,18 +64,21 @@ __device__ static inline void barrier_cluster_arrive(
#if __cccl_ptx_isa >= 800
extern "C" _CCCL_DEVICE void __cuda_ptx_barrier_cluster_arrive_is_not_supported_before_SM_90__();
template <typename = void>
_CCCL_DEVICE static inline void barrier_cluster_arrive(sem_release_t)
_CCCL_DEVICE static inline void barrier_cluster_arrive(
sem_release_t)
{
// __sem == sem_release (due to parameter type constraint)
NV_IF_ELSE_TARGET(
NV_PROVIDES_SM_90,
(asm volatile("barrier.cluster.arrive.release;"
:
:
: "memory");),
(
// Unsupported architectures will have a linker error with a semi-decent error message
__cuda_ptx_barrier_cluster_arrive_is_not_supported_before_SM_90__();));
#if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 900
asm volatile (
"barrier.cluster.arrive.release;"
:
:
: "memory"
);
#else
// Unsupported architectures will have a linker error with a semi-decent error message
__cuda_ptx_barrier_cluster_arrive_is_not_supported_before_SM_90__();
#endif
}
#endif // __cccl_ptx_isa >= 800

Expand All @@ -86,18 +93,21 @@ __device__ static inline void barrier_cluster_arrive(
#if __cccl_ptx_isa >= 800
extern "C" _CCCL_DEVICE void __cuda_ptx_barrier_cluster_arrive_is_not_supported_before_SM_90__();
template <typename = void>
_CCCL_DEVICE static inline void barrier_cluster_arrive(sem_relaxed_t)
_CCCL_DEVICE static inline void barrier_cluster_arrive(
sem_relaxed_t)
{
// __sem == sem_relaxed (due to parameter type constraint)
NV_IF_ELSE_TARGET(
NV_PROVIDES_SM_90,
(asm volatile("barrier.cluster.arrive.relaxed;"
:
:
:);),
(
// Unsupported architectures will have a linker error with a semi-decent error message
__cuda_ptx_barrier_cluster_arrive_is_not_supported_before_SM_90__();));
#if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 900
asm volatile (
"barrier.cluster.arrive.relaxed;"
:
:
:
);
#else
// Unsupported architectures will have a linker error with a semi-decent error message
__cuda_ptx_barrier_cluster_arrive_is_not_supported_before_SM_90__();
#endif
}
#endif // __cccl_ptx_isa >= 800

Expand All @@ -112,18 +122,21 @@ __device__ static inline void barrier_cluster_wait(
#if __cccl_ptx_isa >= 800
extern "C" _CCCL_DEVICE void __cuda_ptx_barrier_cluster_wait_is_not_supported_before_SM_90__();
template <typename = void>
_CCCL_DEVICE static inline void barrier_cluster_wait(sem_acquire_t)
_CCCL_DEVICE static inline void barrier_cluster_wait(
sem_acquire_t)
{
// __sem == sem_acquire (due to parameter type constraint)
NV_IF_ELSE_TARGET(
NV_PROVIDES_SM_90,
(asm volatile("barrier.cluster.wait.acquire;"
:
:
: "memory");),
(
// Unsupported architectures will have a linker error with a semi-decent error message
__cuda_ptx_barrier_cluster_wait_is_not_supported_before_SM_90__();));
#if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 900
asm volatile (
"barrier.cluster.wait.acquire;"
:
:
: "memory"
);
#else
// Unsupported architectures will have a linker error with a semi-decent error message
__cuda_ptx_barrier_cluster_wait_is_not_supported_before_SM_90__();
#endif
}
#endif // __cccl_ptx_isa >= 800

Expand Down
82 changes: 48 additions & 34 deletions libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,8 +4,7 @@
#define _CUDA_PTX_GENERATED_CP_ASYNC_BULK_H_

/*
// cp.async.bulk.dst.src.mbarrier::complete_tx::bytes [dstMem], [srcMem], size, [smem_bar]; // 1a. unicast PTX ISA 80,
SM_90
// cp.async.bulk.dst.src.mbarrier::complete_tx::bytes [dstMem], [srcMem], size, [smem_bar]; // 1a. unicast PTX ISA 80, SM_90
// .dst = { .shared::cluster }
// .src = { .global }
template <typename = void>
Expand All @@ -30,15 +29,20 @@ _CCCL_DEVICE static inline void cp_async_bulk(
{
// __space == space_cluster (due to parameter type constraint)
// __space == space_global (due to parameter type constraint)
NV_IF_ELSE_TARGET(
NV_PROVIDES_SM_90,
(asm("cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3]; // 1a. unicast"
:
: "r"(__as_ptr_smem(__dstMem)), "l"(__as_ptr_gmem(__srcMem)), "r"(__size), "r"(__as_ptr_smem(__smem_bar))
: "memory");),
(
// Unsupported architectures will have a linker error with a semi-decent error message
__cuda_ptx_cp_async_bulk_is_not_supported_before_SM_90__();));
#if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 900
asm (
"cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3]; // 1a. unicast"
:
: "r"(__as_ptr_smem(__dstMem)),
"l"(__as_ptr_gmem(__srcMem)),
"r"(__size),
"r"(__as_ptr_smem(__smem_bar))
: "memory"
);
#else
// Unsupported architectures will have a linker error with a semi-decent error message
__cuda_ptx_cp_async_bulk_is_not_supported_before_SM_90__();
#endif
}
#endif // __cccl_ptx_isa >= 800

Expand Down Expand Up @@ -68,18 +72,20 @@ _CCCL_DEVICE static inline void cp_async_bulk(
{
// __space == space_cluster (due to parameter type constraint)
// __space == space_shared (due to parameter type constraint)
NV_IF_ELSE_TARGET(
NV_PROVIDES_SM_90,
(asm("cp.async.bulk.shared::cluster.shared::cta.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3]; // 2. "
:
: "r"(__as_ptr_remote_dsmem(__dstMem)),
"r"(__as_ptr_smem(__srcMem)),
"r"(__size),
"r"(__as_ptr_remote_dsmem(__rdsmem_bar))
: "memory");),
(
// Unsupported architectures will have a linker error with a semi-decent error message
__cuda_ptx_cp_async_bulk_is_not_supported_before_SM_90__();));
#if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 900
asm (
"cp.async.bulk.shared::cluster.shared::cta.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3]; // 2. "
:
: "r"(__as_ptr_remote_dsmem(__dstMem)),
"r"(__as_ptr_smem(__srcMem)),
"r"(__size),
"r"(__as_ptr_remote_dsmem(__rdsmem_bar))
: "memory"
);
#else
// Unsupported architectures will have a linker error with a semi-decent error message
__cuda_ptx_cp_async_bulk_is_not_supported_before_SM_90__();
#endif
}
#endif // __cccl_ptx_isa >= 800

Expand All @@ -98,20 +104,28 @@ __device__ static inline void cp_async_bulk(
#if __cccl_ptx_isa >= 800
extern "C" _CCCL_DEVICE void __cuda_ptx_cp_async_bulk_is_not_supported_before_SM_90__();
template <typename = void>
_CCCL_DEVICE static inline void
cp_async_bulk(space_global_t, space_shared_t, void* __dstMem, const void* __srcMem, const _CUDA_VSTD::uint32_t& __size)
_CCCL_DEVICE static inline void cp_async_bulk(
space_global_t,
space_shared_t,
void* __dstMem,
const void* __srcMem,
const _CUDA_VSTD::uint32_t& __size)
{
// __space == space_global (due to parameter type constraint)
// __space == space_shared (due to parameter type constraint)
NV_IF_ELSE_TARGET(
NV_PROVIDES_SM_90,
(asm("cp.async.bulk.global.shared::cta.bulk_group [%0], [%1], %2; // 3. "
:
: "l"(__as_ptr_gmem(__dstMem)), "r"(__as_ptr_smem(__srcMem)), "r"(__size)
: "memory");),
(
// Unsupported architectures will have a linker error with a semi-decent error message
__cuda_ptx_cp_async_bulk_is_not_supported_before_SM_90__();));
#if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 900
asm (
"cp.async.bulk.global.shared::cta.bulk_group [%0], [%1], %2; // 3. "
:
: "l"(__as_ptr_gmem(__dstMem)),
"r"(__as_ptr_smem(__srcMem)),
"r"(__size)
: "memory"
);
#else
// Unsupported architectures will have a linker error with a semi-decent error message
__cuda_ptx_cp_async_bulk_is_not_supported_before_SM_90__();
#endif
}
#endif // __cccl_ptx_isa >= 800

Expand Down
Loading
Loading