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
a150da5a
Commit
a150da5a
authored
Jul 31, 2024
by
Astha Rai
Browse files
cleaned up compiler errors
parent
84c65962
Changes
10
Hide whitespace changes
Inline
Side-by-side
Showing
10 changed files
with
32 additions
and
24 deletions
+32
-24
codegen/include/ck/host/types.hpp
codegen/include/ck/host/types.hpp
+3
-0
codegen/src/device_gemm_multiple_d_operation_xdl_cshuffle.cpp
...gen/src/device_gemm_multiple_d_operation_xdl_cshuffle.cpp
+7
-6
codegen/src/device_grouped_conv_fwd_multiple_abd_operation_xdl_cshuffle.cpp
..._grouped_conv_fwd_multiple_abd_operation_xdl_cshuffle.cpp
+9
-8
codegen/src/headers.cpp
codegen/src/headers.cpp
+3
-0
codegen/test/grouped_conv_fwd_multiple_d_v1.cpp
codegen/test/grouped_conv_fwd_multiple_d_v1.cpp
+2
-2
codegen/test/grouped_conv_fwd_multiple_d_v2.cpp
codegen/test/grouped_conv_fwd_multiple_d_v2.cpp
+2
-2
codegen/test/grouped_conv_fwd_multiple_d_v3.cpp
codegen/test/grouped_conv_fwd_multiple_d_v3.cpp
+2
-2
codegen/test/grouped_conv_fwd_multiple_d_v4.cpp
codegen/test/grouped_conv_fwd_multiple_d_v4.cpp
+2
-2
codegen/test/rtc/src/kernel.cpp
codegen/test/rtc/src/kernel.cpp
+1
-1
codegen/test/rtc/src/tmp_dir.cpp
codegen/test/rtc/src/tmp_dir.cpp
+1
-1
No files found.
codegen/include/ck/host/types.hpp
View file @
a150da5a
...
...
@@ -76,8 +76,11 @@ std::string SequenceStr(const std::vector<int>& v);
std
::
string
MakeTuple
(
const
std
::
vector
<
std
::
string
>&
v
);
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wglobal-constructors"
template
<
int
...
xs
>
const
std
::
string
S
=
SequenceStr
({
xs
...});
#pragma clang diagnostic pop
constexpr
const
char
*
PassThrough
=
"ck::tensor_operation::element_wise::PassThrough"
;
constexpr
const
char
*
Bilinear
=
"ck::tensor_operation::element_wise::Bilinear"
;
...
...
codegen/src/device_gemm_multiple_d_operation_xdl_cshuffle.cpp
View file @
a150da5a
...
...
@@ -3,6 +3,7 @@
#include "ck/host/device_gemm_multiple_d/operation.hpp"
#include "ck/host/stringutils.hpp"
#include "ck/host/types.hpp"
#include "ck/host/utils.hpp"
#include <cassert>
...
...
@@ -32,11 +33,11 @@ static std::string GetGemmSpec(const std::size_t m,
}
// function to update prologue/epilogue with user provided operation
void
Operation_Xdl_CShuffle
::
update_prologue
(
const
std
::
string
&
pro
logue
)
void
Operation_Xdl_CShuffle
::
update_prologue
(
const
std
::
string
&
pro
)
{
if
(
!
pro
logue
.
empty
())
if
(
!
pro
.
empty
())
{
this
->
prologue
=
pro
logue
;
this
->
prologue
=
pro
;
this
->
cde_elem_op
=
"CDEElementOp"
;
}
else
...
...
@@ -45,11 +46,11 @@ void Operation_Xdl_CShuffle::update_prologue(const std::string& prologue)
}
}
void
Operation_Xdl_CShuffle
::
update_epilogue
(
const
std
::
string
&
epi
logue
)
void
Operation_Xdl_CShuffle
::
update_epilogue
(
const
std
::
string
&
epi
)
{
if
(
!
epi
logue
.
empty
())
if
(
!
epi
.
empty
())
{
this
->
epilogue
=
epi
logue
;
this
->
epilogue
=
epi
;
this
->
cde_elem_op
=
"CDEElementOp"
;
}
else
...
...
codegen/src/device_grouped_conv_fwd_multiple_abd_operation_xdl_cshuffle.cpp
View file @
a150da5a
...
...
@@ -4,6 +4,7 @@
#include "ck/host/device_grouped_conv_fwd_multiple_d/conv_fwd_op.hpp"
#include <iostream>
#include "ck/host/stringutils.hpp"
#include "ck/host/types.hpp"
#include "ck/host/utils.hpp"
#include <cassert>
...
...
@@ -13,7 +14,7 @@ namespace conv {
// calculate appropriate Gemm Specification based on input tensor dimensions
// NOTE: in CK, MNKPadding is always used for forward convolution
static
std
::
string
GetGemmSpec
(
const
std
::
size_t
m
,
/**
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,
...
...
@@ -31,14 +32,14 @@ static std::string GetGemmSpec(const std::size_t m,
return "ck::tensor_operation::device::GemmSpecialization::Default";
return "ck::tensor_operation::device::GemmSpecialization::" + spec + "Padding";
}
}
**/
// function to update prologue/epilogue with user provided operation
void
Operation_Conv_Fwd_Xdl_Cshuffle
::
update_prologue
(
const
std
::
string
&
pro
logue
)
void
Operation_Conv_Fwd_Xdl_Cshuffle
::
update_prologue
(
const
std
::
string
&
pro
)
{
if
(
!
pro
logue
.
empty
())
if
(
!
pro
.
empty
())
{
this
->
prologue
=
pro
logue
;
this
->
prologue
=
pro
;
this
->
cde_elem_op
=
"CDEElementOp"
;
}
else
...
...
@@ -47,11 +48,11 @@ void Operation_Conv_Fwd_Xdl_Cshuffle::update_prologue(const std::string& prologu
}
}
void
Operation_Conv_Fwd_Xdl_Cshuffle
::
update_epilogue
(
const
std
::
string
&
epi
logue
)
void
Operation_Conv_Fwd_Xdl_Cshuffle
::
update_epilogue
(
const
std
::
string
&
epi
)
{
if
(
!
epi
logue
.
empty
())
if
(
!
epi
.
empty
())
{
this
->
epilogue
=
epi
logue
;
this
->
epilogue
=
epi
;
this
->
cde_elem_op
=
"CDEElementOp"
;
}
else
...
...
codegen/src/headers.cpp
View file @
a150da5a
...
...
@@ -4,7 +4,10 @@
namespace
ck
{
namespace
host
{
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wglobal-constructors"
const
std
::
string
config_header
=
""
;
#pragma clang diagnostic pop
std
::
unordered_map
<
std
::
string_view
,
std
::
string_view
>
GetHeaders
()
{
...
...
codegen/test/grouped_conv_fwd_multiple_d_v1.cpp
View file @
a150da5a
...
...
@@ -92,7 +92,7 @@ struct Epilogue
static_cast
<
int
>
(
prob
.
C
),
static_cast
<
int
>
(
prob
.
Y
),
static_cast
<
int
>
(
prob
.
X
)};
ck
::
Array
<
ck
::
index_t
,
5
>
d_lengths
=
{};
//
ck::Array<ck::index_t, 5> d_lengths = {};
ck
::
Array
<
ck
::
index_t
,
5
>
in_strides
{
static_cast
<
int
>
(
prob
.
C
),
static_cast
<
int
>
(
prob
.
Hi
*
prob
.
Wi
*
prob
.
G
*
prob
.
C
),
...
...
@@ -109,7 +109,7 @@ struct Epilogue
1
,
static_cast
<
int
>
(
prob
.
X
*
prob
.
C
),
static_cast
<
int
>
(
prob
.
C
)};
ck
::
Array
<
ck
::
index_t
,
5
>
d_strides
=
{};
//
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_dilations
=
{
1
,
1
};
...
...
codegen/test/grouped_conv_fwd_multiple_d_v2.cpp
View file @
a150da5a
...
...
@@ -92,7 +92,7 @@ struct Epilogue
static_cast
<
int
>
(
prob
.
C
),
static_cast
<
int
>
(
prob
.
Y
),
static_cast
<
int
>
(
prob
.
X
)};
ck
::
Array
<
ck
::
index_t
,
5
>
d_lengths
=
{};
//
ck::Array<ck::index_t, 5> d_lengths = {};
ck
::
Array
<
ck
::
index_t
,
5
>
in_strides
{
static_cast
<
int
>
(
prob
.
C
),
static_cast
<
int
>
(
prob
.
Hi
*
prob
.
Wi
*
prob
.
G
*
prob
.
C
),
...
...
@@ -109,7 +109,7 @@ struct Epilogue
1
,
static_cast
<
int
>
(
prob
.
X
*
prob
.
C
),
static_cast
<
int
>
(
prob
.
C
)};
ck
::
Array
<
ck
::
index_t
,
5
>
d_strides
=
{};
//
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_dilations
=
{
1
,
1
};
...
...
codegen/test/grouped_conv_fwd_multiple_d_v3.cpp
View file @
a150da5a
...
...
@@ -92,7 +92,7 @@ struct Epilogue
static_cast
<
int
>
(
prob
.
C
),
static_cast
<
int
>
(
prob
.
Y
),
static_cast
<
int
>
(
prob
.
X
)};
ck
::
Array
<
ck
::
index_t
,
5
>
d_lengths
=
{};
//
ck::Array<ck::index_t, 5> d_lengths = {};
ck
::
Array
<
ck
::
index_t
,
5
>
in_strides
{
static_cast
<
int
>
(
prob
.
C
),
static_cast
<
int
>
(
prob
.
Hi
*
prob
.
Wi
*
prob
.
G
*
prob
.
C
),
...
...
@@ -109,7 +109,7 @@ struct Epilogue
1
,
static_cast
<
int
>
(
prob
.
X
*
prob
.
C
),
static_cast
<
int
>
(
prob
.
C
)};
ck
::
Array
<
ck
::
index_t
,
5
>
d_strides
=
{};
//
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_dilations
=
{
1
,
1
};
...
...
codegen/test/grouped_conv_fwd_multiple_d_v4.cpp
View file @
a150da5a
...
...
@@ -92,7 +92,7 @@ struct Epilogue
static_cast
<
int
>
(
prob
.
C
),
static_cast
<
int
>
(
prob
.
Y
),
static_cast
<
int
>
(
prob
.
X
)};
ck
::
Array
<
ck
::
index_t
,
5
>
d_lengths
=
{};
//
ck::Array<ck::index_t, 5> d_lengths = {};
ck
::
Array
<
ck
::
index_t
,
5
>
in_strides
{
static_cast
<
int
>
(
prob
.
C
),
static_cast
<
int
>
(
prob
.
Hi
*
prob
.
Wi
*
prob
.
G
*
prob
.
C
),
...
...
@@ -109,7 +109,7 @@ struct Epilogue
1
,
static_cast
<
int
>
(
prob
.
X
*
prob
.
C
),
static_cast
<
int
>
(
prob
.
C
)};
ck
::
Array
<
ck
::
index_t
,
5
>
d_strides
=
{};
//
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_dilations
=
{
1
,
1
};
...
...
codegen/test/rtc/src/kernel.cpp
View file @
a150da5a
...
...
@@ -118,4 +118,4 @@ void kernel::launch(hipStream_t stream,
launch_kernel
(
impl
->
fun
,
stream
,
global
,
local
,
kernargs
.
data
(),
size
);
}
}
// namespace rtc
\ No newline at end of file
}
// namespace rtc
codegen/test/rtc/src/tmp_dir.cpp
View file @
a150da5a
...
...
@@ -45,4 +45,4 @@ void tmp_dir::execute(const std::string& cmd) const
tmp_dir
::~
tmp_dir
()
{
std
::
filesystem
::
remove_all
(
this
->
path
);
}
}
// namespace rtc
\ No newline at end of file
}
// namespace rtc
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