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
762647e3
Commit
762647e3
authored
Jul 31, 2024
by
Astha Rai
Browse files
removed commented code
parent
a150da5a
Changes
6
Hide whitespace changes
Inline
Side-by-side
Showing
6 changed files
with
4 additions
and
31 deletions
+4
-31
codegen/CMakeLists.txt
codegen/CMakeLists.txt
+2
-2
codegen/src/device_grouped_conv_fwd_multiple_abd_operation_xdl_cshuffle.cpp
..._grouped_conv_fwd_multiple_abd_operation_xdl_cshuffle.cpp
+2
-21
codegen/test/grouped_conv_fwd_multiple_d_v1.cpp
codegen/test/grouped_conv_fwd_multiple_d_v1.cpp
+0
-2
codegen/test/grouped_conv_fwd_multiple_d_v2.cpp
codegen/test/grouped_conv_fwd_multiple_d_v2.cpp
+0
-2
codegen/test/grouped_conv_fwd_multiple_d_v3.cpp
codegen/test/grouped_conv_fwd_multiple_d_v3.cpp
+0
-2
codegen/test/grouped_conv_fwd_multiple_d_v4.cpp
codegen/test/grouped_conv_fwd_multiple_d_v4.cpp
+0
-2
No files found.
codegen/CMakeLists.txt
View file @
762647e3
...
@@ -33,8 +33,8 @@ list(APPEND CMAKE_MODULE_PATH ${CK_ROOT}/cmake)
...
@@ -33,8 +33,8 @@ list(APPEND CMAKE_MODULE_PATH ${CK_ROOT}/cmake)
include
(
Embed
)
include
(
Embed
)
file
(
GLOB_RECURSE KERNEL_FILES CONFIGURE_DEPENDS
file
(
GLOB_RECURSE KERNEL_FILES CONFIGURE_DEPENDS
${
CK_ROOT
}
/include/ck/*.hpp
)
${
CK_ROOT
}
/include/ck/*.hpp
)
message
(
STATUS
"KERNEL_FILES:
${
KERNEL_FILES
}
"
)
#
message(STATUS "KERNEL_FILES: ${KERNEL_FILES}")
message
(
STATUS
"RELATIVE:
${
CK_ROOT
}
/include"
)
#
message(STATUS "RELATIVE: ${CK_ROOT}/include")
add_embed_library
(
ck_headers
${
KERNEL_FILES
}
RELATIVE
${
CK_ROOT
}
/include
)
add_embed_library
(
ck_headers
${
KERNEL_FILES
}
RELATIVE
${
CK_ROOT
}
/include
)
file
(
GLOB SOURCES CONFIGURE_DEPENDS src/*.cpp
)
file
(
GLOB SOURCES CONFIGURE_DEPENDS src/*.cpp
)
...
...
codegen/src/device_grouped_conv_fwd_multiple_abd_operation_xdl_cshuffle.cpp
View file @
762647e3
...
@@ -12,27 +12,8 @@ namespace ck {
...
@@ -12,27 +12,8 @@ namespace ck {
namespace
host
{
namespace
host
{
namespace
conv
{
namespace
conv
{
// calculate appropriate Gemm Specification based on input tensor dimensions
// NOTE: in CK, MNKPadding is always used for forward convolution, so didn't
// NOTE: in CK, MNKPadding is always used for forward convolution
// add GemmSpec function here
/**static std::string GetGemmSpec(const std::size_t m,
const std::size_t n,
const std::size_t k,
const std::size_t m_per_block,
const std::size_t n_per_block,
const std::size_t k_per_block)
{
std::string spec = "";
if(integer_divide_ceil(m, m_per_block) * m_per_block - m != 0)
spec += "M";
if(integer_divide_ceil(n, n_per_block) * n_per_block - n != 0)
spec += "N";
if(integer_divide_ceil(k, k_per_block) * k_per_block - k != 0)
spec += "K";
if(spec == "")
return "ck::tensor_operation::device::GemmSpecialization::Default";
return "ck::tensor_operation::device::GemmSpecialization::" + spec + "Padding";
}**/
// function to update prologue/epilogue with user provided operation
// function to update prologue/epilogue with user provided operation
void
Operation_Conv_Fwd_Xdl_Cshuffle
::
update_prologue
(
const
std
::
string
&
pro
)
void
Operation_Conv_Fwd_Xdl_Cshuffle
::
update_prologue
(
const
std
::
string
&
pro
)
...
...
codegen/test/grouped_conv_fwd_multiple_d_v1.cpp
View file @
762647e3
...
@@ -92,7 +92,6 @@ struct Epilogue
...
@@ -92,7 +92,6 @@ struct Epilogue
static_cast
<
int
>
(
prob
.
C
),
static_cast
<
int
>
(
prob
.
C
),
static_cast
<
int
>
(
prob
.
Y
),
static_cast
<
int
>
(
prob
.
Y
),
static_cast
<
int
>
(
prob
.
X
)};
static_cast
<
int
>
(
prob
.
X
)};
// ck::Array<ck::index_t, 5> d_lengths = {};
ck
::
Array
<
ck
::
index_t
,
5
>
in_strides
{
static_cast
<
int
>
(
prob
.
C
),
ck
::
Array
<
ck
::
index_t
,
5
>
in_strides
{
static_cast
<
int
>
(
prob
.
C
),
static_cast
<
int
>
(
prob
.
Hi
*
prob
.
Wi
*
prob
.
G
*
prob
.
C
),
static_cast
<
int
>
(
prob
.
Hi
*
prob
.
Wi
*
prob
.
G
*
prob
.
C
),
...
@@ -109,7 +108,6 @@ struct Epilogue
...
@@ -109,7 +108,6 @@ struct Epilogue
1
,
1
,
static_cast
<
int
>
(
prob
.
X
*
prob
.
C
),
static_cast
<
int
>
(
prob
.
X
*
prob
.
C
),
static_cast
<
int
>
(
prob
.
C
)};
static_cast
<
int
>
(
prob
.
C
)};
// ck::Array<ck::index_t, 5> d_strides = {};
ck
::
Array
<
ck
::
index_t
,
2
>
conv_filter_strides
=
{
2
,
2
};
ck
::
Array
<
ck
::
index_t
,
2
>
conv_filter_strides
=
{
2
,
2
};
ck
::
Array
<
ck
::
index_t
,
2
>
conv_filter_dilations
=
{
1
,
1
};
ck
::
Array
<
ck
::
index_t
,
2
>
conv_filter_dilations
=
{
1
,
1
};
...
...
codegen/test/grouped_conv_fwd_multiple_d_v2.cpp
View file @
762647e3
...
@@ -92,7 +92,6 @@ struct Epilogue
...
@@ -92,7 +92,6 @@ struct Epilogue
static_cast
<
int
>
(
prob
.
C
),
static_cast
<
int
>
(
prob
.
C
),
static_cast
<
int
>
(
prob
.
Y
),
static_cast
<
int
>
(
prob
.
Y
),
static_cast
<
int
>
(
prob
.
X
)};
static_cast
<
int
>
(
prob
.
X
)};
// ck::Array<ck::index_t, 5> d_lengths = {};
ck
::
Array
<
ck
::
index_t
,
5
>
in_strides
{
static_cast
<
int
>
(
prob
.
C
),
ck
::
Array
<
ck
::
index_t
,
5
>
in_strides
{
static_cast
<
int
>
(
prob
.
C
),
static_cast
<
int
>
(
prob
.
Hi
*
prob
.
Wi
*
prob
.
G
*
prob
.
C
),
static_cast
<
int
>
(
prob
.
Hi
*
prob
.
Wi
*
prob
.
G
*
prob
.
C
),
...
@@ -109,7 +108,6 @@ struct Epilogue
...
@@ -109,7 +108,6 @@ struct Epilogue
1
,
1
,
static_cast
<
int
>
(
prob
.
X
*
prob
.
C
),
static_cast
<
int
>
(
prob
.
X
*
prob
.
C
),
static_cast
<
int
>
(
prob
.
C
)};
static_cast
<
int
>
(
prob
.
C
)};
// ck::Array<ck::index_t, 5> d_strides = {};
ck
::
Array
<
ck
::
index_t
,
2
>
conv_filter_strides
=
{
1
,
1
};
ck
::
Array
<
ck
::
index_t
,
2
>
conv_filter_strides
=
{
1
,
1
};
ck
::
Array
<
ck
::
index_t
,
2
>
conv_filter_dilations
=
{
1
,
1
};
ck
::
Array
<
ck
::
index_t
,
2
>
conv_filter_dilations
=
{
1
,
1
};
...
...
codegen/test/grouped_conv_fwd_multiple_d_v3.cpp
View file @
762647e3
...
@@ -92,7 +92,6 @@ struct Epilogue
...
@@ -92,7 +92,6 @@ struct Epilogue
static_cast
<
int
>
(
prob
.
C
),
static_cast
<
int
>
(
prob
.
C
),
static_cast
<
int
>
(
prob
.
Y
),
static_cast
<
int
>
(
prob
.
Y
),
static_cast
<
int
>
(
prob
.
X
)};
static_cast
<
int
>
(
prob
.
X
)};
// ck::Array<ck::index_t, 5> d_lengths = {};
ck
::
Array
<
ck
::
index_t
,
5
>
in_strides
{
static_cast
<
int
>
(
prob
.
C
),
ck
::
Array
<
ck
::
index_t
,
5
>
in_strides
{
static_cast
<
int
>
(
prob
.
C
),
static_cast
<
int
>
(
prob
.
Hi
*
prob
.
Wi
*
prob
.
G
*
prob
.
C
),
static_cast
<
int
>
(
prob
.
Hi
*
prob
.
Wi
*
prob
.
G
*
prob
.
C
),
...
@@ -109,7 +108,6 @@ struct Epilogue
...
@@ -109,7 +108,6 @@ struct Epilogue
1
,
1
,
static_cast
<
int
>
(
prob
.
X
*
prob
.
C
),
static_cast
<
int
>
(
prob
.
X
*
prob
.
C
),
static_cast
<
int
>
(
prob
.
C
)};
static_cast
<
int
>
(
prob
.
C
)};
// ck::Array<ck::index_t, 5> d_strides = {};
ck
::
Array
<
ck
::
index_t
,
2
>
conv_filter_strides
=
{
2
,
2
};
ck
::
Array
<
ck
::
index_t
,
2
>
conv_filter_strides
=
{
2
,
2
};
ck
::
Array
<
ck
::
index_t
,
2
>
conv_filter_dilations
=
{
1
,
1
};
ck
::
Array
<
ck
::
index_t
,
2
>
conv_filter_dilations
=
{
1
,
1
};
...
...
codegen/test/grouped_conv_fwd_multiple_d_v4.cpp
View file @
762647e3
...
@@ -92,7 +92,6 @@ struct Epilogue
...
@@ -92,7 +92,6 @@ struct Epilogue
static_cast
<
int
>
(
prob
.
C
),
static_cast
<
int
>
(
prob
.
C
),
static_cast
<
int
>
(
prob
.
Y
),
static_cast
<
int
>
(
prob
.
Y
),
static_cast
<
int
>
(
prob
.
X
)};
static_cast
<
int
>
(
prob
.
X
)};
// ck::Array<ck::index_t, 5> d_lengths = {};
ck
::
Array
<
ck
::
index_t
,
5
>
in_strides
{
static_cast
<
int
>
(
prob
.
C
),
ck
::
Array
<
ck
::
index_t
,
5
>
in_strides
{
static_cast
<
int
>
(
prob
.
C
),
static_cast
<
int
>
(
prob
.
Hi
*
prob
.
Wi
*
prob
.
G
*
prob
.
C
),
static_cast
<
int
>
(
prob
.
Hi
*
prob
.
Wi
*
prob
.
G
*
prob
.
C
),
...
@@ -109,7 +108,6 @@ struct Epilogue
...
@@ -109,7 +108,6 @@ struct Epilogue
1
,
1
,
static_cast
<
int
>
(
prob
.
X
*
prob
.
C
),
static_cast
<
int
>
(
prob
.
X
*
prob
.
C
),
static_cast
<
int
>
(
prob
.
C
)};
static_cast
<
int
>
(
prob
.
C
)};
// ck::Array<ck::index_t, 5> d_strides = {};
ck
::
Array
<
ck
::
index_t
,
2
>
conv_filter_strides
=
{
1
,
1
};
ck
::
Array
<
ck
::
index_t
,
2
>
conv_filter_strides
=
{
1
,
1
};
ck
::
Array
<
ck
::
index_t
,
2
>
conv_filter_dilations
=
{
1
,
1
};
ck
::
Array
<
ck
::
index_t
,
2
>
conv_filter_dilations
=
{
1
,
1
};
...
...
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