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
dfbb4122
Commit
dfbb4122
authored
Jun 06, 2023
by
Jing Zhang
Committed by
root
Jun 06, 2023
Browse files
changed layout for examples
parent
d9a9465f
Changes
8
Expand all
Hide whitespace changes
Inline
Side-by-side
Showing
8 changed files
with
198 additions
and
198 deletions
+198
-198
client_example/16_convnd_fwd/conv3d_fwd_fp16.cpp
client_example/16_convnd_fwd/conv3d_fwd_fp16.cpp
+2
-2
client_example/16_convnd_fwd/conv3d_fwd_fp32.cpp
client_example/16_convnd_fwd/conv3d_fwd_fp32.cpp
+2
-2
library/include/ck/library/tensor_operation_instance/gpu/grouped_convolution_forward.hpp
...or_operation_instance/gpu/grouped_convolution_forward.hpp
+14
-14
library/src/tensor_operation_instance/gpu/grouped_conv3d_fwd/CMakeLists.txt
..._operation_instance/gpu/grouped_conv3d_fwd/CMakeLists.txt
+4
-4
library/src/tensor_operation_instance/gpu/grouped_conv3d_fwd/device_grouped_conv3d_fwd_xdl_ndhwgc_gkzyxc_ndhwgk_bf16_instance.cpp
...ped_conv3d_fwd_xdl_ndhwgc_gkzyxc_ndhwgk_bf16_instance.cpp
+44
-44
library/src/tensor_operation_instance/gpu/grouped_conv3d_fwd/device_grouped_conv3d_fwd_xdl_ndhwgc_gkzyxc_ndhwgk_f16_instance.cpp
...uped_conv3d_fwd_xdl_ndhwgc_gkzyxc_ndhwgk_f16_instance.cpp
+44
-44
library/src/tensor_operation_instance/gpu/grouped_conv3d_fwd/device_grouped_conv3d_fwd_xdl_ndhwgc_gkzyxc_ndhwgk_f32_instance.cpp
...uped_conv3d_fwd_xdl_ndhwgc_gkzyxc_ndhwgk_f32_instance.cpp
+44
-44
library/src/tensor_operation_instance/gpu/grouped_conv3d_fwd/device_grouped_conv3d_fwd_xdl_ndhwgc_gkzyxc_ndhwgk_int8_instance.cpp
...ped_conv3d_fwd_xdl_ndhwgc_gkzyxc_ndhwgk_int8_instance.cpp
+44
-44
No files found.
client_example/16_convnd_fwd/conv3d_fwd_fp16.cpp
View file @
dfbb4122
...
@@ -11,7 +11,7 @@ using WeiDataType = ck::half_t;
...
@@ -11,7 +11,7 @@ using WeiDataType = ck::half_t;
using
OutDataType
=
ck
::
half_t
;
using
OutDataType
=
ck
::
half_t
;
using
InLayout
=
ck
::
tensor_layout
::
convolution
::
NDHWGC
;
using
InLayout
=
ck
::
tensor_layout
::
convolution
::
NDHWGC
;
using
WeiLayout
=
ck
::
tensor_layout
::
convolution
::
KZYX
G
C
;
using
WeiLayout
=
ck
::
tensor_layout
::
convolution
::
G
KZYXC
;
using
OutLayout
=
ck
::
tensor_layout
::
convolution
::
NDHWGK
;
using
OutLayout
=
ck
::
tensor_layout
::
convolution
::
NDHWGK
;
static
constexpr
ck
::
index_t
NumDimSpatial
=
3
;
static
constexpr
ck
::
index_t
NumDimSpatial
=
3
;
...
@@ -38,7 +38,7 @@ int main()
...
@@ -38,7 +38,7 @@ int main()
InLayout
,
InLayout
,
WeiLayout
,
WeiLayout
,
OutLayout
>
(
OutLayout
>
(
{
N
,
Di
,
Hi
,
Wi
,
G
,
C
},
{
K
,
Z
,
Y
,
X
,
G
,
C
},
{
N
,
Do
,
Ho
,
Wo
,
G
,
K
})
{
N
,
Di
,
Hi
,
Wi
,
G
,
C
},
{
G
,
K
,
Z
,
Y
,
X
,
C
},
{
N
,
Do
,
Ho
,
Wo
,
G
,
K
})
?
EXIT_SUCCESS
?
EXIT_SUCCESS
:
EXIT_FAILURE
;
:
EXIT_FAILURE
;
}
}
client_example/16_convnd_fwd/conv3d_fwd_fp32.cpp
View file @
dfbb4122
...
@@ -11,7 +11,7 @@ using WeiDataType = float;
...
@@ -11,7 +11,7 @@ using WeiDataType = float;
using
OutDataType
=
float
;
using
OutDataType
=
float
;
using
InLayout
=
ck
::
tensor_layout
::
convolution
::
NDHWGC
;
using
InLayout
=
ck
::
tensor_layout
::
convolution
::
NDHWGC
;
using
WeiLayout
=
ck
::
tensor_layout
::
convolution
::
KZYX
G
C
;
using
WeiLayout
=
ck
::
tensor_layout
::
convolution
::
G
KZYXC
;
using
OutLayout
=
ck
::
tensor_layout
::
convolution
::
NDHWGK
;
using
OutLayout
=
ck
::
tensor_layout
::
convolution
::
NDHWGK
;
static
constexpr
ck
::
index_t
NumDimSpatial
=
3
;
static
constexpr
ck
::
index_t
NumDimSpatial
=
3
;
...
@@ -38,7 +38,7 @@ int main()
...
@@ -38,7 +38,7 @@ int main()
InLayout
,
InLayout
,
WeiLayout
,
WeiLayout
,
OutLayout
>
(
OutLayout
>
(
{
N
,
Di
,
Hi
,
Wi
,
G
,
C
},
{
K
,
Z
,
Y
,
X
,
G
,
C
},
{
N
,
Do
,
Ho
,
Wo
,
G
,
K
})
{
N
,
Di
,
Hi
,
Wi
,
G
,
C
},
{
G
,
K
,
Z
,
Y
,
X
,
C
},
{
N
,
Do
,
Ho
,
Wo
,
G
,
K
})
?
EXIT_SUCCESS
?
EXIT_SUCCESS
:
EXIT_FAILURE
;
:
EXIT_FAILURE
;
}
}
library/include/ck/library/tensor_operation_instance/gpu/grouped_convolution_forward.hpp
View file @
dfbb4122
...
@@ -245,11 +245,11 @@ void add_device_grouped_conv3d_fwd_xdl_gndhwc_gkzyxc_gndhwk_int8_instances(
...
@@ -245,11 +245,11 @@ void add_device_grouped_conv3d_fwd_xdl_gndhwc_gkzyxc_gndhwk_int8_instances(
PassThrough
,
PassThrough
,
PassThrough
>>>&
instances
);
PassThrough
>>>&
instances
);
// grouped conv3d forward, NDHWGC/KZYX
G
C/NDHWGK
// grouped conv3d forward, NDHWGC/
G
KZYXC/NDHWGK
void
add_device_grouped_conv3d_fwd_xdl_ndhwgc_kzyx
g
c_ndhwgk_bf16_instances
(
void
add_device_grouped_conv3d_fwd_xdl_ndhwgc_
g
kzyxc_ndhwgk_bf16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvFwdMultipleD
<
3
,
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvFwdMultipleD
<
3
,
NDHWGC
,
NDHWGC
,
KZYX
G
C
,
G
KZYXC
,
Empty_Tuple
,
Empty_Tuple
,
NDHWGK
,
NDHWGK
,
BF16
,
BF16
,
...
@@ -260,10 +260,10 @@ void add_device_grouped_conv3d_fwd_xdl_ndhwgc_kzyxgc_ndhwgk_bf16_instances(
...
@@ -260,10 +260,10 @@ void add_device_grouped_conv3d_fwd_xdl_ndhwgc_kzyxgc_ndhwgk_bf16_instances(
PassThrough
,
PassThrough
,
PassThrough
>>>&
instances
);
PassThrough
>>>&
instances
);
void
add_device_grouped_conv3d_fwd_xdl_ndhwgc_kzyx
g
c_ndhwgk_f16_instances
(
void
add_device_grouped_conv3d_fwd_xdl_ndhwgc_
g
kzyxc_ndhwgk_f16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvFwdMultipleD
<
3
,
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvFwdMultipleD
<
3
,
NDHWGC
,
NDHWGC
,
KZYX
G
C
,
G
KZYXC
,
Empty_Tuple
,
Empty_Tuple
,
NDHWGK
,
NDHWGK
,
F16
,
F16
,
...
@@ -274,10 +274,10 @@ void add_device_grouped_conv3d_fwd_xdl_ndhwgc_kzyxgc_ndhwgk_f16_instances(
...
@@ -274,10 +274,10 @@ void add_device_grouped_conv3d_fwd_xdl_ndhwgc_kzyxgc_ndhwgk_f16_instances(
PassThrough
,
PassThrough
,
PassThrough
>>>&
instances
);
PassThrough
>>>&
instances
);
void
add_device_grouped_conv3d_fwd_xdl_ndhwgc_kzyx
g
c_ndhwgk_f32_instances
(
void
add_device_grouped_conv3d_fwd_xdl_ndhwgc_
g
kzyxc_ndhwgk_f32_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvFwdMultipleD
<
3
,
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvFwdMultipleD
<
3
,
NDHWGC
,
NDHWGC
,
KZYX
G
C
,
G
KZYXC
,
Empty_Tuple
,
Empty_Tuple
,
NDHWGK
,
NDHWGK
,
F32
,
F32
,
...
@@ -288,10 +288,10 @@ void add_device_grouped_conv3d_fwd_xdl_ndhwgc_kzyxgc_ndhwgk_f32_instances(
...
@@ -288,10 +288,10 @@ void add_device_grouped_conv3d_fwd_xdl_ndhwgc_kzyxgc_ndhwgk_f32_instances(
PassThrough
,
PassThrough
,
PassThrough
>>>&
instances
);
PassThrough
>>>&
instances
);
void
add_device_grouped_conv3d_fwd_xdl_ndhwgc_kzyx
g
c_ndhwgk_int8_instances
(
void
add_device_grouped_conv3d_fwd_xdl_ndhwgc_
g
kzyxc_ndhwgk_int8_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvFwdMultipleD
<
3
,
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvFwdMultipleD
<
3
,
NDHWGC
,
NDHWGC
,
KZYX
G
C
,
G
KZYXC
,
Empty_Tuple
,
Empty_Tuple
,
NDHWGK
,
NDHWGK
,
int8_t
,
int8_t
,
...
@@ -433,28 +433,28 @@ struct DeviceOperationInstanceFactory<ck::tensor_operation::device::DeviceGroupe
...
@@ -433,28 +433,28 @@ struct DeviceOperationInstanceFactory<ck::tensor_operation::device::DeviceGroupe
}
}
}
}
else
if
constexpr
(
NumDimSpatial
==
3
&&
is_same_v
<
InLayout
,
NDHWGC
>
&&
else
if
constexpr
(
NumDimSpatial
==
3
&&
is_same_v
<
InLayout
,
NDHWGC
>
&&
is_same_v
<
WeiLayout
,
KZYX
G
C
>
&&
is_same_v
<
OutLayout
,
NDHWGK
>
)
is_same_v
<
WeiLayout
,
G
KZYXC
>
&&
is_same_v
<
OutLayout
,
NDHWGK
>
)
{
{
if
constexpr
(
is_same_v
<
InDataType
,
float
>
&&
is_same_v
<
WeiDataType
,
float
>
&&
if
constexpr
(
is_same_v
<
InDataType
,
float
>
&&
is_same_v
<
WeiDataType
,
float
>
&&
is_same_v
<
OutDataType
,
float
>
)
is_same_v
<
OutDataType
,
float
>
)
{
{
add_device_grouped_conv3d_fwd_xdl_ndhwgc_kzyx
g
c_ndhwgk_f32_instances
(
op_ptrs
);
add_device_grouped_conv3d_fwd_xdl_ndhwgc_
g
kzyxc_ndhwgk_f32_instances
(
op_ptrs
);
}
}
else
if
constexpr
(
is_same_v
<
InDataType
,
half_t
>
&&
is_same_v
<
WeiDataType
,
half_t
>
&&
else
if
constexpr
(
is_same_v
<
InDataType
,
half_t
>
&&
is_same_v
<
WeiDataType
,
half_t
>
&&
is_same_v
<
OutDataType
,
half_t
>
)
is_same_v
<
OutDataType
,
half_t
>
)
{
{
add_device_grouped_conv3d_fwd_xdl_ndhwgc_kzyx
g
c_ndhwgk_f16_instances
(
op_ptrs
);
add_device_grouped_conv3d_fwd_xdl_ndhwgc_
g
kzyxc_ndhwgk_f16_instances
(
op_ptrs
);
}
}
else
if
constexpr
(
is_same_v
<
InDataType
,
ck
::
bhalf_t
>
&&
else
if
constexpr
(
is_same_v
<
InDataType
,
ck
::
bhalf_t
>
&&
is_same_v
<
WeiDataType
,
ck
::
bhalf_t
>
&&
is_same_v
<
WeiDataType
,
ck
::
bhalf_t
>
&&
is_same_v
<
OutDataType
,
ck
::
bhalf_t
>
)
is_same_v
<
OutDataType
,
ck
::
bhalf_t
>
)
{
{
add_device_grouped_conv3d_fwd_xdl_ndhwgc_kzyx
g
c_ndhwgk_bf16_instances
(
op_ptrs
);
add_device_grouped_conv3d_fwd_xdl_ndhwgc_
g
kzyxc_ndhwgk_bf16_instances
(
op_ptrs
);
}
}
else
if
constexpr
(
is_same_v
<
InDataType
,
int8_t
>
&&
is_same_v
<
WeiDataType
,
int8_t
>
&&
else
if
constexpr
(
is_same_v
<
InDataType
,
int8_t
>
&&
is_same_v
<
WeiDataType
,
int8_t
>
&&
is_same_v
<
OutDataType
,
int8_t
>
)
is_same_v
<
OutDataType
,
int8_t
>
)
{
{
add_device_grouped_conv3d_fwd_xdl_ndhwgc_kzyx
g
c_ndhwgk_int8_instances
(
op_ptrs
);
add_device_grouped_conv3d_fwd_xdl_ndhwgc_
g
kzyxc_ndhwgk_int8_instances
(
op_ptrs
);
}
}
}
}
...
...
library/src/tensor_operation_instance/gpu/grouped_conv3d_fwd/CMakeLists.txt
View file @
dfbb4122
...
@@ -4,8 +4,8 @@ add_instance_library(device_grouped_conv3d_fwd_instance
...
@@ -4,8 +4,8 @@ add_instance_library(device_grouped_conv3d_fwd_instance
device_grouped_conv3d_fwd_xdl_gndhwc_gkzyxc_gndhwk_f32_instance.cpp
device_grouped_conv3d_fwd_xdl_gndhwc_gkzyxc_gndhwk_f32_instance.cpp
device_grouped_conv3d_fwd_xdl_gndhwc_gkzyxc_gndhwk_int8_instance.cpp
device_grouped_conv3d_fwd_xdl_gndhwc_gkzyxc_gndhwk_int8_instance.cpp
device_grouped_conv3d_fwd_xdl_ndhwgc_kzyx
g
c_ndhwgk_bf16_instance.cpp
device_grouped_conv3d_fwd_xdl_ndhwgc_
g
kzyxc_ndhwgk_bf16_instance.cpp
device_grouped_conv3d_fwd_xdl_ndhwgc_kzyx
g
c_ndhwgk_f16_instance.cpp
device_grouped_conv3d_fwd_xdl_ndhwgc_
g
kzyxc_ndhwgk_f16_instance.cpp
device_grouped_conv3d_fwd_xdl_ndhwgc_kzyx
g
c_ndhwgk_f32_instance.cpp
device_grouped_conv3d_fwd_xdl_ndhwgc_
g
kzyxc_ndhwgk_f32_instance.cpp
device_grouped_conv3d_fwd_xdl_ndhwgc_kzyx
g
c_ndhwgk_int8_instance.cpp
device_grouped_conv3d_fwd_xdl_ndhwgc_
g
kzyxc_ndhwgk_int8_instance.cpp
)
)
library/src/tensor_operation_instance/gpu/grouped_conv3d_fwd/device_grouped_conv3d_fwd_xdl_ndhwgc_gkzyxc_ndhwgk_bf16_instance.cpp
View file @
dfbb4122
This diff is collapsed.
Click to expand it.
library/src/tensor_operation_instance/gpu/grouped_conv3d_fwd/device_grouped_conv3d_fwd_xdl_ndhwgc_gkzyxc_ndhwgk_f16_instance.cpp
View file @
dfbb4122
This diff is collapsed.
Click to expand it.
library/src/tensor_operation_instance/gpu/grouped_conv3d_fwd/device_grouped_conv3d_fwd_xdl_ndhwgc_gkzyxc_ndhwgk_f32_instance.cpp
View file @
dfbb4122
This diff is collapsed.
Click to expand it.
library/src/tensor_operation_instance/gpu/grouped_conv3d_fwd/device_grouped_conv3d_fwd_xdl_ndhwgc_gkzyxc_ndhwgk_int8_instance.cpp
View file @
dfbb4122
This diff is collapsed.
Click to expand it.
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