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
63fedbd0
Commit
63fedbd0
authored
Sep 21, 2023
by
Bartlomiej Kocot
Browse files
Disable tests for disabled dtypes
parent
c0c6fa59
Changes
7
Hide whitespace changes
Inline
Side-by-side
Showing
7 changed files
with
119 additions
and
34 deletions
+119
-34
library/src/tensor_operation_instance/gpu/column_to_image/device_column_to_image_nhwc_1d_instance.cpp
...lumn_to_image/device_column_to_image_nhwc_1d_instance.cpp
+16
-0
library/src/tensor_operation_instance/gpu/column_to_image/device_column_to_image_nhwc_2d_instance.cpp
...lumn_to_image/device_column_to_image_nhwc_2d_instance.cpp
+16
-0
library/src/tensor_operation_instance/gpu/column_to_image/device_column_to_image_nhwc_3d_instance.cpp
...lumn_to_image/device_column_to_image_nhwc_3d_instance.cpp
+16
-0
library/src/tensor_operation_instance/gpu/image_to_column/device_image_to_column_nhwc_1d_instance.cpp
...age_to_column/device_image_to_column_nhwc_1d_instance.cpp
+8
-0
library/src/tensor_operation_instance/gpu/image_to_column/device_image_to_column_nhwc_2d_instance.cpp
...age_to_column/device_image_to_column_nhwc_2d_instance.cpp
+8
-0
library/src/tensor_operation_instance/gpu/image_to_column/device_image_to_column_nhwc_3d_instance.cpp
...age_to_column/device_image_to_column_nhwc_3d_instance.cpp
+8
-0
test/conv_tensor_rearrange/test_conv_tensor_rearrange.cpp
test/conv_tensor_rearrange/test_conv_tensor_rearrange.cpp
+47
-34
No files found.
library/src/tensor_operation_instance/gpu/column_to_image/device_column_to_image_nhwc_1d_instance.cpp
View file @
63fedbd0
...
...
@@ -15,28 +15,44 @@ void add_device_column_to_image_nhwc_1d_bf16_instances(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
1
,
GNWC
,
BF16
,
BF16
,
ColumnToImage
>>>&
instances
)
{
#ifdef CK_ENABLE_BF16
add_device_operation_instances
(
instances
,
device_column_to_image_bf16_instances
<
1
,
GNWC
>
{});
#else
ignore
=
instances
;
#endif
}
void
add_device_column_to_image_nhwc_1d_f16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
1
,
GNWC
,
F16
,
F16
,
ColumnToImage
>>>&
instances
)
{
#ifdef CK_ENABLE_FP16
add_device_operation_instances
(
instances
,
device_column_to_image_f16_instances
<
1
,
GNWC
>
{});
#else
ignore
=
instances
;
#endif
}
void
add_device_column_to_image_nhwc_1d_f32_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
1
,
GNWC
,
F32
,
F32
,
ColumnToImage
>>>&
instances
)
{
#ifdef CK_ENABLE_FP32
add_device_operation_instances
(
instances
,
device_column_to_image_f32_instances
<
1
,
GNWC
>
{});
#else
ignore
=
instances
;
#endif
}
void
add_device_column_to_image_nhwc_1d_i8_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
1
,
GNWC
,
int8_t
,
int8_t
,
ColumnToImage
>>>&
instances
)
{
#ifdef CK_ENABLE_INT8
add_device_operation_instances
(
instances
,
device_column_to_image_i8_instances
<
1
,
GNWC
>
{});
#else
ignore
=
instances
;
#endif
}
}
// namespace instance
...
...
library/src/tensor_operation_instance/gpu/column_to_image/device_column_to_image_nhwc_2d_instance.cpp
View file @
63fedbd0
...
...
@@ -15,21 +15,33 @@ void add_device_column_to_image_nhwc_2d_bf16_instances(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
2
,
GNHWC
,
BF16
,
BF16
,
ColumnToImage
>>>&
instances
)
{
#ifdef CK_ENABLE_BF16
add_device_operation_instances
(
instances
,
device_column_to_image_bf16_instances
<
2
,
GNHWC
>
{});
#else
ignore
=
instances
;
#endif
}
void
add_device_column_to_image_nhwc_2d_f16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
2
,
GNHWC
,
F16
,
F16
,
ColumnToImage
>>>&
instances
)
{
#ifdef CK_ENABLE_FP16
add_device_operation_instances
(
instances
,
device_column_to_image_f16_instances
<
2
,
GNHWC
>
{});
#else
ignore
=
instances
;
#endif
}
void
add_device_column_to_image_nhwc_2d_f32_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
2
,
GNHWC
,
F32
,
F32
,
ColumnToImage
>>>&
instances
)
{
#ifdef CK_ENABLE_FP32
add_device_operation_instances
(
instances
,
device_column_to_image_f32_instances
<
2
,
GNHWC
>
{});
#else
ignore
=
instances
;
#endif
}
void
add_device_column_to_image_nhwc_2d_i8_instances
(
...
...
@@ -37,7 +49,11 @@ void add_device_column_to_image_nhwc_2d_i8_instances(
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
2
,
GNHWC
,
int8_t
,
int8_t
,
ColumnToImage
>>>&
instances
)
{
#ifdef CK_ENABLE_INT8
add_device_operation_instances
(
instances
,
device_column_to_image_i8_instances
<
2
,
GNHWC
>
{});
#else
ignore
=
instances
;
#endif
}
}
// namespace instance
...
...
library/src/tensor_operation_instance/gpu/column_to_image/device_column_to_image_nhwc_3d_instance.cpp
View file @
63fedbd0
...
...
@@ -15,21 +15,33 @@ void add_device_column_to_image_nhwc_3d_bf16_instances(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
3
,
GNDHWC
,
BF16
,
BF16
,
ColumnToImage
>>>&
instances
)
{
#ifdef CK_ENABLE_BF16
add_device_operation_instances
(
instances
,
device_column_to_image_bf16_instances
<
3
,
GNDHWC
>
{});
#else
ignore
=
instances
;
#endif
}
void
add_device_column_to_image_nhwc_3d_f16_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
3
,
GNDHWC
,
F16
,
F16
,
ColumnToImage
>>>&
instances
)
{
#ifdef CK_ENABLE_FP16
add_device_operation_instances
(
instances
,
device_column_to_image_f16_instances
<
3
,
GNDHWC
>
{});
#else
ignore
=
instances
;
#endif
}
void
add_device_column_to_image_nhwc_3d_f32_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
3
,
GNDHWC
,
F32
,
F32
,
ColumnToImage
>>>&
instances
)
{
#ifdef CK_ENABLE_FP32
add_device_operation_instances
(
instances
,
device_column_to_image_f32_instances
<
3
,
GNDHWC
>
{});
#else
ignore
=
instances
;
#endif
}
void
add_device_column_to_image_nhwc_3d_i8_instances
(
...
...
@@ -37,7 +49,11 @@ void add_device_column_to_image_nhwc_3d_i8_instances(
std
::
unique_ptr
<
DeviceConvTensorRearrange
<
3
,
GNDHWC
,
int8_t
,
int8_t
,
ColumnToImage
>>>&
instances
)
{
#ifdef CK_ENABLE_INT8
add_device_operation_instances
(
instances
,
device_column_to_image_i8_instances
<
3
,
GNDHWC
>
{});
#else
ignore
=
instances
;
#endif
}
}
// namespace instance
...
...
library/src/tensor_operation_instance/gpu/image_to_column/device_image_to_column_nhwc_1d_instance.cpp
View file @
63fedbd0
...
...
@@ -17,6 +17,8 @@ void add_device_image_to_column_nhwc_1d_bf16_instances(
{
#ifdef CK_ENABLE_BF16
add_device_operation_instances
(
instances
,
device_image_to_column_bf16_instances
<
1
,
GNWC
>
{});
#else
ignore
=
instances
;
#endif
}
...
...
@@ -26,6 +28,8 @@ void add_device_image_to_column_nhwc_1d_f16_instances(
{
#ifdef CK_ENABLE_FP16
add_device_operation_instances
(
instances
,
device_image_to_column_f16_instances
<
1
,
GNWC
>
{});
#else
ignore
=
instances
;
#endif
}
...
...
@@ -35,6 +39,8 @@ void add_device_image_to_column_nhwc_1d_f32_instances(
{
#ifdef CK_ENABLE_FP32
add_device_operation_instances
(
instances
,
device_image_to_column_f32_instances
<
1
,
GNWC
>
{});
#else
ignore
=
instances
;
#endif
}
...
...
@@ -44,6 +50,8 @@ void add_device_image_to_column_nhwc_1d_i8_instances(
{
#ifdef CK_ENABLE_INT8
add_device_operation_instances
(
instances
,
device_image_to_column_i8_instances
<
1
,
GNWC
>
{});
#else
ignore
=
instances
;
#endif
}
...
...
library/src/tensor_operation_instance/gpu/image_to_column/device_image_to_column_nhwc_2d_instance.cpp
View file @
63fedbd0
...
...
@@ -17,6 +17,8 @@ void add_device_image_to_column_nhwc_2d_bf16_instances(
{
#ifdef CK_ENABLE_BF16
add_device_operation_instances
(
instances
,
device_image_to_column_bf16_instances
<
2
,
GNHWC
>
{});
#else
ignore
=
instances
;
#endif
}
...
...
@@ -26,6 +28,8 @@ void add_device_image_to_column_nhwc_2d_f16_instances(
{
#ifdef CK_ENABLE_FP16
add_device_operation_instances
(
instances
,
device_image_to_column_f16_instances
<
2
,
GNHWC
>
{});
#else
ignore
=
instances
;
#endif
}
...
...
@@ -35,6 +39,8 @@ void add_device_image_to_column_nhwc_2d_f32_instances(
{
#ifdef CK_ENABLE_FP32
add_device_operation_instances
(
instances
,
device_image_to_column_f32_instances
<
2
,
GNHWC
>
{});
#else
ignore
=
instances
;
#endif
}
...
...
@@ -45,6 +51,8 @@ void add_device_image_to_column_nhwc_2d_i8_instances(
{
#ifdef CK_ENABLE_INT8
add_device_operation_instances
(
instances
,
device_image_to_column_i8_instances
<
2
,
GNHWC
>
{});
#else
ignore
=
instances
;
#endif
}
...
...
library/src/tensor_operation_instance/gpu/image_to_column/device_image_to_column_nhwc_3d_instance.cpp
View file @
63fedbd0
...
...
@@ -17,6 +17,8 @@ void add_device_image_to_column_nhwc_3d_bf16_instances(
{
#ifdef CK_ENABLE_BF16
add_device_operation_instances
(
instances
,
device_image_to_column_bf16_instances
<
3
,
GNDHWC
>
{});
#else
ignore
=
instances
;
#endif
}
...
...
@@ -26,6 +28,8 @@ void add_device_image_to_column_nhwc_3d_f16_instances(
{
#ifdef CK_ENABLE_FP16
add_device_operation_instances
(
instances
,
device_image_to_column_f16_instances
<
3
,
GNDHWC
>
{});
#else
ignore
=
instances
;
#endif
}
...
...
@@ -35,6 +39,8 @@ void add_device_image_to_column_nhwc_3d_f32_instances(
{
#ifdef CK_ENABLE_FP32
add_device_operation_instances
(
instances
,
device_image_to_column_f32_instances
<
3
,
GNDHWC
>
{});
#else
ignore
=
instances
;
#endif
}
...
...
@@ -45,6 +51,8 @@ void add_device_image_to_column_nhwc_3d_i8_instances(
{
#ifdef CK_ENABLE_INT8
add_device_operation_instances
(
instances
,
device_image_to_column_i8_instances
<
3
,
GNDHWC
>
{});
#else
ignore
=
instances
;
#endif
}
...
...
test/conv_tensor_rearrange/test_conv_tensor_rearrange.cpp
View file @
63fedbd0
...
...
@@ -15,14 +15,12 @@ template <typename Tuple>
class
TestConvTensorRearrange
:
public
::
testing
::
Test
{
protected:
using
InDataType
=
std
::
tuple_element_t
<
0
,
Tuple
>
;
using
OutDataType
=
std
::
tuple_element_t
<
1
,
Tuple
>
;
using
ImLayout
=
std
::
tuple_element_t
<
2
,
Tuple
>
;
using
ConvTensorRearrangeOp
=
std
::
tuple_element_t
<
3
,
Tuple
>
;
using
ImLayout
=
std
::
tuple_element_t
<
0
,
Tuple
>
;
using
ConvTensorRearrangeOp
=
std
::
tuple_element_t
<
1
,
Tuple
>
;
std
::
vector
<
ck
::
utils
::
conv
::
ConvParam
>
conv_params
;
template
<
ck
::
index_t
NDimSpatial
>
template
<
ck
::
index_t
NDimSpatial
,
typename
InDataType
,
typename
OutDataType
>
void
Run
()
{
EXPECT_FALSE
(
conv_params
.
empty
());
...
...
@@ -47,32 +45,14 @@ class TestConvTensorRearrange : public ::testing::Test
using
namespace
ck
::
tensor_layout
::
convolution
;
using
namespace
ck
::
conv_tensor_rearrange_op
;
using
KernelTypes1d
=
::
testing
::
Types
<
std
::
tuple
<
float
,
float
,
GNWC
,
ImageToColumn
>
,
std
::
tuple
<
ck
::
bhalf_t
,
ck
::
bhalf_t
,
GNWC
,
ImageToColumn
>
,
std
::
tuple
<
ck
::
half_t
,
ck
::
half_t
,
GNWC
,
ImageToColumn
>
,
std
::
tuple
<
int8_t
,
int8_t
,
GNWC
,
ImageToColumn
>
,
std
::
tuple
<
float
,
float
,
GNWC
,
ColumnToImage
>
,
std
::
tuple
<
ck
::
bhalf_t
,
ck
::
bhalf_t
,
GNWC
,
ColumnToImage
>
,
std
::
tuple
<
ck
::
half_t
,
ck
::
half_t
,
GNWC
,
ColumnToImage
>
,
std
::
tuple
<
int8_t
,
int8_t
,
GNWC
,
ColumnToImage
>>
;
using
KernelTypes2d
=
::
testing
::
Types
<
std
::
tuple
<
float
,
float
,
GNHWC
,
ImageToColumn
>
,
std
::
tuple
<
ck
::
bhalf_t
,
ck
::
bhalf_t
,
GNHWC
,
ImageToColumn
>
,
std
::
tuple
<
ck
::
half_t
,
ck
::
half_t
,
GNHWC
,
ImageToColumn
>
,
std
::
tuple
<
int8_t
,
int8_t
,
GNHWC
,
ImageToColumn
>
,
std
::
tuple
<
float
,
float
,
GNHWC
,
ColumnToImage
>
,
std
::
tuple
<
ck
::
bhalf_t
,
ck
::
bhalf_t
,
GNHWC
,
ColumnToImage
>
,
std
::
tuple
<
ck
::
half_t
,
ck
::
half_t
,
GNHWC
,
ColumnToImage
>
,
std
::
tuple
<
int8_t
,
int8_t
,
GNHWC
,
ColumnToImage
>>
;
using
KernelTypes3d
=
::
testing
::
Types
<
std
::
tuple
<
float
,
float
,
GNDHWC
,
ImageToColumn
>
,
std
::
tuple
<
ck
::
bhalf_t
,
ck
::
bhalf_t
,
GNDHWC
,
ImageToColumn
>
,
std
::
tuple
<
ck
::
half_t
,
ck
::
half_t
,
GNDHWC
,
ImageToColumn
>
,
std
::
tuple
<
int8_t
,
int8_t
,
GNDHWC
,
ImageToColumn
>
,
std
::
tuple
<
float
,
float
,
GNDHWC
,
ColumnToImage
>
,
std
::
tuple
<
ck
::
bhalf_t
,
ck
::
bhalf_t
,
GNDHWC
,
ColumnToImage
>
,
std
::
tuple
<
ck
::
half_t
,
ck
::
half_t
,
GNDHWC
,
ColumnToImage
>
,
std
::
tuple
<
int8_t
,
int8_t
,
GNDHWC
,
ColumnToImage
>>
;
using
KernelTypes1d
=
::
testing
::
Types
<
std
::
tuple
<
GNWC
,
ImageToColumn
>
,
std
::
tuple
<
GNWC
,
ColumnToImage
>>
;
using
KernelTypes2d
=
::
testing
::
Types
<
std
::
tuple
<
GNHWC
,
ImageToColumn
>
,
std
::
tuple
<
GNHWC
,
ColumnToImage
>>
;
using
KernelTypes3d
=
::
testing
::
Types
<
std
::
tuple
<
GNDHWC
,
ImageToColumn
>
,
std
::
tuple
<
GNDHWC
,
ColumnToImage
>>
;
template
<
typename
Tuple
>
class
TestConvTensorRearrange1d
:
public
TestConvTensorRearrange
<
Tuple
>
...
...
@@ -107,7 +87,18 @@ TYPED_TEST(TestConvTensorRearrange1d, Test1D)
this
->
conv_params
.
push_back
({
1
,
1
,
1
,
1
,
4
,
{
3
},
{
28
},
{
2
},
{
1
},
{
1
},
{
1
}});
// dilation != 1
this
->
conv_params
.
push_back
({
1
,
1
,
1
,
1
,
4
,
{
3
},
{
28
},
{
1
},
{
2
},
{
1
},
{
1
}});
this
->
template
Run
<
1
>();
#ifdef CK_ENABLE_FP32
this
->
template
Run
<
1
,
float
,
float
>();
#endif
#ifdef CK_ENABLE_BF16
this
->
template
Run
<
1
,
ck
::
bhalf_t
,
ck
::
bhalf_t
>();
#endif
#ifdef CK_ENABLE_FP16
this
->
template
Run
<
1
,
ck
::
half_t
,
ck
::
half_t
>();
#endif
#ifdef CK_ENABLE_INT8
this
->
template
Run
<
1
,
int8_t
,
int8_t
>();
#endif
}
TYPED_TEST
(
TestConvTensorRearrange2d
,
Test2D
)
...
...
@@ -122,7 +113,18 @@ TYPED_TEST(TestConvTensorRearrange2d, Test2D)
this
->
conv_params
.
push_back
({
2
,
1
,
64
,
1
,
64
,
{
1
,
1
},
{
3
,
3
},
{
1
,
1
},
{
1
,
1
},
{
0
,
0
},
{
0
,
0
}});
this
->
conv_params
.
push_back
(
{
2
,
1
,
64
,
1
,
64
,
{
3
,
3
},
{
28
,
28
},
{
2
,
2
},
{
2
,
2
},
{
1
,
1
},
{
1
,
1
}});
this
->
template
Run
<
2
>();
#ifdef CK_ENABLE_FP32
this
->
template
Run
<
2
,
float
,
float
>();
#endif
#ifdef CK_ENABLE_BF16
this
->
template
Run
<
2
,
ck
::
bhalf_t
,
ck
::
bhalf_t
>();
#endif
#ifdef CK_ENABLE_FP16
this
->
template
Run
<
2
,
ck
::
half_t
,
ck
::
half_t
>();
#endif
#ifdef CK_ENABLE_INT8
this
->
template
Run
<
2
,
int8_t
,
int8_t
>();
#endif
}
TYPED_TEST
(
TestConvTensorRearrange3d
,
Test3D
)
...
...
@@ -136,5 +138,16 @@ TYPED_TEST(TestConvTensorRearrange3d, Test3D)
{
3
,
1
,
32
,
1
,
64
,
{
1
,
1
,
1
},
{
3
,
3
,
3
},
{
1
,
1
,
1
},
{
1
,
1
,
1
},
{
0
,
0
,
0
},
{
0
,
0
,
0
}});
this
->
conv_params
.
push_back
(
{
3
,
1
,
64
,
1
,
64
,
{
3
,
3
,
3
},
{
14
,
14
,
14
},
{
2
,
2
,
2
},
{
2
,
2
,
2
},
{
1
,
1
,
1
},
{
1
,
1
,
1
}});
this
->
template
Run
<
3
>();
#ifdef CK_ENABLE_FP32
this
->
template
Run
<
3
,
float
,
float
>();
#endif
#ifdef CK_ENABLE_BF16
this
->
template
Run
<
3
,
ck
::
bhalf_t
,
ck
::
bhalf_t
>();
#endif
#ifdef CK_ENABLE_FP16
this
->
template
Run
<
3
,
ck
::
half_t
,
ck
::
half_t
>();
#endif
#ifdef CK_ENABLE_INT8
this
->
template
Run
<
3
,
int8_t
,
int8_t
>();
#endif
}
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