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
f63ca8e8
"...composable_kernel.git" did not exist on "4daedf8ca56f3bd93481708bd9d762045839ec20"
Commit
f63ca8e8
authored
May 19, 2022
by
myamlak
Browse files
Format
parent
a7676df9
Changes
3
Hide whitespace changes
Inline
Side-by-side
Showing
3 changed files
with
13 additions
and
5 deletions
+13
-5
example/20_cgemm/cgemm_xdl_bf16.cpp
example/20_cgemm/cgemm_xdl_bf16.cpp
+10
-2
include/ck/tensor_operation/gpu/device/device_cgemm_4gemm_xdl_cshuffle.hpp
..._operation/gpu/device/device_cgemm_4gemm_xdl_cshuffle.hpp
+1
-1
include/ck/tensor_operation/gpu/element/binary_element_wise_operation.hpp
...r_operation/gpu/element/binary_element_wise_operation.hpp
+2
-2
No files found.
example/20_cgemm/cgemm_xdl_bf16.cpp
View file @
f63ca8e8
...
@@ -278,8 +278,16 @@ int main(int argc, char* argv[])
...
@@ -278,8 +278,16 @@ int main(int argc, char* argv[])
ref_invoker
.
Run
(
ref_argument
);
ref_invoker
.
Run
(
ref_argument
);
ck
::
utils
::
check_err
(
c_m_n_real_device_f32_result
.
mData
,
c_m_n_real_host_result
.
mData
);
ck
::
utils
::
check_err
(
c_m_n_real_device_f32_result
.
mData
,
ck
::
utils
::
check_err
(
c_m_n_imag_device_f32_result
.
mData
,
c_m_n_imag_host_result
.
mData
);
c_m_n_real_host_result
.
mData
,
"Verification error: incorrect results in real part!"
,
1e-2
f
,
1e-3
f
);
ck
::
utils
::
check_err
(
c_m_n_imag_device_f32_result
.
mData
,
c_m_n_imag_host_result
.
mData
,
"Verification error: incorrect results in imaginary part!"
,
1e-2
f
,
1e-3
f
);
}
}
return
0
;
return
0
;
...
...
include/ck/tensor_operation/gpu/device/device_cgemm_4gemm_xdl_cshuffle.hpp
View file @
f63ca8e8
...
@@ -100,7 +100,7 @@ struct DeviceCGemm_4Gemm_Xdl_CShuffle
...
@@ -100,7 +100,7 @@ struct DeviceCGemm_4Gemm_Xdl_CShuffle
make_tuple
(
generate_sequence_v2
([
&
](
auto
I
)
{
return
I
;
},
Number
<
2
>
{})),
make_tuple
(
generate_sequence_v2
([
&
](
auto
I
)
{
return
I
;
},
Number
<
2
>
{})),
make_tuple
(
Sequence
<
0
>
{}));
make_tuple
(
Sequence
<
0
>
{}));
return
PadDescriptor_M0_1d
(
desc_m0
,
gridSize
,
blockSize
);
return
PadDescriptor_M0_1d
(
desc_m0
,
gridSize
,
blockSize
);
}
}
static
auto
MakeAGridDescriptor_AK0_M_AK1
(
index_t
MRaw
,
index_t
KRaw
,
index_t
StrideA
)
static
auto
MakeAGridDescriptor_AK0_M_AK1
(
index_t
MRaw
,
index_t
KRaw
,
index_t
StrideA
)
...
...
include/ck/tensor_operation/gpu/element/binary_element_wise_operation.hpp
View file @
f63ca8e8
...
@@ -42,8 +42,8 @@ struct Substract
...
@@ -42,8 +42,8 @@ struct Substract
{
{
dst
=
src1
-
src2
;
dst
=
src1
-
src2
;
}
}
__host__
__device__
constexpr
void
__host__
__device__
constexpr
void
operator
()(
float
&
dst
,
const
float
&
src1
,
const
float
&
src2
)
const
operator
()(
float
&
dst
,
const
float
&
src1
,
const
float
&
src2
)
const
{
{
dst
=
src1
-
src2
;
dst
=
src1
-
src2
;
...
...
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