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

Jump to bottom

Refactor f8_t and bf8_t as custom types, enable use of custom types #1167

Closed
wants to merge 10 commits into from
Closed

Refactor f8_t and bf8_t as custom types, enable use of custom types #1167

Show file tree
Hide file tree
Changes from 5 commits
Commits
Show all changes
10 commits
Select commit Hold shift + click to select a range
01d8575
Refactor f8_t and bf8_t as custom types and update conversion methods
geyyer Feb 14, 2024
f7170d5
For load use built-in type instead of custom type
geyyer Feb 14, 2024
6aad2f2
Add dynamic buffer get method implementation for custom types
geyyer Feb 14, 2024
59ea031
Update conthreadwise conversion pipeline
geyyer Feb 14, 2024
0291aa6
Merge branch 'develop' into lwpck-1271
geyyer Feb 14, 2024
048d5e3
Add element-wise ops
geyyer Feb 27, 2024
77c5212
Refactor static buffer setastype for custom types
geyyer Feb 27, 2024
cf9568e
Update custom types to support various src/dst types combinations
geyyer Mar 26, 2024
cb5297d
Format
geyyer Mar 28, 2024
00381ea
Fix operator overloading
geyyer Apr 16, 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
12 changes: 8 additions & 4 deletions include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v3r1.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -237,13 +237,17 @@ struct ThreadwiseTensorSliceTransfer_v3r1

constexpr index_t elem_op_vec_len = get_elem_op_vec_len();

using src_elem_op_vec_t = typename vector_type<SrcData, elem_op_vec_len>::type;
using dst_elem_op_vec_t = typename vector_type<DstData, elem_op_vec_len>::type;
using src_elem_op_vec_t =
typename vector_type<SrcData, elem_op_vec_len>::conversion_type;
using dst_elem_op_vec_t =
typename vector_type<DstData, elem_op_vec_len>::conversion_type;

static_for<0, SrcScalarPerVector / elem_op_vec_len, 1>{}([&](auto idx) {
// apply the src elementwise op and convert to DstData under the hood if needed
src_element_op_(op_r_v.template AsType<dst_elem_op_vec_t>()(idx),
src_vector_container.template AsType<src_elem_op_vec_t>()[idx]);
src_element_op_(
op_r_v.template AsType<dst_elem_op_vec_t>()(idx),
bit_cast<src_elem_op_vec_t>(
src_vector_container.template AsType<src_elem_op_vec_t>()[idx]));
});

// copy data from src_vector_container into src_thread_scratch_
Expand Down
5 changes: 3 additions & 2 deletions include/ck/utility/amd_buffer_addressing.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -415,8 +415,9 @@ __device__ typename vector_type<T, N>::type amd_buffer_load_impl(int32x4_t src_w
(is_same<T, half_t>::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) ||
(is_same<T, bhalf_t>::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) ||
(is_same<T, int32_t>::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) ||
(is_same<T, f8_t>::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) ||
(is_same<T, bf8_t>::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) ||
(is_same<T, _BitInt(8)>::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) ||
(is_same<T, unsigned _BitInt(8)>::value &&
(N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) ||
(is_same<T, int8_t>::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)),
"wrong! not implemented");

Expand Down
Loading