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
c521ee62
"configs/vscode:/vscode.git/clone" did not exist on "b7184e9db577a4ba391f5c33eb0a3f4e147204e1"
Commit
c521ee62
authored
Apr 27, 2022
by
Jianfeng yan
Browse files
fixed segm fault caused by a typo in creating a/b_grid_desc
parent
09b9ddb3
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
15 additions
and
3 deletions
+15
-3
include/ck/tensor_operation/gpu/device/device_gemm_xdl_splitk_c_shuffle.hpp
...operation/gpu/device/device_gemm_xdl_splitk_c_shuffle.hpp
+15
-2
test/CMakeLists.txt
test/CMakeLists.txt
+0
-1
No files found.
include/ck/tensor_operation/gpu/device/device_gemm_xdl_splitk_c_shuffle.hpp
View file @
c521ee62
...
@@ -586,8 +586,8 @@ struct DeviceGemmXdlSplitKCShuffle
...
@@ -586,8 +586,8 @@ struct DeviceGemmXdlSplitKCShuffle
const
auto
AKSplitted
=
AKPad
/
k_batch
;
const
auto
AKSplitted
=
AKPad
/
k_batch
;
const
auto
BKSplitted
=
BKPad
/
k_batch
;
const
auto
BKSplitted
=
BKPad
/
k_batch
;
a_grid_desc_ak0_m_ak1_
=
DeviceOp
::
MakeAGridDescriptor_AK0_M_AK1
(
MRaw
,
AK
Pa
d
,
StrideA
);
a_grid_desc_ak0_m_ak1_
=
DeviceOp
::
MakeAGridDescriptor_AK0_M_AK1
(
MRaw
,
AK
Splitte
d
,
StrideA
);
b_grid_desc_bk0_n_bk1_
=
DeviceOp
::
MakeBGridDescriptor_BK0_N_BK1
(
BK
Pa
d
,
NRaw
,
StrideB
);
b_grid_desc_bk0_n_bk1_
=
DeviceOp
::
MakeBGridDescriptor_BK0_N_BK1
(
BK
Splitte
d
,
NRaw
,
StrideB
);
c_grid_desc_m_n_
=
DeviceOp
::
MakeCGridDescriptor_M_N
(
MRaw
,
NRaw
,
StrideC
);
c_grid_desc_m_n_
=
DeviceOp
::
MakeCGridDescriptor_M_N
(
MRaw
,
NRaw
,
StrideC
);
if
(
GridwiseGemm
::
CheckValidity
(
if
(
GridwiseGemm
::
CheckValidity
(
...
@@ -626,6 +626,19 @@ struct DeviceGemmXdlSplitKCShuffle
...
@@ -626,6 +626,19 @@ struct DeviceGemmXdlSplitKCShuffle
block_2_ctile_map_
=
MakeBlock2CTileMap
(
BatchCount_
,
c_grid_desc_m_n_
,
1
,
1
);
block_2_ctile_map_
=
MakeBlock2CTileMap
(
BatchCount_
,
c_grid_desc_m_n_
,
1
,
1
);
}
}
for
(
int
batch
=
0
;
batch
<
BatchCount_
;
++
batch
)
{
printf
(
"batch = %d, ptr_offset = [%ld, %ld, %ld]
\n
"
,
batch
,
compute_ptr_offset_of_batch_
.
GetAPtrOffset
(
batch
),
compute_ptr_offset_of_batch_
.
GetBPtrOffset
(
batch
),
compute_ptr_offset_of_batch_
.
GetCPtrOffset
(
batch
));
}
const
index_t
grid_size
=
GridwiseGemm
::
CalculateGridSize
(
c_grid_desc_m_n_
)
*
BatchCount_
;
for
(
int
blk
=
0
;
blk
<
grid_size
;
++
blk
)
{
const
auto
ctile_idx
=
block_2_ctile_map_
.
CalculateBottomIndex
(
make_multi_index
(
blk
));
printf
(
"blk= %d, ctile_idx = [%d, %d]
\n
"
,
blk
,
ctile_idx
[
I0
],
ctile_idx
[
I1
]);
}
}
}
// private:
// private:
...
...
test/CMakeLists.txt
View file @
c521ee62
...
@@ -45,4 +45,3 @@ add_subdirectory(grouped_gemm)
...
@@ -45,4 +45,3 @@ add_subdirectory(grouped_gemm)
add_subdirectory
(
convnd_fwd
)
add_subdirectory
(
convnd_fwd
)
add_subdirectory
(
reduce
)
add_subdirectory
(
reduce
)
add_subdirectory
(
conv2d_bwd_weight
)
add_subdirectory
(
conv2d_bwd_weight
)
add_subdirectory
(
fp16_transfer_bf16
)
\ No newline at end of file
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