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_ROCM
Commits
51b9abb9
Commit
51b9abb9
authored
Nov 08, 2024
by
Andriy Roshchenko
Browse files
Verify more tests on floating point data
parent
646b8e5c
Changes
11
Hide whitespace changes
Inline
Side-by-side
Showing
11 changed files
with
59 additions
and
25 deletions
+59
-25
CMakePresets.json
CMakePresets.json
+4
-2
example/01_gemm/run_gemm_example_streamk_v2.inc
example/01_gemm/run_gemm_example_streamk_v2.inc
+4
-4
example/01_gemm/run_gemm_example_v2.inc
example/01_gemm/run_gemm_example_v2.inc
+3
-11
example/15_grouped_gemm/grouped_gemm_xdl_fixed_nk_fp16_fp8.cpp
...le/15_grouped_gemm/grouped_gemm_xdl_fixed_nk_fp16_fp8.cpp
+1
-1
example/62_convnd_activ/convinvscale/run_convnd_fwd_convinvscale_example.inc
...ctiv/convinvscale/run_convnd_fwd_convinvscale_example.inc
+1
-1
example/62_convnd_activ/convscale/run_convnd_fwd_convscale_example.inc
...nvnd_activ/convscale/run_convnd_fwd_convscale_example.inc
+1
-1
example/62_convnd_activ/convscale_add/run_convnd_fwd_convscale_add_example.inc
...iv/convscale_add/run_convnd_fwd_convscale_add_example.inc
+1
-1
example/62_convnd_activ/convscale_reduce/run_convnd_fwd_example.inc
..._convnd_activ/convscale_reduce/run_convnd_fwd_example.inc
+1
-1
example/62_convnd_activ/convscale_relu/run_convnd_fwd_convscale_relu_example.inc
.../convscale_relu/run_convnd_fwd_convscale_relu_example.inc
+1
-1
example/65_gemm_multiply_multiply/gemm_multiply_multiply_xdl_fp8.cpp
...gemm_multiply_multiply/gemm_multiply_multiply_xdl_fp8.cpp
+1
-1
example/65_gemm_multiply_multiply/gemm_multiply_multiply_xdl_fp8_ab_scale.cpp
...iply_multiply/gemm_multiply_multiply_xdl_fp8_ab_scale.cpp
+41
-1
No files found.
CMakePresets.json
View file @
51b9abb9
...
@@ -66,7 +66,8 @@
...
@@ -66,7 +66,8 @@
"cacheVariables"
:
{
"cacheVariables"
:
{
"GPU_TARGETS"
:
"gfx90a"
,
"GPU_TARGETS"
:
"gfx90a"
,
"CMAKE_BUILD_TYPE"
:
"Release"
,
"CMAKE_BUILD_TYPE"
:
"Release"
,
"CMAKE_CXX_FLAGS"
:
"-O3"
"CMAKE_CXX_FLAGS"
:
"-O3"
,
"CK_USE_FP8_ON_UNSUPPORTED_ARCH"
:
"ON"
}
}
},
},
{
{
...
@@ -76,7 +77,8 @@
...
@@ -76,7 +77,8 @@
"cacheVariables"
:
{
"cacheVariables"
:
{
"GPU_TARGETS"
:
"gfx90a"
,
"GPU_TARGETS"
:
"gfx90a"
,
"CMAKE_BUILD_TYPE"
:
"Debug"
,
"CMAKE_BUILD_TYPE"
:
"Debug"
,
"CMAKE_CXX_FLAGS"
:
"-O0 -ggdb"
"CMAKE_CXX_FLAGS"
:
"-O0 -ggdb"
,
"CK_USE_FP8_ON_UNSUPPORTED_ARCH"
:
"ON"
}
}
},
},
{
{
...
...
example/01_gemm/run_gemm_example_streamk_v2.inc
View file @
51b9abb9
...
@@ -162,12 +162,12 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config)
...
@@ -162,12 +162,12 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config)
b_k_n
.
GenerateTensorValue
(
GeneratorTensor_2
<
BDataType
>
{
-
2
,
2
});
b_k_n
.
GenerateTensorValue
(
GeneratorTensor_2
<
BDataType
>
{
-
2
,
2
});
break
;
break
;
case
2
:
case
2
:
a_m_k
.
GenerateTensorValue
(
GeneratorTensor_
1
<
ADataType
>
{
1
});
a_m_k
.
GenerateTensorValue
(
GeneratorTensor_
3
<
ADataType
>
{
0.0
,
1.0
});
b_k_n
.
GenerateTensorValue
(
GeneratorTensor_
2
<
BDataType
>
{
-
2
,
2
});
b_k_n
.
GenerateTensorValue
(
GeneratorTensor_
3
<
BDataType
>
{
-
2
,
2
});
break
;
break
;
case
3
:
case
3
:
a_m_k
.
GenerateTensorValue
(
GeneratorTensor_
2
<
ADataType
>
{
-
2
,
2
});
a_m_k
.
GenerateTensorValue
(
GeneratorTensor_
3
<
ADataType
>
{
-
2
,
2
});
b_k_n
.
GenerateTensorValue
(
GeneratorTensor_
1
<
B
DataType
>
{
1
});
b_k_n
.
GenerateTensorValue
(
GeneratorTensor_
3
<
A
DataType
>
{
0.0
,
1.0
});
break
;
break
;
default
:
default
:
a_m_k
.
GenerateTensorValue
(
GeneratorTensor_3
<
ADataType
>
{
0.0
,
1.0
});
a_m_k
.
GenerateTensorValue
(
GeneratorTensor_3
<
ADataType
>
{
0.0
,
1.0
});
...
...
example/01_gemm/run_gemm_example_v2.inc
View file @
51b9abb9
...
@@ -146,19 +146,11 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config)
...
@@ -146,19 +146,11 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config)
b_k_n
.
GenerateTensorValue
(
GeneratorTensor_1
<
BDataType
>
{
1
});
b_k_n
.
GenerateTensorValue
(
GeneratorTensor_1
<
BDataType
>
{
1
});
break
;
break
;
case
1
:
case
1
:
a_m_k
.
GenerateTensorValue
(
GeneratorTensor_2
<
ADataType
>
{
-
2
,
2
});
a_m_k
.
GenerateTensorValue
(
GeneratorTensor_3
<
ADataType
>
{
-
2
,
2
});
b_k_n
.
GenerateTensorValue
(
GeneratorTensor_2
<
BDataType
>
{
-
2
,
2
});
b_k_n
.
GenerateTensorValue
(
GeneratorTensor_3
<
BDataType
>
{
-
3
,
3
});
break
;
case
2
:
a_m_k
.
GenerateTensorValue
(
GeneratorTensor_1
<
ADataType
>
{
1
});
b_k_n
.
GenerateTensorValue
(
GeneratorTensor_2
<
BDataType
>
{
-
2
,
2
});
break
;
case
3
:
a_m_k
.
GenerateTensorValue
(
GeneratorTensor_2
<
ADataType
>
{
-
2
,
2
});
b_k_n
.
GenerateTensorValue
(
GeneratorTensor_1
<
BDataType
>
{
1
});
break
;
break
;
default
:
default
:
a_m_k
.
GenerateTensorValue
(
GeneratorTensor_3
<
ADataType
>
{
0
.0
,
1.0
});
a_m_k
.
GenerateTensorValue
(
GeneratorTensor_3
<
ADataType
>
{
-
1
.0
,
1.0
});
b_k_n
.
GenerateTensorValue
(
GeneratorTensor_3
<
BDataType
>
{
-
0.5
,
0.5
});
b_k_n
.
GenerateTensorValue
(
GeneratorTensor_3
<
BDataType
>
{
-
0.5
,
0.5
});
}
}
...
...
example/15_grouped_gemm/grouped_gemm_xdl_fixed_nk_fp16_fp8.cpp
View file @
51b9abb9
...
@@ -75,7 +75,7 @@ struct ProblemSize final
...
@@ -75,7 +75,7 @@ struct ProblemSize final
struct
ExecutionConfig
final
struct
ExecutionConfig
final
{
{
bool
do_verification
=
true
;
bool
do_verification
=
true
;
int
init_method
=
1
;
int
init_method
=
2
;
int
k_batch
=
1
;
int
k_batch
=
1
;
bool
time_kernel
=
false
;
bool
time_kernel
=
false
;
};
};
...
...
example/62_convnd_activ/convinvscale/run_convnd_fwd_convinvscale_example.inc
View file @
51b9abb9
...
@@ -8,7 +8,7 @@ bool run_convnd_fwd_example(int argc, char* argv[])
...
@@ -8,7 +8,7 @@ bool run_convnd_fwd_example(int argc, char* argv[])
print_helper_msg
();
print_helper_msg
();
bool
do_verification
=
true
;
bool
do_verification
=
true
;
int
init_method
=
1
;
int
init_method
=
2
;
bool
time_kernel
=
false
;
bool
time_kernel
=
false
;
ck
::
utils
::
conv
::
ConvParam
conv_param
{
ck
::
utils
::
conv
::
ConvParam
conv_param
{
...
...
example/62_convnd_activ/convscale/run_convnd_fwd_convscale_example.inc
View file @
51b9abb9
...
@@ -8,7 +8,7 @@ bool run_convnd_fwd_example(int argc, char* argv[])
...
@@ -8,7 +8,7 @@ bool run_convnd_fwd_example(int argc, char* argv[])
print_helper_msg
();
print_helper_msg
();
bool
do_verification
=
true
;
bool
do_verification
=
true
;
int
init_method
=
1
;
int
init_method
=
2
;
bool
time_kernel
=
false
;
bool
time_kernel
=
false
;
ck
::
utils
::
conv
::
ConvParam
conv_param
{
ck
::
utils
::
conv
::
ConvParam
conv_param
{
...
...
example/62_convnd_activ/convscale_add/run_convnd_fwd_convscale_add_example.inc
View file @
51b9abb9
...
@@ -8,7 +8,7 @@ bool run_convnd_fwd_example(int argc, char* argv[])
...
@@ -8,7 +8,7 @@ bool run_convnd_fwd_example(int argc, char* argv[])
print_helper_msg
();
print_helper_msg
();
bool
do_verification
=
true
;
bool
do_verification
=
true
;
int
init_method
=
1
;
int
init_method
=
2
;
bool
time_kernel
=
false
;
bool
time_kernel
=
false
;
ck
::
utils
::
conv
::
ConvParam
conv_param
{
ck
::
utils
::
conv
::
ConvParam
conv_param
{
...
...
example/62_convnd_activ/convscale_reduce/run_convnd_fwd_example.inc
View file @
51b9abb9
...
@@ -8,7 +8,7 @@ bool run_convnd_fwd_example(int argc, char* argv[])
...
@@ -8,7 +8,7 @@ bool run_convnd_fwd_example(int argc, char* argv[])
print_helper_msg
();
print_helper_msg
();
bool
do_verification
=
true
;
bool
do_verification
=
true
;
int
init_method
=
1
;
int
init_method
=
2
;
bool
time_kernel
=
false
;
bool
time_kernel
=
false
;
ck
::
utils
::
conv
::
ConvParam
conv_param
{
ck
::
utils
::
conv
::
ConvParam
conv_param
{
...
...
example/62_convnd_activ/convscale_relu/run_convnd_fwd_convscale_relu_example.inc
View file @
51b9abb9
...
@@ -8,7 +8,7 @@ bool run_convnd_fwd_example(int argc, char* argv[])
...
@@ -8,7 +8,7 @@ bool run_convnd_fwd_example(int argc, char* argv[])
print_helper_msg
();
print_helper_msg
();
bool
do_verification
=
true
;
bool
do_verification
=
true
;
int
init_method
=
1
;
int
init_method
=
2
;
bool
time_kernel
=
false
;
bool
time_kernel
=
false
;
ck
::
utils
::
conv
::
ConvParam
conv_param
{
ck
::
utils
::
conv
::
ConvParam
conv_param
{
...
...
example/65_gemm_multiply_multiply/gemm_multiply_multiply_xdl_fp8.cpp
View file @
51b9abb9
...
@@ -86,7 +86,7 @@ using DeviceOpInstance = ck::tensor_operation::device::DeviceGemmMultiD_Xdl_CShu
...
@@ -86,7 +86,7 @@ using DeviceOpInstance = ck::tensor_operation::device::DeviceGemmMultiD_Xdl_CShu
int
main
(
int
argc
,
char
*
argv
[])
int
main
(
int
argc
,
char
*
argv
[])
{
{
bool
do_verification
=
true
;
bool
do_verification
=
true
;
int
init_method
=
1
;
int
init_method
=
2
;
bool
time_kernel
=
false
;
bool
time_kernel
=
false
;
// GEMM shape
// GEMM shape
...
...
example/65_gemm_multiply_multiply/gemm_multiply_multiply_xdl_fp8_ab_scale.cpp
View file @
51b9abb9
...
@@ -78,7 +78,7 @@ using DeviceOpInstance = ck::tensor_operation::device::DeviceGemmMultiD_ABScale_
...
@@ -78,7 +78,7 @@ using DeviceOpInstance = ck::tensor_operation::device::DeviceGemmMultiD_ABScale_
int
main
(
int
argc
,
char
*
argv
[])
int
main
(
int
argc
,
char
*
argv
[])
{
{
bool
do_verification
=
true
;
bool
do_verification
=
true
;
int
init_method
=
1
;
int
init_method
=
5
;
bool
time_kernel
=
false
;
bool
time_kernel
=
false
;
// GEMM shape
// GEMM shape
...
@@ -186,6 +186,20 @@ int main(int argc, char* argv[])
...
@@ -186,6 +186,20 @@ int main(int argc, char* argv[])
b0_k_n
.
GenerateTensorValue
(
GeneratorTensor_1
<
B0DataType
>
{});
b0_k_n
.
GenerateTensorValue
(
GeneratorTensor_1
<
B0DataType
>
{});
a1_m_k
.
GenerateTensorValue
(
GeneratorTensor_3
<
A1DataType
>
{
0
,
1.0
});
a1_m_k
.
GenerateTensorValue
(
GeneratorTensor_3
<
A1DataType
>
{
0
,
1.0
});
b1_k_n
.
GenerateTensorValue
(
GeneratorTensor_3
<
B1DataType
>
{
0
,
1.0
});
b1_k_n
.
GenerateTensorValue
(
GeneratorTensor_3
<
B1DataType
>
{
0
,
1.0
});
break
;
case
6
:
a0_m_k
.
GenerateTensorValue
(
GeneratorTensor_PI
<
ADataType
>
{});
b0_k_n
.
GenerateTensorValue
(
GeneratorTensor_1
<
BDataType
>
{
0.5
f
});
a1_m_k
.
GenerateTensorValue
(
GeneratorTensor_1
<
A1DataType
>
{
0.5
});
b1_k_n
.
GenerateTensorValue
(
GeneratorTensor_1
<
B1DataType
>
{
4
});
break
;
case
7
:
a0_m_k
.
GenerateTensorValue
(
GeneratorTensor_PI_A
<
ADataType
>
{});
b0_k_n
.
GenerateTensorValue
(
GeneratorTensor_PI_B
<
BDataType
>
{});
a1_m_k
.
GenerateTensorValue
(
GeneratorTensor_1
<
A1DataType
>
{
2
});
b1_k_n
.
GenerateTensorValue
(
GeneratorTensor_1
<
B1DataType
>
{
0.5
});
break
;
break
;
default:
default:
a0_m_k
.
GenerateTensorValue
(
GeneratorTensor_3
<
A0DataType
>
{
-
0.5
,
0.5
});
a0_m_k
.
GenerateTensorValue
(
GeneratorTensor_3
<
A0DataType
>
{
-
0.5
,
0.5
});
...
@@ -306,6 +320,32 @@ int main(int argc, char* argv[])
...
@@ -306,6 +320,32 @@ int main(int argc, char* argv[])
e_device_buf
.
FromDevice
(
e_m_n_device_result
.
mData
.
data
());
e_device_buf
.
FromDevice
(
e_m_n_device_result
.
mData
.
data
());
if
(
init_method
==
6
||
init_method
==
7
)
{
std
::
cout
<<
std
::
fixed
<<
std
::
setprecision
(
16
);
float
a
=
ck
::
type_convert
<
float
>
(
a0_device_buf
(
0
,
10
));
float
b
=
ck
::
type_convert
<
float
>
(
b0_device_buf
(
0
,
10
));
std
::
cout
<<
"a(0,10): "
<<
a
<<
std
::
endl
;
std
::
cout
<<
"b(0,10): "
<<
b
<<
std
::
endl
;
std
::
cout
<<
"a: "
<<
ck
::
type_convert
<
float
>
(
a0_device_buf
(
0
,
0
))
<<
std
::
endl
;
std
::
cout
<<
"a: "
<<
ck
::
type_convert
<
float
>
(
a0_device_buf
(
0
,
1
))
<<
std
::
endl
;
std
::
cout
<<
"a: "
<<
ck
::
type_convert
<
float
>
(
a0_device_buf
(
0
,
2
))
<<
std
::
endl
;
std
::
cout
<<
"b: "
<<
ck
::
type_convert
<
float
>
(
b0_device_buf
(
0
,
0
))
<<
std
::
endl
;
std
::
cout
<<
"b: "
<<
ck
::
type_convert
<
float
>
(
b0_device_buf
(
1
,
0
))
<<
std
::
endl
;
std
::
cout
<<
"b: "
<<
ck
::
type_convert
<
float
>
(
b0_device_buf
(
2
,
0
))
<<
std
::
endl
;
float
d
=
ck
::
type_convert
<
float
>
(
e_m_n_device_result
(
0
,
10
));
float
h
=
ck
::
type_convert
<
float
>
(
e_m_n_host_result
(
10
,
0
));
std
::
cout
<<
"device result: "
<<
d
<<
std
::
endl
;
std
::
cout
<<
"host result: "
<<
h
<<
std
::
endl
;
std
::
cout
<<
"expected result: "
<<
M_PI
<<
std
::
endl
;
std
::
cout
<<
"device - host: "
<<
std
::
abs
(
d
-
h
)
<<
std
::
endl
;
std
::
cout
<<
"device - expected: "
<<
std
::
abs
(
d
-
M_PI
)
<<
std
::
endl
;
std
::
cout
<<
"atol: "
<<
5e-2
<<
std
::
endl
;
std
::
cout
<<
std
::
endl
<<
std
::
endl
;
}
return
ck
::
utils
::
check_err
(
return
ck
::
utils
::
check_err
(
e_m_n_device_result
,
e_m_n_host_result
,
"Error: Incorrect results!"
,
5e-2
,
5e-2
)
e_m_n_device_result
,
e_m_n_host_result
,
"Error: Incorrect results!"
,
5e-2
,
5e-2
)
?
0
?
0
...
...
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