CodeQL / Analyze (actions) (push) Canceled after 0s
CodeQL / Analyze (python) (push) Canceled after 0s
pre-commit / pre-commit (push) Canceled after 0s
Multi-Arch CI / setup (push) Canceled after 0s
Publish build_manylinux_x86_64 images / publish_build_manylinux_x86_64 (push) Canceled after 0s
Setup (Multi-Arch) / setup (push) Canceled after 0s
Multi-Arch CI / Linux::${{ fromJSON(needs.setup.outputs.linux_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubi10 image / publish_no_rocm_image_ubi10 (push) Canceled after 0s
Multi-Arch CI / Windows::${{ fromJSON(needs.setup.outputs.windows_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Multi-Arch CI / Manifest Diff (push) Canceled after 0s
TheRock Manifest Diff Report / Generate Manifest Diff Report (push) Canceled after 0s
Multi-Arch CI / CI Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_ocl_rt images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_rocgdb images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Unit Tests / Unit Tests :: ubuntu-24.04 (push) Canceled after 0s
Unit Tests / Unit Tests :: windows-2022 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_media (push) Canceled after 0s
Unit Tests / Unit Tests Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_openmpi (push) Canceled after 0s
106 lines
2.7 KiB
C++
106 lines
2.7 KiB
C++
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
|
|
// SPDX-License-Identifier: MIT
|
|
|
|
#ifndef CK_ARRAY_HPP
|
|
#define CK_ARRAY_HPP
|
|
|
|
#include "functional2.hpp"
|
|
#include "sequence.hpp"
|
|
#include <type_traits>
|
|
#include <cassert>
|
|
|
|
#if __clang_major__ >= 23
|
|
#pragma clang diagnostic push
|
|
#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
|
|
#endif
|
|
namespace ck {
|
|
|
|
template <typename TData, index_t NSize>
|
|
struct Array
|
|
{
|
|
using type = Array;
|
|
using data_type = TData;
|
|
|
|
TData mData[NSize];
|
|
|
|
__host__ __device__ static constexpr index_t Size() { return NSize; }
|
|
|
|
__host__ __device__ constexpr const TData& At(index_t i) const [[clang::lifetimebound]]
|
|
{
|
|
return mData[i];
|
|
}
|
|
|
|
__host__ __device__ constexpr TData& At(index_t i) [[clang::lifetimebound]] { return mData[i]; }
|
|
|
|
__host__ __device__ constexpr const TData& operator[](index_t i) const [[clang::lifetimebound]]
|
|
{
|
|
return At(i);
|
|
}
|
|
|
|
__host__ __device__ constexpr TData& operator()(index_t i) [[clang::lifetimebound]]
|
|
{
|
|
return At(i);
|
|
}
|
|
|
|
template <typename... Args>
|
|
__host__ constexpr auto Emplace(index_t i, Args&&... args)
|
|
-> std::enable_if_t<std::is_nothrow_constructible_v<TData, Args&&...>>
|
|
{
|
|
assert(i >= 0 && i < NSize);
|
|
mData[i].~TData();
|
|
new(mData + i) TData(ck::forward<Args>(args)...);
|
|
}
|
|
|
|
template <typename T>
|
|
__host__ __device__ constexpr auto operator=(const T& a)
|
|
{
|
|
static_assert(T::Size() == Size(), "wrong! size not the same");
|
|
|
|
static_for<0, Size(), 1>{}([&](auto i) { operator()(i) = a[i]; });
|
|
|
|
return *this;
|
|
}
|
|
__host__ __device__ constexpr const TData* begin() const [[clang::lifetimebound]]
|
|
{
|
|
return &mData[0];
|
|
}
|
|
__host__ __device__ constexpr const TData* end() const [[clang::lifetimebound]]
|
|
{
|
|
return &mData[NSize];
|
|
}
|
|
__host__ __device__ constexpr TData* begin() [[clang::lifetimebound]] { return &mData[0]; }
|
|
__host__ __device__ constexpr TData* end() [[clang::lifetimebound]] { return &mData[NSize]; }
|
|
};
|
|
|
|
// empty Array
|
|
template <typename TData>
|
|
struct Array<TData, 0>
|
|
{
|
|
using type = Array;
|
|
using data_type = TData;
|
|
|
|
__host__ __device__ static constexpr index_t Size() { return 0; }
|
|
};
|
|
|
|
template <typename X, typename... Xs>
|
|
__host__ __device__ constexpr auto make_array(X&& x, Xs&&... xs)
|
|
{
|
|
using data_type = remove_cvref_t<X>;
|
|
return Array<data_type, sizeof...(Xs) + 1>{ck::forward<X>(x), ck::forward<Xs>(xs)...};
|
|
}
|
|
|
|
// make empty array
|
|
template <typename X>
|
|
__host__ __device__ constexpr auto make_array()
|
|
{
|
|
return Array<X, 0>{};
|
|
}
|
|
|
|
} // namespace ck
|
|
|
|
#if __clang_major__ >= 23
|
|
#pragma clang diagnostic pop
|
|
#endif
|
|
|
|
#endif
|