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
6790b8f3
Commit
6790b8f3
authored
Nov 20, 2018
by
Chao Liu
Browse files
rename
parent
d2a488dd
Changes
7
Show whitespace changes
Inline
Side-by-side
Showing
7 changed files
with
24 additions
and
23 deletions
+24
-23
driver/conv.cu
driver/conv.cu
+3
-2
driver/device_direct_convolution_1.cuh
driver/device_direct_convolution_1.cuh
+5
-12
driver/device_direct_convolution_2.cuh
driver/device_direct_convolution_2.cuh
+12
-5
src/include/blockwise_direct_convolution.cuh
src/include/blockwise_direct_convolution.cuh
+1
-1
src/include/gridwise_direct_convolution_1.cuh
src/include/gridwise_direct_convolution_1.cuh
+1
-1
src/include/gridwise_direct_convolution_2.cuh
src/include/gridwise_direct_convolution_2.cuh
+2
-2
src/include/threadwise_direct_convolution.cuh
src/include/threadwise_direct_convolution.cuh
+0
-0
No files found.
driver/conv.cu
View file @
6790b8f3
...
@@ -5,7 +5,8 @@
...
@@ -5,7 +5,8 @@
#include "nvToolsExt.h"
#include "nvToolsExt.h"
#include "tensor.hpp"
#include "tensor.hpp"
#include "constant_tensor_descriptor.cuh"
#include "constant_tensor_descriptor.cuh"
#include "device_direct_convolution_3.cuh"
#include "device_direct_convolution_1.cuh"
#include "device_direct_convolution_2.cuh"
template
<
class
T
>
template
<
class
T
>
struct
GeneratorConstant
struct
GeneratorConstant
...
@@ -177,7 +178,7 @@ int main()
...
@@ -177,7 +178,7 @@ int main()
for
(
int
i
=
0
;
i
<
20
;
++
i
)
for
(
int
i
=
0
;
i
<
20
;
++
i
)
{
{
device_convolution
(
in_desc
,
in
,
wei_desc
,
wei
,
out_desc
,
out_device
);
device_
direct_
convolution
_2
(
in_desc
,
in
,
wei_desc
,
wei
,
out_desc
,
out_device
);
}
}
#if 0
#if 0
...
...
driver/device_direct_convolution_
3
.cuh
→
driver/device_direct_convolution_
1
.cuh
View file @
6790b8f3
#pragma once
#pragma once
#include "direct_convolution_
3
.cuh"
#include "
gridwise_
direct_convolution_
1
.cuh"
template
<
class
T
,
class
InDesc
,
class
WeiDesc
,
class
OutDesc
>
template
<
class
T
,
class
InDesc
,
class
WeiDesc
,
class
OutDesc
>
void
device_convolution
(
void
device_
direct_
convolution
_1
(
InDesc
,
const
Tensor
<
T
>&
in
,
WeiDesc
,
const
Tensor
<
T
>&
wei
,
OutDesc
,
Tensor
<
T
>&
out
)
InDesc
,
const
Tensor
<
T
>&
in
,
WeiDesc
,
const
Tensor
<
T
>&
wei
,
OutDesc
,
Tensor
<
T
>&
out
)
{
{
std
::
size_t
data_sz
=
sizeof
(
T
);
std
::
size_t
data_sz
=
sizeof
(
T
);
...
@@ -26,16 +26,12 @@ void device_convolution(
...
@@ -26,16 +26,12 @@ void device_convolution(
constexpr
auto
out_desc
=
OutDesc
{};
constexpr
auto
out_desc
=
OutDesc
{};
constexpr
unsigned
OutTileSizeH
=
2
;
constexpr
unsigned
OutTileSizeH
=
2
;
constexpr
unsigned
OutTileSizeW
=
2
;
constexpr
unsigned
OutTileSizeW
=
2
;
constexpr
unsigned
NPerBlock
=
2
;
constexpr
unsigned
NPerBlock
=
1
;
constexpr
unsigned
KPerBlock
=
32
;
constexpr
unsigned
KPerBlock
=
4
;
constexpr
unsigned
CPerBlock
=
2
;
constexpr
unsigned
CPerBlock
=
2
;
constexpr
unsigned
YPerBlock
=
1
;
constexpr
unsigned
YPerBlock
=
8
;
constexpr
unsigned
XPerBlock
=
16
;
constexpr
unsigned
XPerBlock
=
16
;
constexpr
unsigned
NPerThread
=
2
;
constexpr
unsigned
KPerThread
=
4
;
constexpr
unsigned
CPerThread
=
2
;
constexpr
unsigned
NBlockOpLen0
=
1
;
constexpr
unsigned
NBlockOpLen0
=
1
;
constexpr
unsigned
NBlockOpLen1
=
1
;
constexpr
unsigned
NBlockOpLen1
=
1
;
constexpr
unsigned
NBlockOpLen2
=
4
;
constexpr
unsigned
NBlockOpLen2
=
4
;
...
@@ -70,9 +66,6 @@ void device_convolution(
...
@@ -70,9 +66,6 @@ void device_convolution(
CPerBlock
,
CPerBlock
,
YPerBlock
,
YPerBlock
,
XPerBlock
,
XPerBlock
,
NPerThread
,
KPerThread
,
CPerThread
,
NBlockOpLen0
,
NBlockOpLen0
,
NBlockOpLen1
,
NBlockOpLen1
,
NBlockOpLen2
,
NBlockOpLen2
,
...
...
driver/device_direct_convolution_2.cuh
View file @
6790b8f3
#pragma once
#pragma once
#include "direct_convolution_2.cuh"
#include "
gridwise_
direct_convolution_2.cuh"
template
<
class
T
,
class
InDesc
,
class
WeiDesc
,
class
OutDesc
>
template
<
class
T
,
class
InDesc
,
class
WeiDesc
,
class
OutDesc
>
void
device_convolution
(
void
device_
direct_
convolution
_2
(
InDesc
,
const
Tensor
<
T
>&
in
,
WeiDesc
,
const
Tensor
<
T
>&
wei
,
OutDesc
,
Tensor
<
T
>&
out
)
InDesc
,
const
Tensor
<
T
>&
in
,
WeiDesc
,
const
Tensor
<
T
>&
wei
,
OutDesc
,
Tensor
<
T
>&
out
)
{
{
std
::
size_t
data_sz
=
sizeof
(
T
);
std
::
size_t
data_sz
=
sizeof
(
T
);
...
@@ -26,12 +26,16 @@ void device_convolution(
...
@@ -26,12 +26,16 @@ void device_convolution(
constexpr
auto
out_desc
=
OutDesc
{};
constexpr
auto
out_desc
=
OutDesc
{};
constexpr
unsigned
OutTileSizeH
=
2
;
constexpr
unsigned
OutTileSizeH
=
2
;
constexpr
unsigned
OutTileSizeW
=
2
;
constexpr
unsigned
OutTileSizeW
=
2
;
constexpr
unsigned
NPerBlock
=
1
;
constexpr
unsigned
NPerBlock
=
2
;
constexpr
unsigned
KPerBlock
=
4
;
constexpr
unsigned
KPerBlock
=
32
;
constexpr
unsigned
CPerBlock
=
2
;
constexpr
unsigned
CPerBlock
=
2
;
constexpr
unsigned
YPerBlock
=
8
;
constexpr
unsigned
YPerBlock
=
1
;
constexpr
unsigned
XPerBlock
=
16
;
constexpr
unsigned
XPerBlock
=
16
;
constexpr
unsigned
NPerThread
=
2
;
constexpr
unsigned
KPerThread
=
4
;
constexpr
unsigned
CPerThread
=
2
;
constexpr
unsigned
NBlockOpLen0
=
1
;
constexpr
unsigned
NBlockOpLen0
=
1
;
constexpr
unsigned
NBlockOpLen1
=
1
;
constexpr
unsigned
NBlockOpLen1
=
1
;
constexpr
unsigned
NBlockOpLen2
=
4
;
constexpr
unsigned
NBlockOpLen2
=
4
;
...
@@ -66,6 +70,9 @@ void device_convolution(
...
@@ -66,6 +70,9 @@ void device_convolution(
CPerBlock
,
CPerBlock
,
YPerBlock
,
YPerBlock
,
XPerBlock
,
XPerBlock
,
NPerThread
,
KPerThread
,
CPerThread
,
NBlockOpLen0
,
NBlockOpLen0
,
NBlockOpLen1
,
NBlockOpLen1
,
NBlockOpLen2
,
NBlockOpLen2
,
...
...
src/include/blockwise_convolution.cuh
→
src/include/blockwise_
direct_
convolution.cuh
View file @
6790b8f3
#pragma once
#pragma once
#include "constant_tensor_descriptor.cuh"
#include "constant_tensor_descriptor.cuh"
#include "threadwise_tensor_op.cuh"
#include "threadwise_tensor_op.cuh"
#include "threadwise_convolution.cuh"
#include "threadwise_
direct_
convolution.cuh"
template
<
class
TFloat
,
template
<
class
TFloat
,
class
InBlockDesc
,
class
InBlockDesc
,
...
...
src/include/direct_convolution_
2
.cuh
→
src/include/
gridwise_
direct_convolution_
1
.cuh
View file @
6790b8f3
#pragma once
#pragma once
#include "constant_tensor_descriptor.cuh"
#include "constant_tensor_descriptor.cuh"
#include "blockwise_tensor_op.cuh"
#include "blockwise_tensor_op.cuh"
#include "blockwise_convolution.cuh"
#include "blockwise_
direct_
convolution.cuh"
template
<
class
TFloat
,
template
<
class
TFloat
,
class
InGlobalDesc
,
class
InGlobalDesc
,
...
...
src/include/direct_convolution_
3
.cuh
→
src/include/
gridwise_
direct_convolution_
2
.cuh
View file @
6790b8f3
#pragma once
#pragma once
#include "constant_tensor_descriptor.cuh"
#include "constant_tensor_descriptor.cuh"
#include "blockwise_tensor_op.cuh"
#include "blockwise_tensor_op.cuh"
#include "blockwise_convolution.cuh"
#include "blockwise_
direct_
convolution.cuh"
#include "threadwise_tensor_op.cuh"
#include "threadwise_tensor_op.cuh"
#include "threadwise_convolution.cuh"
#include "threadwise_
direct_
convolution.cuh"
template
<
class
TFloat
,
template
<
class
TFloat
,
class
InGlobalDesc
,
class
InGlobalDesc
,
...
...
src/include/threadwise_convolution.cuh
→
src/include/threadwise_
direct_
convolution.cuh
View file @
6790b8f3
File moved
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