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
yangql
composable_kernel-1
Commits
a414e3fd
Commit
a414e3fd
authored
Feb 15, 2019
by
Chao Liu
Browse files
update build
parent
67c6f73f
Changes
6
Hide whitespace changes
Inline
Side-by-side
Showing
6 changed files
with
17 additions
and
11 deletions
+17
-11
CMakeLists.txt
CMakeLists.txt
+2
-1
driver/CMakeLists.txt
driver/CMakeLists.txt
+5
-1
driver/driver.cu
driver/driver.cu
+1
-0
driver/driver.hip.cpp
driver/driver.hip.cpp
+7
-7
src/CMakeLists.txt
src/CMakeLists.txt
+1
-1
src/include/device.hpp
src/include/device.hpp
+1
-1
No files found.
CMakeLists.txt
View file @
a414e3fd
cmake_minimum_required
(
VERSION 2.8.3
)
#cmake_minimum_required(VERSION 2.8.3)
cmake_minimum_required
(
VERSION 3.9
)
project
(
modular_convolution
)
#c++
...
...
driver/CMakeLists.txt
View file @
a414e3fd
set
(
DRIVER_SOURCE driver.cpp
)
if
(
DEVICE_BACKEND STREQUAL
"HIP"
)
set
(
DRIVER_SOURCE driver.hip.cpp
)
elseif
(
DEVICE_BACKEND STREQUAL
"CUDA"
)
set
(
DRIVER_SOURCE driver.cu
)
endif
()
add_executable
(
driver
${
DRIVER_SOURCE
}
)
target_link_libraries
(
driver PRIVATE tensor
)
driver/driver.cu
0 → 120000
View file @
a414e3fd
driver
.
hip
.
cpp
\ No newline at end of file
driver/driver.cpp
→
driver/driver.
hip.
cpp
View file @
a414e3fd
...
...
@@ -50,7 +50,7 @@ struct GeneratorTensor_3
std
::
initializer_list
<
std
::
size_t
>
ids
=
{
static_cast
<
std
::
size_t
>
(
is
)...};
std
::
vector
<
std
::
size_t
>
lens
(
sizeof
...(
Is
),
100
);
std
::
vector
<
std
::
size_t
>
strides
(
sizeof
...(
Is
),
1
);
std
::
partial_sum
(
lens
.
rbegin
(),
lens
.
rbegin
()
+
(
sizeof
...(
Is
)
-
1
),
strides
.
rbegin
()
+
1
);
std
::
partial_sum
(
lens
.
rbegin
(),
lens
.
rbegin
()
+
(
sizeof
...(
Is
)
-
1
),
strides
.
rbegin
()
+
1
);
return
std
::
inner_product
(
ids
.
begin
(),
ids
.
end
(),
strides
.
begin
(),
std
::
size_t
(
0
))
+
1
;
#endif
}
...
...
@@ -340,7 +340,7 @@ void host_winograd_3x3_convolution(
std
::
size_t
ho
=
OutTileSizeH
*
y
+
j
;
for
(
int
i
=
0
;
i
<
OutTileSizeW
;
++
i
)
{
std
::
size_t
wo
=
OutTileSizeW
*
x
+
i
;
std
::
size_t
wo
=
OutTileSizeW
*
x
+
i
;
out
(
n
,
k
,
ho
,
wo
)
=
out_hold
(
n
,
k
,
y
,
x
,
j
,
i
);
}
}
...
...
@@ -393,13 +393,13 @@ int main()
constexpr unsigned WPad = 0;
#elif
0
// 3x3, 34x34
constexpr
unsigned
N
=
64
;
constexpr
unsigned
C
=
256
;
constexpr
unsigned
N
=
64
;
constexpr
unsigned
C
=
256
;
constexpr
unsigned
HI
=
34
;
constexpr
unsigned
WI
=
34
;
constexpr
unsigned
K
=
64
;
constexpr
unsigned
S
=
3
;
constexpr
unsigned
R
=
3
;
constexpr
unsigned
K
=
64
;
constexpr
unsigned
S
=
3
;
constexpr
unsigned
R
=
3
;
constexpr
unsigned
HPad
=
0
;
constexpr
unsigned
WPad
=
0
;
...
...
src/CMakeLists.txt
View file @
a414e3fd
...
...
@@ -11,7 +11,7 @@ set_target_properties(tensor PROPERTIES POSITION_INDEPENDENT_CODE ON)
if
(
DEVICE_BACKEND STREQUAL
"CUDA"
)
target_link_libraries
(
device
nvToolsExt cudart
)
target_link_libraries
(
tensor
nvToolsExt cudart
)
endif
()
install
(
TARGETS tensor LIBRARY DESTINATION lib
)
src/include/device.hpp
View file @
a414e3fd
...
...
@@ -43,7 +43,7 @@ float launch_kernel(F kernel, dim3 grid_dim, dim3 block_dim, Args... args)
hipGetErrorString
(
hipGetLastError
());
#elif DEVICE_BACKEND_CUDA
const
void
*
f
=
reinterpret_cast
<
const
void
*>
(
kernel
);
void
*
p_args
=
{
&
args
...};
void
*
p_args
[]
=
{
&
args
...};
timer
.
Start
();
...
...
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