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
16ec937e
Commit
16ec937e
authored
Jul 18, 2023
by
Bartlomiej Kocot
Browse files
Fix comments
parent
5076b38e
Changes
15
Expand all
Hide whitespace changes
Inline
Side-by-side
Showing
15 changed files
with
76 additions
and
75 deletions
+76
-75
library/include/ck/library/tensor_operation_instance/gpu/grouped_conv_bwd_data/device_grouped_conv_bwd_data_xdl_instance.hpp
...nv_bwd_data/device_grouped_conv_bwd_data_xdl_instance.hpp
+39
-39
library/include/ck/library/tensor_operation_instance/gpu/grouped_convolution_backward_data.hpp
...ration_instance/gpu/grouped_convolution_backward_data.hpp
+24
-24
library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/device_grouped_conv2d_bwd_data_xdl_gnhwc_gkyxc_gnhwk_bf16_instance.cpp
...d_conv2d_bwd_data_xdl_gnhwc_gkyxc_gnhwk_bf16_instance.cpp
+1
-1
library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/device_grouped_conv2d_bwd_data_xdl_gnhwc_gkyxc_gnhwk_f16_instance.cpp
...ed_conv2d_bwd_data_xdl_gnhwc_gkyxc_gnhwk_f16_instance.cpp
+1
-1
library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/device_grouped_conv2d_bwd_data_xdl_gnhwc_gkyxc_gnhwk_f32_instance.cpp
...ed_conv2d_bwd_data_xdl_gnhwc_gkyxc_gnhwk_f32_instance.cpp
+1
-1
library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/device_grouped_conv2d_bwd_data_xdl_nhwgc_gkyxc_nhwgk_bf16_instance.cpp
...d_conv2d_bwd_data_xdl_nhwgc_gkyxc_nhwgk_bf16_instance.cpp
+1
-1
library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/device_grouped_conv2d_bwd_data_xdl_nhwgc_gkyxc_nhwgk_f16_instance.cpp
...ed_conv2d_bwd_data_xdl_nhwgc_gkyxc_nhwgk_f16_instance.cpp
+1
-1
library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/device_grouped_conv2d_bwd_data_xdl_nhwgc_gkyxc_nhwgk_f32_instance.cpp
...ed_conv2d_bwd_data_xdl_nhwgc_gkyxc_nhwgk_f32_instance.cpp
+1
-1
library/src/tensor_operation_instance/gpu/grouped_conv3d_bwd_data/device_grouped_conv3d_bwd_data_xdl_gndhwc_gkzyxc_gndhwk_bf16_instance.cpp
...onv3d_bwd_data_xdl_gndhwc_gkzyxc_gndhwk_bf16_instance.cpp
+1
-1
library/src/tensor_operation_instance/gpu/grouped_conv3d_bwd_data/device_grouped_conv3d_bwd_data_xdl_gndhwc_gkzyxc_gndhwk_f16_instance.cpp
...conv3d_bwd_data_xdl_gndhwc_gkzyxc_gndhwk_f16_instance.cpp
+1
-1
library/src/tensor_operation_instance/gpu/grouped_conv3d_bwd_data/device_grouped_conv3d_bwd_data_xdl_gndhwc_gkzyxc_gndhwk_f32_instance.cpp
...conv3d_bwd_data_xdl_gndhwc_gkzyxc_gndhwk_f32_instance.cpp
+1
-1
library/src/tensor_operation_instance/gpu/grouped_conv3d_bwd_data/device_grouped_conv3d_bwd_data_xdl_ndhwgc_gkzyxc_ndhwgk_bf16_instance.cpp
...onv3d_bwd_data_xdl_ndhwgc_gkzyxc_ndhwgk_bf16_instance.cpp
+1
-1
library/src/tensor_operation_instance/gpu/grouped_conv3d_bwd_data/device_grouped_conv3d_bwd_data_xdl_ndhwgc_gkzyxc_ndhwgk_f16_instance.cpp
...conv3d_bwd_data_xdl_ndhwgc_gkzyxc_ndhwgk_f16_instance.cpp
+1
-1
library/src/tensor_operation_instance/gpu/grouped_conv3d_bwd_data/device_grouped_conv3d_bwd_data_xdl_ndhwgc_gkzyxc_ndhwgk_f32_instance.cpp
...conv3d_bwd_data_xdl_ndhwgc_gkzyxc_ndhwgk_f32_instance.cpp
+1
-1
profiler/src/CMakeLists.txt
profiler/src/CMakeLists.txt
+1
-0
No files found.
library/include/ck/library/tensor_operation_instance/gpu/grouped_conv_bwd_data/device_grouped_conv_bwd_data_xdl_instance.hpp
View file @
16ec937e
This diff is collapsed.
Click to expand it.
library/include/ck/library/tensor_operation_instance/gpu/grouped_convolution_backward_data.hpp
View file @
16ec937e
...
...
@@ -16,7 +16,7 @@ namespace device {
namespace
instance
{
// conv2d backward data
void
add_device_grouped_conv2d_bwd_data_xdl_gnhw
c
_gkyxc_gnhw
k
_f16_instances
(
void
add_device_grouped_conv2d_bwd_data_xdl_gnhw
k
_gkyxc_gnhw
c
_f16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
2
,
GNHWK
,
GKYXC
,
...
...
@@ -30,7 +30,7 @@ void add_device_grouped_conv2d_bwd_data_xdl_gnhwc_gkyxc_gnhwk_f16_instances(
PassThrough
,
PassThrough
>>>&
instances
);
void
add_device_grouped_conv2d_bwd_data_xdl_gnhw
c
_gkyxc_gnhw
k
_f32_instances
(
void
add_device_grouped_conv2d_bwd_data_xdl_gnhw
k
_gkyxc_gnhw
c
_f32_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
2
,
GNHWK
,
GKYXC
,
...
...
@@ -44,7 +44,7 @@ void add_device_grouped_conv2d_bwd_data_xdl_gnhwc_gkyxc_gnhwk_f32_instances(
PassThrough
,
PassThrough
>>>&
instances
);
void
add_device_grouped_conv2d_bwd_data_xdl_gnhw
c
_gkyxc_gnhw
k
_bf16_instances
(
void
add_device_grouped_conv2d_bwd_data_xdl_gnhw
k
_gkyxc_gnhw
c
_bf16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
2
,
GNHWK
,
GKYXC
,
...
...
@@ -58,7 +58,7 @@ void add_device_grouped_conv2d_bwd_data_xdl_gnhwc_gkyxc_gnhwk_bf16_instances(
PassThrough
,
PassThrough
>>>&
instances
);
void
add_device_grouped_conv2d_bwd_data_xdl_nhwg
c
_gkyxc_nhwg
k
_f16_instances
(
void
add_device_grouped_conv2d_bwd_data_xdl_nhwg
k
_gkyxc_nhwg
c
_f16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
2
,
NHWGK
,
GKYXC
,
...
...
@@ -72,7 +72,7 @@ void add_device_grouped_conv2d_bwd_data_xdl_nhwgc_gkyxc_nhwgk_f16_instances(
PassThrough
,
PassThrough
>>>&
instances
);
void
add_device_grouped_conv2d_bwd_data_xdl_nhwg
c
_gkyxc_nhwg
k
_f32_instances
(
void
add_device_grouped_conv2d_bwd_data_xdl_nhwg
k
_gkyxc_nhwg
c
_f32_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
2
,
NHWGK
,
GKYXC
,
...
...
@@ -86,7 +86,7 @@ void add_device_grouped_conv2d_bwd_data_xdl_nhwgc_gkyxc_nhwgk_f32_instances(
PassThrough
,
PassThrough
>>>&
instances
);
void
add_device_grouped_conv2d_bwd_data_xdl_nhwg
c
_gkyxc_nhwg
k
_bf16_instances
(
void
add_device_grouped_conv2d_bwd_data_xdl_nhwg
k
_gkyxc_nhwg
c
_bf16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
2
,
NHWGK
,
GKYXC
,
...
...
@@ -101,7 +101,7 @@ void add_device_grouped_conv2d_bwd_data_xdl_nhwgc_gkyxc_nhwgk_bf16_instances(
PassThrough
>>>&
instances
);
// conv3d backward data
void
add_device_grouped_conv3d_bwd_data_xdl_gndhw
c
_gkzyxc_gndhw
k
_f16_instances
(
void
add_device_grouped_conv3d_bwd_data_xdl_gndhw
k
_gkzyxc_gndhw
c
_f16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
3
,
GNDHWK
,
GKZYXC
,
...
...
@@ -115,7 +115,7 @@ void add_device_grouped_conv3d_bwd_data_xdl_gndhwc_gkzyxc_gndhwk_f16_instances(
PassThrough
,
PassThrough
>>>&
instances
);
void
add_device_grouped_conv3d_bwd_data_xdl_gndhw
c
_gkzyxc_gndhw
k
_f32_instances
(
void
add_device_grouped_conv3d_bwd_data_xdl_gndhw
k
_gkzyxc_gndhw
c
_f32_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
3
,
GNDHWK
,
GKZYXC
,
...
...
@@ -129,7 +129,7 @@ void add_device_grouped_conv3d_bwd_data_xdl_gndhwc_gkzyxc_gndhwk_f32_instances(
PassThrough
,
PassThrough
>>>&
instances
);
void
add_device_grouped_conv3d_bwd_data_xdl_gndhw
c
_gkzyxc_gndhw
k
_bf16_instances
(
void
add_device_grouped_conv3d_bwd_data_xdl_gndhw
k
_gkzyxc_gndhw
c
_bf16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
3
,
GNDHWK
,
GKZYXC
,
...
...
@@ -143,7 +143,7 @@ void add_device_grouped_conv3d_bwd_data_xdl_gndhwc_gkzyxc_gndhwk_bf16_instances(
PassThrough
,
PassThrough
>>>&
instances
);
void
add_device_grouped_conv3d_bwd_data_xdl_ndhwg
c
_gkzyxc_ndhwg
k
_f16_instances
(
void
add_device_grouped_conv3d_bwd_data_xdl_ndhwg
k
_gkzyxc_ndhwg
c
_f16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
3
,
NDHWGK
,
GKZYXC
,
...
...
@@ -157,7 +157,7 @@ void add_device_grouped_conv3d_bwd_data_xdl_ndhwgc_gkzyxc_ndhwgk_f16_instances(
PassThrough
,
PassThrough
>>>&
instances
);
void
add_device_grouped_conv3d_bwd_data_xdl_ndhwg
c
_gkzyxc_ndhwg
k
_f32_instances
(
void
add_device_grouped_conv3d_bwd_data_xdl_ndhwg
k
_gkzyxc_ndhwg
c
_f32_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
3
,
NDHWGK
,
GKZYXC
,
...
...
@@ -171,7 +171,7 @@ void add_device_grouped_conv3d_bwd_data_xdl_ndhwgc_gkzyxc_ndhwgk_f32_instances(
PassThrough
,
PassThrough
>>>&
instances
);
void
add_device_grouped_conv3d_bwd_data_xdl_ndhwg
c
_gkzyxc_ndhwg
k
_bf16_instances
(
void
add_device_grouped_conv3d_bwd_data_xdl_ndhwg
k
_gkzyxc_ndhwg
c
_bf16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
3
,
NDHWGK
,
GKZYXC
,
...
...
@@ -233,17 +233,17 @@ struct DeviceOperationInstanceFactory<
if
constexpr
(
is_same_v
<
InDataType
,
F16
>
&&
is_same_v
<
WeiDataType
,
F16
>
&&
is_same_v
<
OutDataType
,
F16
>
)
{
add_device_grouped_conv2d_bwd_data_xdl_gnhw
c
_gkyxc_gnhw
k
_f16_instances
(
op_ptrs
);
add_device_grouped_conv2d_bwd_data_xdl_gnhw
k
_gkyxc_gnhw
c
_f16_instances
(
op_ptrs
);
}
else
if
constexpr
(
is_same_v
<
InDataType
,
F32
>
&&
is_same_v
<
WeiDataType
,
F32
>
&&
is_same_v
<
OutDataType
,
F32
>
)
{
add_device_grouped_conv2d_bwd_data_xdl_gnhw
c
_gkyxc_gnhw
k
_f32_instances
(
op_ptrs
);
add_device_grouped_conv2d_bwd_data_xdl_gnhw
k
_gkyxc_gnhw
c
_f32_instances
(
op_ptrs
);
}
else
if
constexpr
(
is_same_v
<
InDataType
,
BF16
>
&&
is_same_v
<
WeiDataType
,
BF16
>
&&
is_same_v
<
OutDataType
,
BF16
>
)
{
add_device_grouped_conv2d_bwd_data_xdl_gnhw
c
_gkyxc_gnhw
k
_bf16_instances
(
add_device_grouped_conv2d_bwd_data_xdl_gnhw
k
_gkyxc_gnhw
c
_bf16_instances
(
op_ptrs
);
}
}
...
...
@@ -253,17 +253,17 @@ struct DeviceOperationInstanceFactory<
if
constexpr
(
is_same_v
<
InDataType
,
F16
>
&&
is_same_v
<
WeiDataType
,
F16
>
&&
is_same_v
<
OutDataType
,
F16
>
)
{
add_device_grouped_conv2d_bwd_data_xdl_nhwg
c
_gkyxc_nhwg
k
_f16_instances
(
op_ptrs
);
add_device_grouped_conv2d_bwd_data_xdl_nhwg
k
_gkyxc_nhwg
c
_f16_instances
(
op_ptrs
);
}
else
if
constexpr
(
is_same_v
<
InDataType
,
F32
>
&&
is_same_v
<
WeiDataType
,
F32
>
&&
is_same_v
<
OutDataType
,
F32
>
)
{
add_device_grouped_conv2d_bwd_data_xdl_nhwg
c
_gkyxc_nhwg
k
_f32_instances
(
op_ptrs
);
add_device_grouped_conv2d_bwd_data_xdl_nhwg
k
_gkyxc_nhwg
c
_f32_instances
(
op_ptrs
);
}
else
if
constexpr
(
is_same_v
<
InDataType
,
BF16
>
&&
is_same_v
<
WeiDataType
,
BF16
>
&&
is_same_v
<
OutDataType
,
BF16
>
)
{
add_device_grouped_conv2d_bwd_data_xdl_nhwg
c
_gkyxc_nhwg
k
_bf16_instances
(
add_device_grouped_conv2d_bwd_data_xdl_nhwg
k
_gkyxc_nhwg
c
_bf16_instances
(
op_ptrs
);
}
}
...
...
@@ -277,19 +277,19 @@ struct DeviceOperationInstanceFactory<
if
constexpr
(
is_same_v
<
InDataType
,
F16
>
&&
is_same_v
<
WeiDataType
,
F16
>
&&
is_same_v
<
OutDataType
,
F16
>
)
{
add_device_grouped_conv3d_bwd_data_xdl_gndhw
c
_gkzyxc_gndhw
k
_f16_instances
(
add_device_grouped_conv3d_bwd_data_xdl_gndhw
k
_gkzyxc_gndhw
c
_f16_instances
(
op_ptrs
);
}
else
if
constexpr
(
is_same_v
<
InDataType
,
F32
>
&&
is_same_v
<
WeiDataType
,
F32
>
&&
is_same_v
<
OutDataType
,
F32
>
)
{
add_device_grouped_conv3d_bwd_data_xdl_gndhw
c
_gkzyxc_gndhw
k
_f32_instances
(
add_device_grouped_conv3d_bwd_data_xdl_gndhw
k
_gkzyxc_gndhw
c
_f32_instances
(
op_ptrs
);
}
else
if
constexpr
(
is_same_v
<
InDataType
,
BF16
>
&&
is_same_v
<
WeiDataType
,
BF16
>
&&
is_same_v
<
OutDataType
,
BF16
>
)
{
add_device_grouped_conv3d_bwd_data_xdl_gndhw
c
_gkzyxc_gndhw
k
_bf16_instances
(
add_device_grouped_conv3d_bwd_data_xdl_gndhw
k
_gkzyxc_gndhw
c
_bf16_instances
(
op_ptrs
);
}
}
...
...
@@ -299,19 +299,19 @@ struct DeviceOperationInstanceFactory<
if
constexpr
(
is_same_v
<
InDataType
,
F16
>
&&
is_same_v
<
WeiDataType
,
F16
>
&&
is_same_v
<
OutDataType
,
F16
>
)
{
add_device_grouped_conv3d_bwd_data_xdl_ndhwg
c
_gkzyxc_ndhwg
k
_f16_instances
(
add_device_grouped_conv3d_bwd_data_xdl_ndhwg
k
_gkzyxc_ndhwg
c
_f16_instances
(
op_ptrs
);
}
else
if
constexpr
(
is_same_v
<
InDataType
,
F32
>
&&
is_same_v
<
WeiDataType
,
F32
>
&&
is_same_v
<
OutDataType
,
F32
>
)
{
add_device_grouped_conv3d_bwd_data_xdl_ndhwg
c
_gkzyxc_ndhwg
k
_f32_instances
(
add_device_grouped_conv3d_bwd_data_xdl_ndhwg
k
_gkzyxc_ndhwg
c
_f32_instances
(
op_ptrs
);
}
else
if
constexpr
(
is_same_v
<
InDataType
,
BF16
>
&&
is_same_v
<
WeiDataType
,
BF16
>
&&
is_same_v
<
OutDataType
,
BF16
>
)
{
add_device_grouped_conv3d_bwd_data_xdl_ndhwg
c
_gkzyxc_ndhwg
k
_bf16_instances
(
add_device_grouped_conv3d_bwd_data_xdl_ndhwg
k
_gkzyxc_ndhwg
c
_bf16_instances
(
op_ptrs
);
}
}
...
...
library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/device_grouped_conv2d_bwd_data_xdl_gnhwc_gkyxc_gnhwk_bf16_instance.cpp
View file @
16ec937e
...
...
@@ -9,7 +9,7 @@ namespace tensor_operation {
namespace
device
{
namespace
instance
{
// Compilation parameters for out[g, n, hi, wi, c] * wei[g, k, y, x, c] = in[g, n, ho, wo, k]
void
add_device_grouped_conv2d_bwd_data_xdl_gnhw
c
_gkyxc_gnhw
k
_bf16_instances
(
void
add_device_grouped_conv2d_bwd_data_xdl_gnhw
k
_gkyxc_gnhw
c
_bf16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
2
,
GNHWK
,
GKYXC
,
...
...
library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/device_grouped_conv2d_bwd_data_xdl_gnhwc_gkyxc_gnhwk_f16_instance.cpp
View file @
16ec937e
...
...
@@ -9,7 +9,7 @@ namespace tensor_operation {
namespace
device
{
namespace
instance
{
// Compilation parameters for out[g, n, hi, wi, c] * wei[g, k, y, x, c] = in[g, n, ho, wo, k]
void
add_device_grouped_conv2d_bwd_data_xdl_gnhw
c
_gkyxc_gnhw
k
_f16_instances
(
void
add_device_grouped_conv2d_bwd_data_xdl_gnhw
k
_gkyxc_gnhw
c
_f16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
2
,
GNHWK
,
GKYXC
,
...
...
library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/device_grouped_conv2d_bwd_data_xdl_gnhwc_gkyxc_gnhwk_f32_instance.cpp
View file @
16ec937e
...
...
@@ -9,7 +9,7 @@ namespace tensor_operation {
namespace
device
{
namespace
instance
{
// Compilation parameters for out[g, n, hi, wi, c] * wei[g, k, y, x, c] = in[g, n, ho, wo, k]
void
add_device_grouped_conv2d_bwd_data_xdl_gnhw
c
_gkyxc_gnhw
k
_f32_instances
(
void
add_device_grouped_conv2d_bwd_data_xdl_gnhw
k
_gkyxc_gnhw
c
_f32_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
2
,
GNHWK
,
GKYXC
,
...
...
library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/device_grouped_conv2d_bwd_data_xdl_nhwgc_gkyxc_nhwgk_bf16_instance.cpp
View file @
16ec937e
...
...
@@ -9,7 +9,7 @@ namespace tensor_operation {
namespace
device
{
namespace
instance
{
// Compilation parameters for out[n, hi, wi, g, c] * wei[g, k, y, x, c] = in[n, ho, wo, g, k]
void
add_device_grouped_conv2d_bwd_data_xdl_nhwg
c
_gkyxc_nhwg
k
_bf16_instances
(
void
add_device_grouped_conv2d_bwd_data_xdl_nhwg
k
_gkyxc_nhwg
c
_bf16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
2
,
NHWGK
,
GKYXC
,
...
...
library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/device_grouped_conv2d_bwd_data_xdl_nhwgc_gkyxc_nhwgk_f16_instance.cpp
View file @
16ec937e
...
...
@@ -9,7 +9,7 @@ namespace tensor_operation {
namespace
device
{
namespace
instance
{
// Compilation parameters for out[n, hi, wi, g, c] * wei[g, k, y, x, c] = in[n, ho, wo, g, k]
void
add_device_grouped_conv2d_bwd_data_xdl_nhwg
c
_gkyxc_nhwg
k
_f16_instances
(
void
add_device_grouped_conv2d_bwd_data_xdl_nhwg
k
_gkyxc_nhwg
c
_f16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
2
,
NHWGK
,
GKYXC
,
...
...
library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/device_grouped_conv2d_bwd_data_xdl_nhwgc_gkyxc_nhwgk_f32_instance.cpp
View file @
16ec937e
...
...
@@ -9,7 +9,7 @@ namespace tensor_operation {
namespace
device
{
namespace
instance
{
// Compilation parameters for out[n, hi, wi, g, c] * wei[g, k, y, x, c] = in[n, ho, wo, g, k]
void
add_device_grouped_conv2d_bwd_data_xdl_nhwg
c
_gkyxc_nhwg
k
_f32_instances
(
void
add_device_grouped_conv2d_bwd_data_xdl_nhwg
k
_gkyxc_nhwg
c
_f32_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
2
,
NHWGK
,
GKYXC
,
...
...
library/src/tensor_operation_instance/gpu/grouped_conv3d_bwd_data/device_grouped_conv3d_bwd_data_xdl_gndhwc_gkzyxc_gndhwk_bf16_instance.cpp
View file @
16ec937e
...
...
@@ -10,7 +10,7 @@ namespace device {
namespace
instance
{
// Compilation parameters for out[g, n, di, hi, wi, c] * wei[g, k, z, y, x, c] = in[g, n, do, ho,
// wo, k]
void
add_device_grouped_conv3d_bwd_data_xdl_gndhw
c
_gkzyxc_gndhw
k
_bf16_instances
(
void
add_device_grouped_conv3d_bwd_data_xdl_gndhw
k
_gkzyxc_gndhw
c
_bf16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
3
,
GNDHWK
,
GKZYXC
,
...
...
library/src/tensor_operation_instance/gpu/grouped_conv3d_bwd_data/device_grouped_conv3d_bwd_data_xdl_gndhwc_gkzyxc_gndhwk_f16_instance.cpp
View file @
16ec937e
...
...
@@ -10,7 +10,7 @@ namespace device {
namespace
instance
{
// Compilation parameters for out[g, n, di, hi, wi, c] * wei[g, k, z, y, x, c] = in[g, n, do, ho,
// wo, k]
void
add_device_grouped_conv3d_bwd_data_xdl_gndhw
c
_gkzyxc_gndhw
k
_f16_instances
(
void
add_device_grouped_conv3d_bwd_data_xdl_gndhw
k
_gkzyxc_gndhw
c
_f16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
3
,
GNDHWK
,
GKZYXC
,
...
...
library/src/tensor_operation_instance/gpu/grouped_conv3d_bwd_data/device_grouped_conv3d_bwd_data_xdl_gndhwc_gkzyxc_gndhwk_f32_instance.cpp
View file @
16ec937e
...
...
@@ -10,7 +10,7 @@ namespace device {
namespace
instance
{
// Compilation parameters for out[g, n, di, hi, wi, c] * wei[g, k, z, y, x, c] = in[g, n, do, ho,
// wo, k]
void
add_device_grouped_conv3d_bwd_data_xdl_gndhw
c
_gkzyxc_gndhw
k
_f32_instances
(
void
add_device_grouped_conv3d_bwd_data_xdl_gndhw
k
_gkzyxc_gndhw
c
_f32_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
3
,
GNDHWK
,
GKZYXC
,
...
...
library/src/tensor_operation_instance/gpu/grouped_conv3d_bwd_data/device_grouped_conv3d_bwd_data_xdl_ndhwgc_gkzyxc_ndhwgk_bf16_instance.cpp
View file @
16ec937e
...
...
@@ -10,7 +10,7 @@ namespace device {
namespace
instance
{
// Compilation parameters for out[n, di, hi, wi, g, c] * wei[g, k, z, y, x, c] = in[n, do, ho, wo,
// g, k]
void
add_device_grouped_conv3d_bwd_data_xdl_ndhwg
c
_gkzyxc_ndhwg
k
_bf16_instances
(
void
add_device_grouped_conv3d_bwd_data_xdl_ndhwg
k
_gkzyxc_ndhwg
c
_bf16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
3
,
NDHWGK
,
GKZYXC
,
...
...
library/src/tensor_operation_instance/gpu/grouped_conv3d_bwd_data/device_grouped_conv3d_bwd_data_xdl_ndhwgc_gkzyxc_ndhwgk_f16_instance.cpp
View file @
16ec937e
...
...
@@ -10,7 +10,7 @@ namespace device {
namespace
instance
{
// Compilation parameters for out[n, di, hi, wi, g, c] * wei[g, k, z, y, x, c] = in[n, do, ho, wo,
// g, k]
void
add_device_grouped_conv3d_bwd_data_xdl_ndhwg
c
_gkzyxc_ndhwg
k
_f16_instances
(
void
add_device_grouped_conv3d_bwd_data_xdl_ndhwg
k
_gkzyxc_ndhwg
c
_f16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
3
,
NDHWGK
,
GKZYXC
,
...
...
library/src/tensor_operation_instance/gpu/grouped_conv3d_bwd_data/device_grouped_conv3d_bwd_data_xdl_ndhwgc_gkzyxc_ndhwgk_f32_instance.cpp
View file @
16ec937e
...
...
@@ -10,7 +10,7 @@ namespace device {
namespace
instance
{
// Compilation parameters for out[n, di, hi, wi, g, c] * wei[g, k, z, y, x, c] = in[n, do, ho, wo,
// g, k]
void
add_device_grouped_conv3d_bwd_data_xdl_ndhwg
c
_gkzyxc_ndhwg
k
_f32_instances
(
void
add_device_grouped_conv3d_bwd_data_xdl_ndhwg
k
_gkzyxc_ndhwg
c
_f32_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGroupedConvBwdDataMultipleD
<
3
,
NDHWGK
,
GKZYXC
,
...
...
profiler/src/CMakeLists.txt
View file @
16ec937e
...
...
@@ -81,4 +81,5 @@ target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_contraction_scale_in
target_link_libraries
(
${
PROFILER_EXECUTABLE
}
PRIVATE device_pool_fwd_instance
)
target_link_libraries
(
${
PROFILER_EXECUTABLE
}
PRIVATE device_batched_gemm_multi_d_instance
)
target_link_libraries
(
${
PROFILER_EXECUTABLE
}
PRIVATE device_grouped_conv2d_bwd_data_instance
)
target_link_libraries
(
${
PROFILER_EXECUTABLE
}
PRIVATE device_grouped_conv3d_bwd_data_instance
)
rocm_install
(
TARGETS
${
PROFILER_EXECUTABLE
}
COMPONENT profiler
)
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