include/ck/utility/amd_lds.hpp Source File

• include
• ck
• utility

amd_lds.hpp

Go to the documentation of this file.

1 // SPDX-License-Identifier: MIT

2 // Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.

3 

4 #pragma once

5 

6 #include "ck/utility/amd_address_space.hpp"

7 #include "ck/utility/dynamic_buffer.hpp"

8 #include "ck/utility/math.hpp"

9 

10 namespace ck {

11 

12 namespace lds_utils {

13 

24 template <typename DataType, index_t NumBuffers>

25 __device__ static auto AllocateLdsBuffers(void* lds_ptr,

26  int32_t num_elems_per_buffer,

27  int32_t start_offset_elems,

28  int32_t lds_alignment)

29 {

30  const DataType* lds_start = static_cast<DataType*>(lds_ptr) + start_offset_elems;

31  const int32_t single_buffer_offset =

32  math::integer_least_multiple(num_elems_per_buffer, lds_alignment);

33  return generate_tuple(

34  [&](auto i) {

35  const int32_t local_offset = i * single_buffer_offset;

36  return make_dynamic_buffer<AddressSpaceEnum::Lds>(lds_start + local_offset,

37  num_elems_per_buffer);

38  },

39  Number<NumBuffers>{});

40 }

41 

42 } // namespace lds_utils

43 } // namespace ck

amd_address_space.hpp

dynamic_buffer.hpp

ck::math::integer_least_multiple

__host__ constexpr __device__ auto integer_least_multiple(X x, Y y)

Definition: math.hpp:78

ck

Definition: ck.hpp:264

ck::generate_tuple

__host__ constexpr __device__ auto generate_tuple(F &&f, Number< N >)

Definition: tuple_helper.hpp:15

ck::integral_constant

Definition: integral_constant.hpp:10

math.hpp