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
d53443d5
Commit
d53443d5
authored
Sep 14, 2022
by
Po-Yen, Chen
Browse files
Use more meaningful names in permute bundle example
parent
7b135645
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
46 additions
and
32 deletions
+46
-32
example/37_permute/run_permute_bundle_example.inc
example/37_permute/run_permute_bundle_example.inc
+46
-32
No files found.
example/37_permute/run_permute_bundle_example.inc
View file @
d53443d5
...
@@ -9,39 +9,51 @@ bool run_permute_bundle(const Problem& problem)
...
@@ -9,39 +9,51 @@ bool run_permute_bundle(const Problem& problem)
using
std
::
begin
,
std
::
end
;
using
std
::
begin
,
std
::
end
;
const
auto
&
shape
=
problem
.
shape
;
const
auto
&
input_bundle_shape
=
problem
.
shape
;
ck
::
remove_cvref_t
<
decltype
(
shape
)
>
transposed_shape
;
const
auto
&
input_bundle_axes
=
problem
.
axes
;
transpose_shape
(
problem
.
shape
,
problem
.
axes
,
begin
(
transposed_shape
));
Tensor
<
BundleType
>
a
(
shape
);
ck
::
remove_cvref_t
<
decltype
(
input_bundle_shape
)
>
output_bundle_shape
;
Tensor
<
BundleType
>
b
(
transposed_shape
);
transpose_shape
(
input_bundle_shape
,
input_bundle_axes
,
begin
(
output_bundle_shape
));
Tensor
<
BundleType
>
input_bundle_tensor
(
input_bundle_shape
);
Tensor
<
BundleType
>
output_bundle_tensor
(
output_bundle_shape
);
// initialize tensor by assigning DataType values
// initialize tensor by assigning DataType values
using
std
::
data
,
std
::
size
;
using
std
::
data
,
std
::
size
;
ck
::
utils
::
FillUniformDistribution
<
DataType
>
{
-
1.
f
,
1.
f
}(
ck
::
span
<
DataType
>
{
ck
::
utils
::
FillUniformDistribution
<
DataType
>
{
-
1.
f
,
1.
f
}(
reinterpret_cast
<
DataType
*>
(
data
(
a
)),
a
.
GetElementSpaceSize
()
*
NumElemsInBundle
});
ck
::
span
<
DataType
>
{
reinterpret_cast
<
DataType
*>
(
data
(
input_bundle_tensor
)),
input_bundle_tensor
.
GetElementSpaceSize
()
*
NumElemsInBundle
});
DeviceMem
a
_device_buf
(
a
.
GetElementSpaceSizeInBytes
());
DeviceMem
input
_device_buf
(
input_bundle_tensor
.
GetElementSpaceSizeInBytes
());
DeviceMem
b
_device_buf
(
b
.
GetElementSpaceSizeInBytes
());
DeviceMem
output
_device_buf
(
output_bundle_tensor
.
GetElementSpaceSizeInBytes
());
a
_device_buf
.
ToDevice
(
data
(
a
));
input
_device_buf
.
ToDevice
(
data
(
input_bundle_tensor
));
std
::
array
<
ck
::
index_t
,
Problem
::
NumDim
>
a_lengths
,
b
_lengths
;
std
::
array
<
ck
::
index_t
,
Problem
::
NumDim
>
input_bundle_lengths
,
output_bundle
_lengths
;
std
::
array
<
ck
::
index_t
,
Problem
::
NumDim
>
a_strides
,
b
_strides
;
std
::
array
<
ck
::
index_t
,
Problem
::
NumDim
>
input_bundle_strides
,
output_bundle
_strides
;
const
void
*
input
=
a
_device_buf
.
GetDeviceBuffer
();
const
void
*
input
_bundle_data
=
input
_device_buf
.
GetDeviceBuffer
();
void
*
output
=
b
_device_buf
.
GetDeviceBuffer
();
void
*
output
_bundle_data
=
output
_device_buf
.
GetDeviceBuffer
();
std
::
copy
(
begin
(
shape
),
end
(
shape
),
begin
(
a_lengths
));
std
::
copy
(
begin
(
input_bundle_shape
),
end
(
input_bundle_shape
),
begin
(
input_bundle_lengths
));
std
::
copy
(
begin
(
a
.
GetStrides
()),
end
(
a
.
GetStrides
()),
begin
(
a_strides
));
std
::
copy
(
begin
(
input_bundle_tensor
.
GetStrides
()),
std
::
copy
(
begin
(
transposed_shape
),
end
(
transposed_shape
),
begin
(
b_lengths
));
end
(
input_bundle_tensor
.
GetStrides
()),
std
::
copy
(
begin
(
b
.
GetStrides
()),
end
(
b
.
GetStrides
()),
begin
(
b_strides
));
begin
(
input_bundle_strides
));
std
::
copy
(
begin
(
output_bundle_shape
),
end
(
output_bundle_shape
),
begin
(
output_bundle_lengths
));
std
::
copy
(
begin
(
output_bundle_tensor
.
GetStrides
()),
end
(
output_bundle_tensor
.
GetStrides
()),
begin
(
output_bundle_strides
));
static_assert
(
std
::
is_default_constructible_v
<
DevicePermuteInstance
>
);
static_assert
(
std
::
is_default_constructible_v
<
DevicePermuteInstance
>
);
auto
permute
=
DevicePermuteInstance
{};
auto
permute
=
DevicePermuteInstance
{};
auto
argument
=
permute
.
MakeArgument
(
auto
argument
=
permute
.
MakeArgument
(
input_bundle_lengths
,
a_lengths
,
a_strides
,
b_lengths
,
b_strides
,
input
,
output
,
PassThrough
{});
input_bundle_strides
,
output_bundle_lengths
,
output_bundle_strides
,
input_bundle_data
,
output_bundle_data
,
PassThrough
{});
if
(
!
permute
.
IsSupportedArgument
(
argument
))
if
(
!
permute
.
IsSupportedArgument
(
argument
))
{
{
...
@@ -55,28 +67,30 @@ bool run_permute_bundle(const Problem& problem)
...
@@ -55,28 +67,30 @@ bool run_permute_bundle(const Problem& problem)
std
::
cout
<<
"Perf: "
<<
ave_time
<<
" ms"
<<
std
::
endl
;
std
::
cout
<<
"Perf: "
<<
ave_time
<<
" ms"
<<
std
::
endl
;
b
_device_buf
.
FromDevice
(
data
(
b
));
output
_device_buf
.
FromDevice
(
data
(
output_bundle_tensor
));
// extend tensor shape from [N, H, W] to [N, H, W, NumElemsInBundle]
// extend tensor shape from [N, H, W] to [N, H, W, NumElemsInBundle]
const
auto
extended
_shape
=
extend_shape
(
shape
,
NumElemsInBundle
);
const
auto
input
_shape
=
extend_shape
(
input_bundle_
shape
,
NumElemsInBundle
);
const
auto
extended
_axes
=
extend_axes
(
problem
.
axes
);
const
auto
input
_axes
=
extend_axes
(
input_bundle_
axes
);
ck
::
remove_cvref_t
<
decltype
(
extended
_shape
)
>
transposed_extended
_shape
;
ck
::
remove_cvref_t
<
decltype
(
input
_shape
)
>
output
_shape
;
transpose_shape
(
extended_shape
,
extended_axes
,
begin
(
transposed_extended
_shape
));
transpose_shape
(
input_shape
,
input_axes
,
begin
(
output
_shape
));
Tensor
<
DataType
>
extended_a
(
extended_shape
);
Tensor
<
DataType
>
input_tensor
(
input_shape
);
std
::
memcpy
(
data
(
extended_a
),
data
(
a
),
a
.
GetElementSpaceSizeInBytes
());
std
::
memcpy
(
data
(
input_tensor
),
data
(
input_bundle_tensor
),
input_bundle_tensor
.
GetElementSpaceSizeInBytes
());
Tensor
<
DataType
>
extended_host_b
(
transposed_extended
_shape
);
Tensor
<
DataType
>
output_tensor
(
output
_shape
);
if
(
!
host_permute
(
extended_a
,
extended
_axes
,
PassThrough
{},
extended_host_b
))
if
(
!
host_permute
(
input_tensor
,
input
_axes
,
PassThrough
{},
output_tensor
))
{
{
return
false
;
return
false
;
}
}
return
ck
::
utils
::
check_err
(
return
ck
::
utils
::
check_err
(
ck
::
span
<
const
DataType
>
{
reinterpret_cast
<
DataType
*>
(
data
(
b
)),
ck
::
span
<
const
DataType
>
{
reinterpret_cast
<
DataType
*>
(
data
(
output_bundle_tensor
)),
b
.
GetElementSpaceSize
()
*
NumElemsInBundle
},
output_bundle_tensor
.
GetElementSpaceSize
()
*
NumElemsInBundle
},
ck
::
span
<
const
DataType
>
{
extended_host_b
.
mData
},
ck
::
span
<
const
DataType
>
{
output_tensor
},
"Error: incorrect results in output tensor"
,
"Error: incorrect results in output tensor"
,
1
e
-
6
,
1
e
-
6
,
1
e
-
6
);
1
e
-
6
);
...
...
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