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

10namespace ck {

11

12namespace lds_utils {

13

24template <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::lds_utils

Definition amd_lds.hpp:12

ck::math::integer_least_multiple

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

Definition utility/math.hpp:78

ck

Definition ck.hpp:268

ck::Number

integral_constant< index_t, N > Number

Definition number.hpp:12

ck::generate_tuple

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

Definition tuple_helper.hpp:21

ck::make_dynamic_buffer

__host__ __device__ constexpr auto make_dynamic_buffer(T *p, ElementSpaceSize element_space_size)

Definition dynamic_buffer.hpp:472

int32_t

signed int int32_t

Definition stdint.h:123

math.hpp