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
a82804a4
Commit
a82804a4
authored
Nov 11, 2022
by
Anthony Chang
Browse files
helper function for transposing host tensor
parent
0345963e
Changes
4
Show whitespace changes
Inline
Side-by-side
Showing
4 changed files
with
132 additions
and
0 deletions
+132
-0
library/include/ck/library/utility/host_tensor.hpp
library/include/ck/library/utility/host_tensor.hpp
+23
-0
test/CMakeLists.txt
test/CMakeLists.txt
+1
-0
test/host_tensor/CMakeLists.txt
test/host_tensor/CMakeLists.txt
+2
-0
test/host_tensor/test_host_tensor.cpp
test/host_tensor/test_host_tensor.cpp
+106
-0
No files found.
library/include/ck/library/utility/host_tensor.hpp
View file @
a82804a4
...
@@ -433,6 +433,29 @@ struct Tensor
...
@@ -433,6 +433,29 @@ struct Tensor
return
mData
[
mDesc
.
GetOffsetFromMultiIndex
(
idx
)];
return
mData
[
mDesc
.
GetOffsetFromMultiIndex
(
idx
)];
}
}
Tensor
<
T
>
Transpose
(
std
::
vector
<
size_t
>
axes
=
{})
{
if
(
axes
.
empty
())
{
axes
.
resize
(
this
->
GetNumOfDimension
());
std
::
iota
(
axes
.
rbegin
(),
axes
.
rend
(),
0
);
}
if
(
axes
.
size
()
!=
mDesc
.
GetNumOfDimension
())
{
throw
std
::
runtime_error
(
"Tensor::Transpose(): size of axes must match tensor dimension"
);
}
std
::
vector
<
size_t
>
tlengths
,
tstrides
;
for
(
const
auto
&
axis
:
axes
)
{
tlengths
.
push_back
(
GetLengths
()[
axis
]);
tstrides
.
push_back
(
GetStrides
()[
axis
]);
}
Tensor
<
T
>
ret
(
*
this
);
ret
.
mDesc
=
HostTensorDescriptor
(
tlengths
,
tstrides
);
return
ret
;
}
typename
Data
::
iterator
begin
()
{
return
mData
.
begin
();
}
typename
Data
::
iterator
begin
()
{
return
mData
.
begin
();
}
typename
Data
::
iterator
end
()
{
return
mData
.
end
();
}
typename
Data
::
iterator
end
()
{
return
mData
.
end
();
}
...
...
test/CMakeLists.txt
View file @
a82804a4
...
@@ -58,3 +58,4 @@ add_subdirectory(batchnorm)
...
@@ -58,3 +58,4 @@ add_subdirectory(batchnorm)
if
(
GPU_TARGETS MATCHES
"gfx1100"
)
if
(
GPU_TARGETS MATCHES
"gfx1100"
)
add_subdirectory
(
wmma_op
)
add_subdirectory
(
wmma_op
)
endif
()
endif
()
add_subdirectory
(
host_tensor
)
test/host_tensor/CMakeLists.txt
0 → 100644
View file @
a82804a4
add_gtest_executable
(
test_host_tensor test_host_tensor.cpp
)
target_link_libraries
(
test_host_tensor PRIVATE utility
)
\ No newline at end of file
test/host_tensor/test_host_tensor.cpp
0 → 100644
View file @
a82804a4
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
#include <iostream>
#include <vector>
#include <gtest/gtest.h>
#include "ck/ck.hpp"
#include "ck/library/utility/host_tensor.hpp"
using
namespace
ck
;
TEST
(
HostTensorTranspose
,
TestBadArugment
)
{
Tensor
<
float
>
tensor
({
13
,
7
});
EXPECT_THROW
(
tensor
.
Transpose
({
0
}),
std
::
runtime_error
);
EXPECT_THROW
(
tensor
.
Transpose
({
0
,
1
,
2
}),
std
::
runtime_error
);
}
TEST
(
HostTensorTranspose
,
Test2D
)
{
std
::
vector
<
size_t
>
lengths
=
{
13
,
7
};
std
::
vector
<
size_t
>
tlengths
=
{
7
,
13
};
Tensor
<
float
>
tensor
(
lengths
);
tensor
(
0
,
0
)
=
0.
f
;
tensor
(
3
,
4
)
=
34.
f
;
EXPECT_EQ
(
tensor
.
GetLengths
(),
lengths
);
EXPECT_EQ
(
tensor
(
0
,
0
),
0.
f
);
EXPECT_EQ
(
tensor
(
3
,
4
),
34.
f
);
EXPECT_EQ
(
tensor
(
4
,
3
),
0.
f
);
EXPECT_EQ
(
tensor
.
Transpose
().
GetLengths
(),
tlengths
);
EXPECT_EQ
(
tensor
.
Transpose
()(
0
,
0
),
0.
f
);
EXPECT_EQ
(
tensor
.
Transpose
()(
4
,
3
),
34.
f
);
EXPECT_EQ
(
tensor
.
Transpose
()(
3
,
4
),
0.
f
);
}
TEST
(
HostTensorTranspose
,
Test3D
)
{
std
::
vector
<
size_t
>
lengths
=
{
13
,
7
,
5
};
std
::
vector
<
size_t
>
tlengths
=
{
5
,
7
,
13
};
Tensor
<
float
>
tensor
(
lengths
);
tensor
(
0
,
0
,
0
)
=
0.
f
;
tensor
(
3
,
4
,
2
)
=
342.
f
;
EXPECT_EQ
(
tensor
.
GetLengths
(),
lengths
);
EXPECT_EQ
(
tensor
(
0
,
0
,
0
),
0.
f
);
EXPECT_EQ
(
tensor
(
3
,
4
,
2
),
342.
f
);
EXPECT_EQ
(
tensor
(
4
,
3
,
2
),
0.
f
);
EXPECT_EQ
(
tensor
.
Transpose
().
GetLengths
(),
tlengths
);
EXPECT_EQ
(
tensor
.
Transpose
()(
0
,
0
,
0
),
0.
f
);
EXPECT_EQ
(
tensor
.
Transpose
()(
2
,
4
,
3
),
342.
f
);
EXPECT_EQ
(
tensor
.
Transpose
()(
2
,
3
,
4
),
0.
f
);
}
TEST
(
HostTensorTranspose
,
Test3D_021
)
{
std
::
vector
<
size_t
>
lengths
=
{
13
,
7
,
5
};
std
::
vector
<
size_t
>
tlengths
=
{
13
,
5
,
7
};
Tensor
<
float
>
tensor
(
lengths
);
tensor
(
0
,
0
,
0
)
=
0.
f
;
tensor
(
3
,
4
,
2
)
=
342.
f
;
EXPECT_EQ
(
tensor
.
GetLengths
(),
lengths
);
EXPECT_EQ
(
tensor
(
0
,
0
,
0
),
0.
f
);
EXPECT_EQ
(
tensor
(
3
,
4
,
2
),
342.
f
);
EXPECT_EQ
(
tensor
(
4
,
3
,
2
),
0.
f
);
// transpose last two dimensions
EXPECT_EQ
(
tensor
.
Transpose
({
0
,
2
,
1
}).
GetLengths
(),
tlengths
);
EXPECT_EQ
(
tensor
.
Transpose
({
0
,
2
,
1
})(
0
,
0
,
0
),
0.
f
);
EXPECT_EQ
(
tensor
.
Transpose
({
0
,
2
,
1
})(
2
,
4
,
3
),
0.
f
);
EXPECT_EQ
(
tensor
.
Transpose
({
0
,
2
,
1
})(
3
,
2
,
4
),
342.
f
);
EXPECT_EQ
(
tensor
.
Transpose
({
0
,
2
,
1
})(
2
,
3
,
4
),
0.
f
);
// transpose last two dimensions back again
EXPECT_EQ
(
tensor
.
Transpose
({
0
,
2
,
1
}).
Transpose
({
0
,
2
,
1
}).
GetLengths
(),
lengths
);
EXPECT_EQ
(
tensor
.
Transpose
({
0
,
2
,
1
}).
Transpose
({
0
,
2
,
1
})(
3
,
4
,
2
),
342.
f
);
}
TEST
(
HostTensorTranspose
,
TestNonpacked2D
)
{
std
::
vector
<
size_t
>
lengths
=
{
13
,
7
};
std
::
vector
<
size_t
>
strides
=
{
100
,
1
};
std
::
vector
<
size_t
>
tlengths
=
{
7
,
13
};
Tensor
<
float
>
tensor
(
lengths
,
strides
);
tensor
(
0
,
0
)
=
0.
f
;
tensor
(
3
,
4
)
=
34.
f
;
EXPECT_EQ
(
tensor
.
GetLengths
(),
lengths
);
EXPECT_EQ
(
tensor
(
0
,
0
),
0.
f
);
EXPECT_EQ
(
tensor
(
3
,
4
),
34.
f
);
EXPECT_EQ
(
tensor
(
4
,
3
),
0.
f
);
EXPECT_EQ
(
tensor
.
Transpose
().
GetLengths
(),
tlengths
);
EXPECT_EQ
(
tensor
.
Transpose
()(
0
,
0
),
0.
f
);
EXPECT_EQ
(
tensor
.
Transpose
()(
4
,
3
),
34.
f
);
EXPECT_EQ
(
tensor
.
Transpose
()(
3
,
4
),
0.
f
);
}
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