"sgl-kernel/csrc/vscode:/vscode.git/clone" did not exist on "b43263307f40a206f1371e4064d410a136d4e004"
dynamic_buffer.hpp 1.94 KB
Newer Older
Chao Liu's avatar
Chao Liu committed
1
// SPDX-License-Identifier: MIT
Illia Silin's avatar
Illia Silin committed
2
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
Chao Liu's avatar
Chao Liu committed
3

4
#pragma once
Chao Liu's avatar
Chao Liu committed
5

Chao Liu's avatar
Chao Liu committed
6
7
8
#include "buffer_view.hpp"

// FIXME: deprecate DynamicBuffer, use BufferView instead
9
10
11

namespace ck {

Chao Liu's avatar
Chao Liu committed
12
// FIXME: deprecate DynamicBuffer, use BufferView instead
13
template <AddressSpaceEnum BufferAddressSpace,
14
15
          typename T,
          typename ElementSpaceSize,
16
          bool InvalidElementUseNumericalZeroValue,
Chao Liu's avatar
Chao Liu committed
17
18
19
20
21
22
23
24
          AmdBufferCoherenceEnum Coherence = AmdBufferCoherenceEnum::DefaultCoherence>
using DynamicBuffer = BufferView<BufferAddressSpace,
                                 T,
                                 ElementSpaceSize,
                                 InvalidElementUseNumericalZeroValue,
                                 Coherence>;

// FIXME: deprecate make_dynamic_buffer, use make_buffer_view instead
25
template <AddressSpaceEnum BufferAddressSpace,
Chao Liu's avatar
Chao Liu committed
26
          AmdBufferCoherenceEnum Coherence = AmdBufferCoherenceEnum::DefaultCoherence,
27
28
          typename T,
          typename ElementSpaceSize>
29
30
__host__ __device__ constexpr auto make_dynamic_buffer(T* p, ElementSpaceSize element_space_size)
{
Chao Liu's avatar
Chao Liu committed
31
32
    return make_buffer_view<BufferAddressSpace, Coherence, T, ElementSpaceSize>(p,
                                                                                element_space_size);
33
34
}

Chao Liu's avatar
Chao Liu committed
35
// FIXME: deprecate make_dynamic_buffer, use make_buffer_view instead
36
template <
37
    AddressSpaceEnum BufferAddressSpace,
Chao Liu's avatar
Chao Liu committed
38
    AmdBufferCoherenceEnum Coherence = AmdBufferCoherenceEnum::DefaultCoherence,
39
40
41
42
    typename T,
    typename ElementSpaceSize,
    typename X,
    typename enable_if<is_same<remove_cvref_t<T>, remove_cvref_t<X>>::value, bool>::type = false>
43
__host__ __device__ constexpr auto
44
make_dynamic_buffer(T* p, ElementSpaceSize element_space_size, X invalid_element_value)
45
{
Chao Liu's avatar
Chao Liu committed
46
47
    return make_buffer_view<BufferAddressSpace, Coherence, T, ElementSpaceSize>(
        p, element_space_size, invalid_element_value);
48
49
50
}

} // namespace ck