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
36f54418
Commit
36f54418
authored
Oct 30, 2023
by
Bartlomiej Kocot
Browse files
Fixes
parent
d9eadda9
Changes
11
Hide whitespace changes
Inline
Side-by-side
Showing
11 changed files
with
28 additions
and
28 deletions
+28
-28
client_example/22_im2col_col2im/column_to_image.cpp
client_example/22_im2col_col2im/column_to_image.cpp
+2
-2
client_example/22_im2col_col2im/image_to_column.cpp
client_example/22_im2col_col2im/image_to_column.cpp
+2
-2
example/52_im2col_col2im/image_to_column_f32.cpp
example/52_im2col_col2im/image_to_column_f32.cpp
+3
-3
include/ck/tensor_operation/gpu/device/device_conv_tensor_rearrange.hpp
...sor_operation/gpu/device/device_conv_tensor_rearrange.hpp
+3
-3
library/include/ck/library/tensor_operation_instance/gpu/conv_tensor_rearrange.hpp
...y/tensor_operation_instance/gpu/conv_tensor_rearrange.hpp
+12
-12
library/src/tensor_operation_instance/gpu/image_to_column/device_image_to_column_gndhwc_3d_instance.cpp
...e_to_column/device_image_to_column_gndhwc_3d_instance.cpp
+1
-1
library/src/tensor_operation_instance/gpu/image_to_column/device_image_to_column_gnhwc_2d_instance.cpp
...ge_to_column/device_image_to_column_gnhwc_2d_instance.cpp
+1
-1
library/src/tensor_operation_instance/gpu/image_to_column/device_image_to_column_gnwc_1d_instance.cpp
...age_to_column/device_image_to_column_gnwc_1d_instance.cpp
+1
-1
library/src/tensor_operation_instance/gpu/image_to_column/device_image_to_column_ndhwgc_3d_instance.cpp
...e_to_column/device_image_to_column_ndhwgc_3d_instance.cpp
+1
-1
library/src/tensor_operation_instance/gpu/image_to_column/device_image_to_column_nhwgc_2d_instance.cpp
...ge_to_column/device_image_to_column_nhwgc_2d_instance.cpp
+1
-1
library/src/tensor_operation_instance/gpu/image_to_column/device_image_to_column_nwgc_1d_instance.cpp
...age_to_column/device_image_to_column_nwgc_1d_instance.cpp
+1
-1
No files found.
client_example/22_im2col_col2im/column_to_image.cpp
View file @
36f54418
...
@@ -53,8 +53,8 @@ int main()
...
@@ -53,8 +53,8 @@ int main()
std
::
array
<
ck
::
index_t
,
2
>
out_spatial_lengths
{
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
2
>
out_spatial_lengths
{
Ho
,
Wo
};
// We have NHWGC in memory space
// We have NHWGC in memory space
// However, CK's API only accept length and stride with order of GNCHW
// However, CK's API only accept
s
length
s
and stride
s
with order of GNCHW
.
// Hence, we need to adjust the order of stride
// Hence, we need to adjust the order of stride
s.
std
::
array
<
ck
::
index_t
,
5
>
image_strides
{
C
,
Hi
*
Wi
*
G
*
C
,
1
,
Wi
*
G
*
C
,
G
*
C
};
std
::
array
<
ck
::
index_t
,
5
>
image_strides
{
C
,
Hi
*
Wi
*
G
*
C
,
1
,
Wi
*
G
*
C
,
G
*
C
};
std
::
array
<
ck
::
index_t
,
3
>
gemm_strides
{
Y
*
X
*
C
,
G
*
Y
*
X
*
C
,
1
};
std
::
array
<
ck
::
index_t
,
3
>
gemm_strides
{
Y
*
X
*
C
,
G
*
Y
*
X
*
C
,
1
};
...
...
client_example/22_im2col_col2im/image_to_column.cpp
View file @
36f54418
...
@@ -53,8 +53,8 @@ int main()
...
@@ -53,8 +53,8 @@ int main()
std
::
array
<
ck
::
index_t
,
2
>
out_spatial_lengths
{
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
2
>
out_spatial_lengths
{
Ho
,
Wo
};
// We have NHWGC in memory space
// We have NHWGC in memory space
// However, CK's API only accept length and stride with order of GNCHW
// However, CK's API only accept
s
length
s
and stride
s
with order of GNCHW
.
// Hence, we need to adjust the order of stride
// Hence, we need to adjust the order of stride
s.
std
::
array
<
ck
::
index_t
,
5
>
image_strides
{
C
,
Hi
*
Wi
*
G
*
C
,
1
,
Wi
*
G
*
C
,
G
*
C
};
std
::
array
<
ck
::
index_t
,
5
>
image_strides
{
C
,
Hi
*
Wi
*
G
*
C
,
1
,
Wi
*
G
*
C
,
G
*
C
};
std
::
array
<
ck
::
index_t
,
3
>
gemm_strides
{
Y
*
X
*
C
,
G
*
Y
*
X
*
C
,
1
};
std
::
array
<
ck
::
index_t
,
3
>
gemm_strides
{
Y
*
X
*
C
,
G
*
Y
*
X
*
C
,
1
};
...
...
example/52_im2col_col2im/image_to_column_f32.cpp
View file @
36f54418
...
@@ -39,7 +39,7 @@ bool RunImageToColumn(const ExecutionConfig& config, const ck::utils::conv::Conv
...
@@ -39,7 +39,7 @@ bool RunImageToColumn(const ExecutionConfig& config, const ck::utils::conv::Conv
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
filter_spatial_lengths
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
filter_spatial_lengths
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
output_spatial_lengths
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
output_spatial_lengths
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
+
3
>
image_g_n_c_wis_strides
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
+
3
>
image_g_n_c_wis_strides
{};
std
::
array
<
ck
::
index_t
,
3
>
gemm_m_k_strides
{};
std
::
array
<
ck
::
index_t
,
3
>
gemm_
g_
m_k_strides
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
conv_filter_strides
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
conv_filter_strides
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
conv_filter_dilations
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
conv_filter_dilations
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
input_left_pads
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
input_left_pads
{};
...
@@ -51,7 +51,7 @@ bool RunImageToColumn(const ExecutionConfig& config, const ck::utils::conv::Conv
...
@@ -51,7 +51,7 @@ bool RunImageToColumn(const ExecutionConfig& config, const ck::utils::conv::Conv
copy
(
conv_params
.
filter_spatial_lengths_
,
filter_spatial_lengths
);
copy
(
conv_params
.
filter_spatial_lengths_
,
filter_spatial_lengths
);
copy
(
conv_params
.
output_spatial_lengths_
,
output_spatial_lengths
);
copy
(
conv_params
.
output_spatial_lengths_
,
output_spatial_lengths
);
copy
(
in_desc
.
GetStrides
(),
image_g_n_c_wis_strides
);
copy
(
in_desc
.
GetStrides
(),
image_g_n_c_wis_strides
);
copy
(
out_desc
.
GetStrides
(),
gemm_m_k_strides
);
copy
(
out_desc
.
GetStrides
(),
gemm_
g_
m_k_strides
);
copy
(
conv_params
.
conv_filter_strides_
,
conv_filter_strides
);
copy
(
conv_params
.
conv_filter_strides_
,
conv_filter_strides
);
copy
(
conv_params
.
conv_filter_dilations_
,
conv_filter_dilations
);
copy
(
conv_params
.
conv_filter_dilations_
,
conv_filter_dilations
);
copy
(
conv_params
.
input_left_pads_
,
input_left_pads
);
copy
(
conv_params
.
input_left_pads_
,
input_left_pads
);
...
@@ -93,7 +93,7 @@ bool RunImageToColumn(const ExecutionConfig& config, const ck::utils::conv::Conv
...
@@ -93,7 +93,7 @@ bool RunImageToColumn(const ExecutionConfig& config, const ck::utils::conv::Conv
filter_spatial_lengths
,
filter_spatial_lengths
,
output_spatial_lengths
,
output_spatial_lengths
,
image_g_n_c_wis_strides
,
image_g_n_c_wis_strides
,
gemm_m_k_strides
,
gemm_
g_
m_k_strides
,
conv_filter_strides
,
conv_filter_strides
,
conv_filter_dilations
,
conv_filter_dilations
,
input_left_pads
,
input_left_pads
,
...
...
include/ck/tensor_operation/gpu/device/device_conv_tensor_rearrange.hpp
View file @
36f54418
...
@@ -14,9 +14,9 @@ namespace device {
...
@@ -14,9 +14,9 @@ namespace device {
/**
/**
* \brief Convolution Tensor Rearrange.
* \brief Convolution Tensor Rearrange.
*
*
* This Device operator supports conver
sio
n image to
* This Device operator supports conver
ting a
n image to
* the
gemm problem
(Image to Column) and
* the
GEMM representation
(Image to Column) and
* conver
sion gemm
form to the image (Column to Image).
* conver
ting a GEMM
form to the image (Column to Image).
* Supported layouts:
* Supported layouts:
* [G, N, Di, Hi, Wi, C] <-> [G, N * Do * Ho * Wo, Z * Y * X * C]
* [G, N, Di, Hi, Wi, C] <-> [G, N * Do * Ho * Wo, Z * Y * X * C]
* [N, Di, Hi, Wi, G, C] <-> [N * Do * Ho * Wo, G, Z * Y * X * C]
* [N, Di, Hi, Wi, G, C] <-> [N * Do * Ho * Wo, G, Z * Y * X * C]
...
...
library/include/ck/library/tensor_operation_instance/gpu/conv_tensor_rearrange.hpp
View file @
36f54418
...
@@ -21,7 +21,7 @@ using namespace ck::conv_tensor_rearrange_op;
...
@@ -21,7 +21,7 @@ using namespace ck::conv_tensor_rearrange_op;
// GNWC/GNHWC/GNDHWC
// GNWC/GNHWC/GNDHWC
// Image to Column
// Image to Column
//
nhwc
, 1d
//
GNWC
, 1d
void
add_device_image_to_column_gnwc_1d_bf16_instances
(
void
add_device_image_to_column_gnwc_1d_bf16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
1
,
GNWC
,
BF16
,
BF16
,
ImageToColumn
>>>&
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
1
,
GNWC
,
BF16
,
BF16
,
ImageToColumn
>>>&
instances
);
instances
);
...
@@ -37,7 +37,7 @@ void add_device_image_to_column_gnwc_1d_f32_instances(
...
@@ -37,7 +37,7 @@ void add_device_image_to_column_gnwc_1d_f32_instances(
void
add_device_image_to_column_gnwc_1d_i8_instances
(
void
add_device_image_to_column_gnwc_1d_i8_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
1
,
GNWC
,
int8_t
,
int8_t
,
ImageToColumn
>>>&
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
1
,
GNWC
,
int8_t
,
int8_t
,
ImageToColumn
>>>&
instances
);
instances
);
//
nhwc
, 2d
//
GNHWC
, 2d
void
add_device_image_to_column_gnhwc_2d_bf16_instances
(
void
add_device_image_to_column_gnhwc_2d_bf16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
2
,
GNHWC
,
BF16
,
BF16
,
ImageToColumn
>>>&
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
2
,
GNHWC
,
BF16
,
BF16
,
ImageToColumn
>>>&
instances
);
instances
);
...
@@ -54,7 +54,7 @@ void add_device_image_to_column_gnhwc_2d_i8_instances(
...
@@ -54,7 +54,7 @@ void add_device_image_to_column_gnhwc_2d_i8_instances(
std
::
vector
<
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
2
,
GNHWC
,
int8_t
,
int8_t
,
ImageToColumn
>>>&
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
2
,
GNHWC
,
int8_t
,
int8_t
,
ImageToColumn
>>>&
instances
);
instances
);
//
nhwc
, 3d
//
GNDHWC
, 3d
void
add_device_image_to_column_gndhwc_3d_bf16_instances
(
void
add_device_image_to_column_gndhwc_3d_bf16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
3
,
GNDHWC
,
BF16
,
BF16
,
ImageToColumn
>>>&
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
3
,
GNDHWC
,
BF16
,
BF16
,
ImageToColumn
>>>&
instances
);
instances
);
...
@@ -73,7 +73,7 @@ void add_device_image_to_column_gndhwc_3d_i8_instances(
...
@@ -73,7 +73,7 @@ void add_device_image_to_column_gndhwc_3d_i8_instances(
instances
);
instances
);
// Column to Image
// Column to Image
//
nhwc
, 1d
//
GNWC
, 1d
void
add_device_column_to_image_gnwc_1d_bf16_instances
(
void
add_device_column_to_image_gnwc_1d_bf16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
1
,
GNWC
,
BF16
,
BF16
,
ColumnToImage
>>>&
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
1
,
GNWC
,
BF16
,
BF16
,
ColumnToImage
>>>&
instances
);
instances
);
...
@@ -89,7 +89,7 @@ void add_device_column_to_image_gnwc_1d_f32_instances(
...
@@ -89,7 +89,7 @@ void add_device_column_to_image_gnwc_1d_f32_instances(
void
add_device_column_to_image_gnwc_1d_i8_instances
(
void
add_device_column_to_image_gnwc_1d_i8_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
1
,
GNWC
,
int8_t
,
int8_t
,
ColumnToImage
>>>&
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
1
,
GNWC
,
int8_t
,
int8_t
,
ColumnToImage
>>>&
instances
);
instances
);
//
nhwc
, 2d
//
GNHWC
, 2d
void
add_device_column_to_image_gnhwc_2d_bf16_instances
(
void
add_device_column_to_image_gnhwc_2d_bf16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
2
,
GNHWC
,
BF16
,
BF16
,
ColumnToImage
>>>&
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
2
,
GNHWC
,
BF16
,
BF16
,
ColumnToImage
>>>&
instances
);
instances
);
...
@@ -106,7 +106,7 @@ void add_device_column_to_image_gnhwc_2d_i8_instances(
...
@@ -106,7 +106,7 @@ void add_device_column_to_image_gnhwc_2d_i8_instances(
std
::
vector
<
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
2
,
GNHWC
,
int8_t
,
int8_t
,
ColumnToImage
>>>&
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
2
,
GNHWC
,
int8_t
,
int8_t
,
ColumnToImage
>>>&
instances
);
instances
);
//
nhwc
, 3d
//
GNDHWC
, 3d
void
add_device_column_to_image_gndhwc_3d_bf16_instances
(
void
add_device_column_to_image_gndhwc_3d_bf16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
3
,
GNDHWC
,
BF16
,
BF16
,
ColumnToImage
>>>&
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
3
,
GNDHWC
,
BF16
,
BF16
,
ColumnToImage
>>>&
instances
);
instances
);
...
@@ -125,7 +125,7 @@ void add_device_column_to_image_gndhwc_3d_i8_instances(
...
@@ -125,7 +125,7 @@ void add_device_column_to_image_gndhwc_3d_i8_instances(
instances
);
instances
);
// NWGC/NHWGC/NDHWGC
// NWGC/NHWGC/NDHWGC
// Image to Column
// Image to Column
//
nhwc
, 1d
//
NWGC
, 1d
void
add_device_image_to_column_nwgc_1d_bf16_instances
(
void
add_device_image_to_column_nwgc_1d_bf16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
1
,
NWGC
,
BF16
,
BF16
,
ImageToColumn
>>>&
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
1
,
NWGC
,
BF16
,
BF16
,
ImageToColumn
>>>&
instances
);
instances
);
...
@@ -141,7 +141,7 @@ void add_device_image_to_column_nwgc_1d_f32_instances(
...
@@ -141,7 +141,7 @@ void add_device_image_to_column_nwgc_1d_f32_instances(
void
add_device_image_to_column_nwgc_1d_i8_instances
(
void
add_device_image_to_column_nwgc_1d_i8_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
1
,
NWGC
,
int8_t
,
int8_t
,
ImageToColumn
>>>&
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
1
,
NWGC
,
int8_t
,
int8_t
,
ImageToColumn
>>>&
instances
);
instances
);
//
nhwc
, 2d
//
NHWGC
, 2d
void
add_device_image_to_column_nhwgc_2d_bf16_instances
(
void
add_device_image_to_column_nhwgc_2d_bf16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
2
,
NHWGC
,
BF16
,
BF16
,
ImageToColumn
>>>&
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
2
,
NHWGC
,
BF16
,
BF16
,
ImageToColumn
>>>&
instances
);
instances
);
...
@@ -158,7 +158,7 @@ void add_device_image_to_column_nhwgc_2d_i8_instances(
...
@@ -158,7 +158,7 @@ void add_device_image_to_column_nhwgc_2d_i8_instances(
std
::
vector
<
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
2
,
NHWGC
,
int8_t
,
int8_t
,
ImageToColumn
>>>&
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
2
,
NHWGC
,
int8_t
,
int8_t
,
ImageToColumn
>>>&
instances
);
instances
);
//
nhwc
, 3d
//
NDHWGC
, 3d
void
add_device_image_to_column_ndhwgc_3d_bf16_instances
(
void
add_device_image_to_column_ndhwgc_3d_bf16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
3
,
NDHWGC
,
BF16
,
BF16
,
ImageToColumn
>>>&
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
3
,
NDHWGC
,
BF16
,
BF16
,
ImageToColumn
>>>&
instances
);
instances
);
...
@@ -177,7 +177,7 @@ void add_device_image_to_column_ndhwgc_3d_i8_instances(
...
@@ -177,7 +177,7 @@ void add_device_image_to_column_ndhwgc_3d_i8_instances(
instances
);
instances
);
// Column to Image
// Column to Image
//
nhwc
, 1d
//
NWGC
, 1d
void
add_device_column_to_image_nwgc_1d_bf16_instances
(
void
add_device_column_to_image_nwgc_1d_bf16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
1
,
NWGC
,
BF16
,
BF16
,
ColumnToImage
>>>&
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
1
,
NWGC
,
BF16
,
BF16
,
ColumnToImage
>>>&
instances
);
instances
);
...
@@ -193,7 +193,7 @@ void add_device_column_to_image_nwgc_1d_f32_instances(
...
@@ -193,7 +193,7 @@ void add_device_column_to_image_nwgc_1d_f32_instances(
void
add_device_column_to_image_nwgc_1d_i8_instances
(
void
add_device_column_to_image_nwgc_1d_i8_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
1
,
NWGC
,
int8_t
,
int8_t
,
ColumnToImage
>>>&
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
1
,
NWGC
,
int8_t
,
int8_t
,
ColumnToImage
>>>&
instances
);
instances
);
//
nhwc
, 2d
//
NHWGC
, 2d
void
add_device_column_to_image_nhwgc_2d_bf16_instances
(
void
add_device_column_to_image_nhwgc_2d_bf16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
2
,
NHWGC
,
BF16
,
BF16
,
ColumnToImage
>>>&
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
2
,
NHWGC
,
BF16
,
BF16
,
ColumnToImage
>>>&
instances
);
instances
);
...
@@ -210,7 +210,7 @@ void add_device_column_to_image_nhwgc_2d_i8_instances(
...
@@ -210,7 +210,7 @@ void add_device_column_to_image_nhwgc_2d_i8_instances(
std
::
vector
<
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
2
,
NHWGC
,
int8_t
,
int8_t
,
ColumnToImage
>>>&
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
2
,
NHWGC
,
int8_t
,
int8_t
,
ColumnToImage
>>>&
instances
);
instances
);
//
nhwc
, 3d
//
NDHWGC
, 3d
void
add_device_column_to_image_ndhwgc_3d_bf16_instances
(
void
add_device_column_to_image_ndhwgc_3d_bf16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
3
,
NDHWGC
,
BF16
,
BF16
,
ColumnToImage
>>>&
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
3
,
NDHWGC
,
BF16
,
BF16
,
ColumnToImage
>>>&
instances
);
instances
);
...
...
library/src/tensor_operation_instance/gpu/image_to_column/device_image_to_column_gndhwc_3d_instance.cpp
View file @
36f54418
// SPDX-License-Identifier: MIT
// SPDX-License-Identifier: MIT
// Copyright (c)
2018-
2023, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2023, Advanced Micro Devices, Inc. All rights reserved.
#include "ck/library/tensor_operation_instance/gpu/conv_tensor_rearrange/device_image_to_column_instance.hpp"
#include "ck/library/tensor_operation_instance/gpu/conv_tensor_rearrange/device_image_to_column_instance.hpp"
#include "ck/library/tensor_operation_instance/add_device_operation_instance.hpp"
#include "ck/library/tensor_operation_instance/add_device_operation_instance.hpp"
...
...
library/src/tensor_operation_instance/gpu/image_to_column/device_image_to_column_gnhwc_2d_instance.cpp
View file @
36f54418
// SPDX-License-Identifier: MIT
// SPDX-License-Identifier: MIT
// Copyright (c)
2018-
2023, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2023, Advanced Micro Devices, Inc. All rights reserved.
#include "ck/library/tensor_operation_instance/gpu/conv_tensor_rearrange/device_image_to_column_instance.hpp"
#include "ck/library/tensor_operation_instance/gpu/conv_tensor_rearrange/device_image_to_column_instance.hpp"
#include "ck/library/tensor_operation_instance/add_device_operation_instance.hpp"
#include "ck/library/tensor_operation_instance/add_device_operation_instance.hpp"
...
...
library/src/tensor_operation_instance/gpu/image_to_column/device_image_to_column_gnwc_1d_instance.cpp
View file @
36f54418
// SPDX-License-Identifier: MIT
// SPDX-License-Identifier: MIT
// Copyright (c)
2018-
2023, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2023, Advanced Micro Devices, Inc. All rights reserved.
#include "ck/library/tensor_operation_instance/gpu/conv_tensor_rearrange/device_image_to_column_instance.hpp"
#include "ck/library/tensor_operation_instance/gpu/conv_tensor_rearrange/device_image_to_column_instance.hpp"
#include "ck/library/tensor_operation_instance/add_device_operation_instance.hpp"
#include "ck/library/tensor_operation_instance/add_device_operation_instance.hpp"
...
...
library/src/tensor_operation_instance/gpu/image_to_column/device_image_to_column_ndhwgc_3d_instance.cpp
View file @
36f54418
// SPDX-License-Identifier: MIT
// SPDX-License-Identifier: MIT
// Copyright (c)
2018-
2023, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2023, Advanced Micro Devices, Inc. All rights reserved.
#include "ck/library/tensor_operation_instance/gpu/conv_tensor_rearrange/device_image_to_column_instance.hpp"
#include "ck/library/tensor_operation_instance/gpu/conv_tensor_rearrange/device_image_to_column_instance.hpp"
#include "ck/library/tensor_operation_instance/add_device_operation_instance.hpp"
#include "ck/library/tensor_operation_instance/add_device_operation_instance.hpp"
...
...
library/src/tensor_operation_instance/gpu/image_to_column/device_image_to_column_nhwgc_2d_instance.cpp
View file @
36f54418
// SPDX-License-Identifier: MIT
// SPDX-License-Identifier: MIT
// Copyright (c)
2018-
2023, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2023, Advanced Micro Devices, Inc. All rights reserved.
#include "ck/library/tensor_operation_instance/gpu/conv_tensor_rearrange/device_image_to_column_instance.hpp"
#include "ck/library/tensor_operation_instance/gpu/conv_tensor_rearrange/device_image_to_column_instance.hpp"
#include "ck/library/tensor_operation_instance/add_device_operation_instance.hpp"
#include "ck/library/tensor_operation_instance/add_device_operation_instance.hpp"
...
...
library/src/tensor_operation_instance/gpu/image_to_column/device_image_to_column_nwgc_1d_instance.cpp
View file @
36f54418
// SPDX-License-Identifier: MIT
// SPDX-License-Identifier: MIT
// Copyright (c)
2018-
2023, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2023, Advanced Micro Devices, Inc. All rights reserved.
#include "ck/library/tensor_operation_instance/gpu/conv_tensor_rearrange/device_image_to_column_instance.hpp"
#include "ck/library/tensor_operation_instance/gpu/conv_tensor_rearrange/device_image_to_column_instance.hpp"
#include "ck/library/tensor_operation_instance/add_device_operation_instance.hpp"
#include "ck/library/tensor_operation_instance/add_device_operation_instance.hpp"
...
...
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