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
06fd9eaa
"library/src/host_tensor/device.cpp" did not exist on "1264925422920f24b3bb4fa34f178e31a23c97b5"
Commit
06fd9eaa
authored
Sep 20, 2023
by
Astha Rai
Browse files
updating formatting
parent
e73a2cb7
Changes
1
Show whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
23 additions
and
25 deletions
+23
-25
example/44_elementwise_permute/elementwise_permute_5D.cpp
example/44_elementwise_permute/elementwise_permute_5D.cpp
+23
-25
No files found.
example/44_elementwise_permute/elementwise_permute_5D.cpp
View file @
06fd9eaa
...
@@ -46,8 +46,6 @@ int main()
...
@@ -46,8 +46,6 @@ int main()
bool
do_verification
=
true
;
bool
do_verification
=
true
;
bool
time_kernel
=
true
;
bool
time_kernel
=
true
;
//std::vector<std::size_t> ncdhw = {16, 128, 32, 64, 16};
//std::vector<std::size_t> nchwd = {16, 128, 64, 16, 32};
std
::
vector
<
std
::
size_t
>
ncdhw
=
{
16
,
8
,
8
,
8
,
8
};
std
::
vector
<
std
::
size_t
>
ncdhw
=
{
16
,
8
,
8
,
8
,
8
};
std
::
vector
<
std
::
size_t
>
nchwd
=
{
16
,
8
,
8
,
8
,
8
};
std
::
vector
<
std
::
size_t
>
nchwd
=
{
16
,
8
,
8
,
8
,
8
};
Tensor
<
ADataType
>
a
(
ncdhw
);
Tensor
<
ADataType
>
a
(
ncdhw
);
...
@@ -64,12 +62,14 @@ int main()
...
@@ -64,12 +62,14 @@ int main()
std
::
array
<
void
*
,
1
>
output
=
{
b_device_buf
.
GetDeviceBuffer
()};
std
::
array
<
void
*
,
1
>
output
=
{
b_device_buf
.
GetDeviceBuffer
()};
std
::
array
<
ck
::
index_t
,
5
>
ab_lengths
;
std
::
array
<
ck
::
index_t
,
5
>
ab_lengths
;
std
::
array
<
ck
::
index_t
,
5
>
a_strides
=
{
static_cast
<
int
>
(
ncdhw
[
1
]
*
ncdhw
[
2
]
*
ncdhw
[
3
]
*
ncdhw
[
4
]),
std
::
array
<
ck
::
index_t
,
5
>
a_strides
=
{
static_cast
<
int
>
(
ncdhw
[
1
]
*
ncdhw
[
2
]
*
ncdhw
[
3
]
*
ncdhw
[
4
]),
static_cast
<
int
>
(
ncdhw
[
2
]
*
ncdhw
[
3
]
*
ncdhw
[
4
]),
static_cast
<
int
>
(
ncdhw
[
2
]
*
ncdhw
[
3
]
*
ncdhw
[
4
]),
static_cast
<
int
>
(
ncdhw
[
3
]
*
ncdhw
[
4
]),
static_cast
<
int
>
(
ncdhw
[
3
]
*
ncdhw
[
4
]),
static_cast
<
int
>
(
ncdhw
[
4
]),
static_cast
<
int
>
(
ncdhw
[
4
]),
1
};
1
};
std
::
array
<
ck
::
index_t
,
5
>
b_strides
=
{
static_cast
<
int
>
(
nchwd
[
1
]
*
nchwd
[
2
]
*
nchwd
[
3
]
*
nchwd
[
4
]),
std
::
array
<
ck
::
index_t
,
5
>
b_strides
=
{
static_cast
<
int
>
(
nchwd
[
1
]
*
nchwd
[
2
]
*
nchwd
[
3
]
*
nchwd
[
4
]),
static_cast
<
int
>
(
nchwd
[
2
]
*
nchwd
[
3
]
*
nchwd
[
4
]),
static_cast
<
int
>
(
nchwd
[
2
]
*
nchwd
[
3
]
*
nchwd
[
4
]),
1
,
1
,
static_cast
<
int
>
(
nchwd
[
3
]
*
nchwd
[
4
]),
static_cast
<
int
>
(
nchwd
[
3
]
*
nchwd
[
4
]),
...
@@ -95,17 +95,15 @@ int main()
...
@@ -95,17 +95,15 @@ int main()
broadcastPermute_invoker_ptr
->
Run
(
argument
.
get
(),
StreamConfig
{
nullptr
,
time_kernel
});
broadcastPermute_invoker_ptr
->
Run
(
argument
.
get
(),
StreamConfig
{
nullptr
,
time_kernel
});
std
::
size_t
flop
=
std
::
size_t
(
2
)
*
ncdhw
[
0
]
*
ncdhw
[
1
]
*
ncdhw
[
2
]
*
ncdhw
[
3
]
*
ncdhw
[
4
];
std
::
size_t
flop
=
std
::
size_t
(
2
)
*
ncdhw
[
0
]
*
ncdhw
[
1
]
*
ncdhw
[
2
]
*
ncdhw
[
3
]
*
ncdhw
[
4
];
std
::
size_t
num_btype
=
sizeof
(
ADataType
)
*
(
ncdhw
[
0
]
*
ncdhw
[
1
]
*
ncdhw
[
2
]
*
ncdhw
[
3
]
*
ncdhw
[
4
])
+
std
::
size_t
num_btype
=
sizeof
(
ADataType
)
*
(
ncdhw
[
0
]
*
ncdhw
[
1
]
*
ncdhw
[
2
]
*
ncdhw
[
3
]
*
ncdhw
[
4
])
+
sizeof
(
BDataType
)
*
(
ncdhw
[
0
]
*
ncdhw
[
1
]
*
ncdhw
[
2
]
*
ncdhw
[
3
]
*
ncdhw
[
4
]);
sizeof
(
BDataType
)
*
(
ncdhw
[
0
]
*
ncdhw
[
1
]
*
ncdhw
[
2
]
*
ncdhw
[
3
]
*
ncdhw
[
4
]);
float
tflops
=
static_cast
<
float
>
(
flop
)
/
1.E9
/
ave_time
;
float
tflops
=
static_cast
<
float
>
(
flop
)
/
1.E9
/
ave_time
;
float
gb_per_sec
=
num_btype
/
1.E6
/
ave_time
;
float
gb_per_sec
=
num_btype
/
1.E6
/
ave_time
;
//LogRangeAsType<float>(std::cout << "A : ", a.mData, ",") << std::endl;
// LogRangeAsType<float>(std::cout << "A : ", a.mData, ",") << std::endl;
//LogRangeAsType<float>(std::cout << "B : ", b.mData, ",") << std::endl;
//std::cout << "A: " << a.mData.data() << std::endl;
std
::
cout
<<
"Perf: "
<<
ave_time
<<
" ms, "
<<
tflops
<<
" TFlops, "
<<
gb_per_sec
<<
" GB/s"
std
::
cout
<<
"Perf: "
<<
ave_time
<<
" ms, "
<<
tflops
<<
" TFlops, "
<<
gb_per_sec
<<
" GB/s"
<<
std
::
endl
;
<<
std
::
endl
;
...
@@ -117,8 +115,8 @@ int main()
...
@@ -117,8 +115,8 @@ int main()
Tensor
<
BDataType
>
host_b
(
nchwd
);
Tensor
<
BDataType
>
host_b
(
nchwd
);
host_elementwise4D
(
host_b
,
a
,
PassThrough
{});
host_elementwise4D
(
host_b
,
a
,
PassThrough
{});
//LogRangeAsType<float>(std::cout << "B : ", b.mData, ",") << std::endl;
//
LogRangeAsType<float>(std::cout << "B : ", b.mData, ",") << std::endl;
//LogRangeAsType<float>(std::cout << "Host B : ", host_b.mData, ",") << std::endl;
//
LogRangeAsType<float>(std::cout << "Host B : ", host_b.mData, ",") << std::endl;
pass
&=
pass
&=
ck
::
utils
::
check_err
(
b
.
mData
,
host_b
.
mData
,
"Error: Incorrect results b"
,
1e-3
,
1e-3
);
ck
::
utils
::
check_err
(
b
.
mData
,
host_b
.
mData
,
"Error: Incorrect results b"
,
1e-3
,
1e-3
);
...
...
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