Skip to content
GitLab
Menu
Projects
Groups
Snippets
Loading...
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in / Register
Toggle navigation
Menu
Open sidebar
gaoqiong
composable_kernel
Commits
32bac6f3
Commit
32bac6f3
authored
Aug 01, 2023
by
aska-0096
Browse files
Temp save
parent
66e61076
Changes
3
Hide whitespace changes
Inline
Side-by-side
Showing
3 changed files
with
13 additions
and
8 deletions
+13
-8
include/ck/tensor_operation/gpu/block/blockwise_fpAintB_gemm_wmma.hpp
...ensor_operation/gpu/block/blockwise_fpAintB_gemm_wmma.hpp
+1
-1
include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer.hpp
...operation/gpu/thread/threadwise_tensor_slice_transfer.hpp
+4
-5
include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v3r1.hpp
...tion/gpu/thread/threadwise_tensor_slice_transfer_v3r1.hpp
+8
-2
No files found.
include/ck/tensor_operation/gpu/block/blockwise_fpAintB_gemm_wmma.hpp
View file @
32bac6f3
...
@@ -419,7 +419,7 @@ struct Blockwise_fpAintB_GemmWMMA
...
@@ -419,7 +419,7 @@ struct Blockwise_fpAintB_GemmWMMA
// convert B from int8 to fp16, multiply scale
// convert B from int8 to fp16, multiply scale
static_for
<
0
,
b_thread_buf
.
Size
(),
1
>
{}([
&
](
auto
i
)
{
static_for
<
0
,
b_thread_buf
.
Size
(),
1
>
{}([
&
](
auto
i
)
{
converted_b_thread_buf
(
i
)
=
scale_thread_buf
[
i
/
WmmaK
]
*
converted_b_thread_buf
(
i
)
=
scale_thread_buf
[
i
/
WmmaK
]
*
type_convert
<
ADataType
>
(
b_thread_buf
[
i
]);
type_convert
<
ADataType
>
(
b_thread_buf
[
i
]);
// call byte permute
});
});
// read A
// read A
a_thread_copy_
.
Run
(
a_thread_copy_
.
Run
(
...
...
include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer.hpp
View file @
32bac6f3
...
@@ -1143,7 +1143,9 @@ struct ThreadwiseTensorSliceTransfer_v4
...
@@ -1143,7 +1143,9 @@ struct ThreadwiseTensorSliceTransfer_v4
const
bool
is_src_valid
=
coordinate_has_valid_offset_assuming_visible_index_is_valid
(
const
bool
is_src_valid
=
coordinate_has_valid_offset_assuming_visible_index_is_valid
(
src_desc
,
src_data_coord
);
src_desc
,
src_data_coord
);
#if 0
printf("Tid: %03d, LDS read offset: %d\n", get_thread_local_1d_id(), src_data_coord.GetOffset());
#endif
// copy data from src_buf into src_tmp_vector
// copy data from src_buf into src_tmp_vector
if
constexpr
(
SrcBuffer
::
IsDynamicBuffer
())
if
constexpr
(
SrcBuffer
::
IsDynamicBuffer
())
{
{
...
@@ -1417,10 +1419,7 @@ struct ThreadwiseTensorSliceTransfer_StaticToStatic_InterRow
...
@@ -1417,10 +1419,7 @@ struct ThreadwiseTensorSliceTransfer_StaticToStatic_InterRow
1
,
1
,
0
);
0
);
v_theother_row
=
type_convert_sp
<
SrcData
>
(
temp
);
v_theother_row
=
type_convert_sp
<
SrcData
>
(
temp
);
// if (get_thread_local_1d_id() == 0){
// printf("src_offset:%d, dst_offset for this row: %d, dst_offset
// for the other row: %d \n",
// src_offset, dst_offset, dst_offset+DstScalarPerVector);}
if
(
get_thread_local_1d_id
()
%
32
<
16
)
if
(
get_thread_local_1d_id
()
%
32
<
16
)
{
{
// apply type convert
// apply type convert
...
...
include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v3r1.hpp
View file @
32bac6f3
...
@@ -207,7 +207,11 @@ struct ThreadwiseTensorSliceTransfer_v3r1
...
@@ -207,7 +207,11 @@ struct ThreadwiseTensorSliceTransfer_v3r1
// copy data from src_buf into src_vector_container
// copy data from src_buf into src_vector_container
auto
src_vector_container
=
src_vector_type
{
auto
src_vector_container
=
src_vector_type
{
src_buf
.
template
Get
<
src_vector_t
>(
src_coord_
.
GetOffset
(),
is_src_valid
)};
src_buf
.
template
Get
<
src_vector_t
>(
src_coord_
.
GetOffset
(),
is_src_valid
)};
if
(
false
){
printf
(
"Tid: %03d, a_grid_buf: %04x
\n
"
,
get_thread_local_1d_id
(),
*
(
reinterpret_cast
<
const
uint16_t
*>
(
&
src_vector_container
.
template
AsType
<
SrcData
>()[
Number
<
0
>
{}])));
}
// copy data from src_vector_container into src_thread_scratch_
// copy data from src_vector_container into src_thread_scratch_
src_thread_scratch_tuple_
(
thread_scratch_id
)
src_thread_scratch_tuple_
(
thread_scratch_id
)
.
template
SetAsType
<
src_vector_t
>(
.
template
SetAsType
<
src_vector_t
>(
...
@@ -442,7 +446,9 @@ struct ThreadwiseTensorSliceTransfer_v3r1
...
@@ -442,7 +446,9 @@ struct ThreadwiseTensorSliceTransfer_v3r1
const
bool
is_dst_valid
=
const
bool
is_dst_valid
=
coordinate_has_valid_offset_assuming_visible_index_is_valid
(
dst_desc
,
dst_coord_
);
coordinate_has_valid_offset_assuming_visible_index_is_valid
(
dst_desc
,
dst_coord_
);
#if 0
printf("Tid: %03d, LDS write offset: %d\n", get_thread_local_1d_id(), dst_coord_.GetOffset());
#endif
using
dst_vector_type
=
vector_type_maker_t
<
DstData
,
DstScalarPerVector
>
;
using
dst_vector_type
=
vector_type_maker_t
<
DstData
,
DstScalarPerVector
>
;
using
dst_vector_t
=
typename
dst_vector_type
::
type
;
using
dst_vector_t
=
typename
dst_vector_type
::
type
;
...
...
Write
Preview
Markdown
is supported
0%
Try again
or
attach a new file
.
Attach a file
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment