Skip to content

P
Paddle-Lite

PaddlePaddle / Paddle-Lite

通知 331
Star 4
Fork 1
P
Paddle-Lite
提交 b4c5fdb8 编写于 作者: X xiebaiyuan 提交者： GitHub
浏览文件
操作

[LITE][OPENCL][Image] conv 1x1 5x5 7x7 suite (#2998) 

* [LITE][OPENCL][Image] conv 1x1 5x5 7x7 suite

# Conflicts:
#	lite/kernels/opencl/conv_image_compute_test.cc

* [LITE][OPENCL][Image] conv 1x1 5x5 7x7 suite,test=develop

* [LITE][OPENCL][Image] conv 1x1 5x5 7x7 suite,test=develop

* [LITE][OPENCL][Image] conv 1x1 5x5 7x7 suite,test=develop

* [LITE][OPENCL][Image] conv 1x1 5x5 7x7 suite,rm1x1 old,test=develop
上级 2c229275
显示空白变更内容
内联 并排
Showing with 522 addition and 346 deletion
lite/backends/opencl/cl_kernel/image/conv2d_1x1_kernel.cl
浏览文件 @ b4c5fdb8
......@@ -15,7 +15,7 @@ __kernel void conv2d_1x1(__private const int global_size_dim0,
__write_only image2d_t output_image,
__private const int stride,
__private const int offset,
__private const int input_c,
__private const int input_c_block,
__private const int input_c_origin,
__private const int dilation,
__private const int input_width, /* of one block */
......@@ -79,14 +79,14 @@ __kernel void conv2d_1x1(__private const int global_size_dim0,
CL_DTYPE4 output3 = 0.0f;
#endif
int max_w_bound = input_c * input_width;
int burndary_index = input_c * 4 - input_c_origin;
int max_w_bound = input_c_block * input_width;
int burndary_index = input_c_block * 4 - input_c_origin;
bool burndary_index_w =
burndary_index == 1 || burndary_index == 2 || burndary_index == 3;
bool burndary_index_z = burndary_index == 2 || burndary_index == 3;
bool burndary_index_y = burndary_index == 3;
for (int i = 0; i < input_c; ++i) {
for (int i = 0; i < input_c_block; ++i) {
// ------------0---------------
int2 pos_in = (int2)(i * input_width + in_pos_in_one_block0.x,
in_pos_in_one_block0.y);
......@@ -107,11 +107,81 @@ __kernel void conv2d_1x1(__private const int global_size_dim0,
input0.w = select(input0.w, zero, outof_bound && burndary_index_w);
input0.z = select(input0.z, zero, outof_bound && burndary_index_z);
input0.y = select(input0.y, zero, outof_bound && burndary_index_y);
#ifdef DEBUG
if (output_pos0.x == 0 && output_pos0.y == 0) {
printf("i ={ %d, }\n", i);
printf("in={ %f , %f , %f , %f } \n",
convert_float(input0.x),
convert_float(input0.y),
convert_float(input0.z),
convert_float(input0.w));
printf("filter0={ %f , %f , %f , %f } \n",
convert_float(weight0.x),
convert_float(weight0.y),
convert_float(weight0.z),
convert_float(weight0.w));
printf("filter1={ %f , %f , %f , %f } \n",
convert_float(weight1.x),
convert_float(weight1.y),
convert_float(weight1.z),
convert_float(weight1.w));
printf("filter2={ %f , %f , %f , %f } \n",
convert_float(weight2.x),
convert_float(weight2.y),
convert_float(weight2.z),
convert_float(weight2.w));
printf("filter3={ %f , %f , %f , %f } \n",
convert_float(weight3.x),
convert_float(weight3.y),
convert_float(weight3.z),
convert_float(weight3.w));
printf("000---- output={ %f , %f , %f , %f } \n",
convert_float(output0.x),
convert_float(output0.y),
convert_float(output0.z),
convert_float(output0.w));
}
#endif
output0 = mad(input0.x, weight0, output0);
#ifdef DEBUG
if (output_pos0.x == 0 && output_pos0.y == 0) {
printf("111---- output={ %f , %f , %f , %f } \n",
convert_float(output0.x),
convert_float(output0.y),
convert_float(output0.z),
convert_float(output0.w));
}
#endif
output0 = mad(input0.y, weight1, output0);
#ifdef DEBUG
if (output_pos0.x == 0 && output_pos0.y == 0) {
printf("222---- output={ %f , %f , %f , %f } \n",
convert_float(output0.x),
convert_float(output0.y),
convert_float(output0.z),
convert_float(output0.w));
}
#endif
output0 = mad(input0.z, weight2, output0);
#ifdef DEBUG
if (output_pos0.x == 0 && output_pos0.y == 0) {
printf("333---- output={ %f , %f , %f , %f } \n",
convert_float(output0.x),
convert_float(output0.y),
convert_float(output0.z),
convert_float(output0.w));
}
#endif
output0 = mad(input0.w, weight3, output0);
#ifdef DEBUG
if (output_pos0.x == 0 && output_pos0.y == 0) {
printf("444---- output={ %f , %f , %f , %f } \n",
convert_float(output0.x),
convert_float(output0.y),
convert_float(output0.z),
convert_float(output0.w));
}
#endif
// -------------1--------------
pos_in = (int2)(i * input_width + in_pos_in_one_block1.x,
in_pos_in_one_block1.y);
......@@ -171,6 +241,43 @@ __kernel void conv2d_1x1(__private const int global_size_dim0,
output3 = mad(input3.y, weight1, output3);
output3 = mad(input3.z, weight2, output3);
output3 = mad(input3.w, weight3, output3);
#ifdef DEBUG
if (output_pos0.x == 0 && output_pos0.y == 0) {
// printf("i,j,k ={ %d, %d , %d }\n", i,j,k);
printf("i ={ %d, }\n", i);
printf("in={ %f , %f , %f , %f } \n",
convert_float(input0.x),
convert_float(input0.y),
convert_float(input0.z),
convert_float(input0.w));
printf("filter0={ %f , %f , %f , %f } \n",
convert_float(weight0.x),
convert_float(weight0.y),
convert_float(weight0.z),
convert_float(weight0.w));
printf("filter1={ %f , %f , %f , %f } \n",
convert_float(weight1.x),
convert_float(weight1.y),
convert_float(weight1.z),
convert_float(weight1.w));
printf("filter2={ %f , %f , %f , %f } \n",
convert_float(weight2.x),
convert_float(weight2.y),
convert_float(weight2.z),
convert_float(weight2.w));
printf("filter3={ %f , %f , %f , %f } \n",
convert_float(weight3.x),
convert_float(weight3.y),
convert_float(weight3.z),
convert_float(weight3.w));
printf("output={ %f , %f , %f , %f } \n",
convert_float(output0.x),
convert_float(output0.y),
convert_float(output0.z),
convert_float(output0.w));
}
#endif
}
#ifdef BATCH_NORM
......@@ -195,7 +302,6 @@ __kernel void conv2d_1x1(__private const int global_size_dim0,
output1 = activation_type4(output1);
output2 = activation_type4(output2);
output3 = activation_type4(output3);
if (out_w0 < old_w) {
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output_image, output_pos0, output0);
}
......@@ -213,7 +319,8 @@ __kernel void conv2d_1x1(__private const int global_size_dim0,
}
}
__kernel void conv2d_1x1_simple(__private const int global_size_dim0,
__kernel void conv2d_1x1_simple(
__private const int global_size_dim0,
__private const int global_size_dim1,
__private const int global_size_dim2,
__read_only image2d_t input_image,
......@@ -222,7 +329,7 @@ __kernel void conv2d_1x1_simple(__private const int global_size_dim0,
__read_only image2d_t bias,
#endif
#ifdef BATCH_NORM
__read_only image2d_t new_scale,
__read_only image2d_t new_scale,
__read_only image2d_t new_biase,
#endif
__write_only image2d_t output_image,
......@@ -358,13 +465,11 @@ __read_only image2d_t new_scale,
READ_IMG_TYPE(CL_DTYPE_CHAR, new_biase, sampler, (int2)(out_c, 0));
#endif
output0 = activation_type4(output0);
output1 = activation_type4(output1);
output2 = activation_type4(output2);
output3 = activation_type4(output3);
if (out_w0 < old_w) {
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output_image, output_pos0, output0);
}
......
lite/backends/opencl/cl_kernel/image/conv2d_7x7_kernel.cl
浏览文件 @ b4c5fdb8
......@@ -36,10 +36,10 @@ __kernel void conv2d_7x7(__private const int global_size_dim0,
const int batch_index = out_nh / output_height;
const int out_nh_in_one_batch = out_nh % output_height;
const filter_n0 = 4 * out_c + 0;
const filter_n1 = 4 * out_c + 1;
const filter_n2 = 4 * out_c + 2;
const filter_n3 = 4 * out_c + 3;
const int filter_n0 = 4 * out_c + 0;
const int filter_n1 = 4 * out_c + 1;
const int filter_n2 = 4 * out_c + 2;
const int filter_n3 = 4 * out_c + 3;
int2 stride_xy;
stride_xy.x = stride;
......
lite/kernels/opencl/conv_image_compute_test.cc
浏览文件 @ b4c5fdb8
......@@ -15,6 +15,7 @@
#include <gtest/gtest.h>
#include <random>
#include "lite/backends/opencl/cl_image_converter.h"
#include "lite/backends/opencl/target_wrapper.h"
#include "lite/core/op_registry.h"
......@@ -26,7 +27,12 @@ namespace lite {
// #define SHADOW_LOG LOG(INFO)
#define SHADOW_LOG VLOG(4)
#define FP16_MAX_DIFF (1e0)
#define FP16_ABS_DIFF (1e-1)
// #define TEST_CONV_IMAGE_ALL_1
#define TEST_CONV_IMAGE_1x1
#define TEST_CONV_IMAGE_3x3
#define TEST_CONV_IMAGE_5x5
#define TEST_CONV_IMAGE_7x7
template <typename Dtype1, typename Dtype2>
static void conv_basic(const Dtype1* din,
Dtype2* dout,
......@@ -127,6 +133,8 @@ int ConvOutputSize(int input_size,
return output_size;
}
#ifdef TEST_CONV_IMAGE_1x1
// #define PRINT_RESULT
// #define LOOP_TEST
TEST(conv2d, compute_image2d_1x1) {
......@@ -140,24 +148,33 @@ TEST(conv2d, compute_image2d_1x1) {
#ifdef LOOP_TEST
for (int batch_size = 1; batch_size < 4; ++batch_size) {
for (int oc = 4; oc < 10; oc += 1) { // oc
for (int ih = 4; ih < 9; ih += 1) { // ih
for (int oc = 2; oc < 10; oc += 1) { // oc
for (int ih = 2; ih < 9; ih += 1) { // ih
int iw = ih;
for (int iw = 4; iw < 10; iw += 1) { // iw
for (int ic = 4; ic < 10; ic += 1) { // ic
for (int ic = 2; ic < 10; ic += 1) { // ic
for (bool bias_flag : {true, false}) {
for (std::string relu_flag : {"relu"}) {
for (std::string relu_flag : {""}) {
#else
const int batch_size = 1;
const int oc = 4;
const int ih = 8;
const int iw = 8;
const int ic = 4;
const bool bias_flag = true;
const std::string relu_flag = "relu";
const int oc = 2;
const int ih = 3;
const int iw = 3;
const int ic = 2;
const bool bias_flag = false;
const std::string relu_flag = "";
#endif
LOG(INFO) << "---------------------------- "
"conv1x1----------------------- "
"run---------------";
const int oh = ih;
const int ow = iw;
LOG(INFO) << "batch_size: " << batch_size;
LOG(INFO) << "ic: " << ic;
LOG(INFO) << "ih: " << ih;
LOG(INFO) << "iw: " << iw;
LOG(INFO) << "oc: " << oc;
LOG(INFO) << "bias_flag: " << bias_flag;
LOG(INFO) << "relu_flag: " << relu_flag;
SHADOW_LOG << "to get kernel ...";
auto kernels =
......@@ -253,9 +270,15 @@ TEST(conv2d, compute_image2d_1x1) {
for (auto& i : input_v) {
i = gen(engine);
#ifdef TEST_CONV_IMAGE_ALL_1
i = 0.01;
#endif
}
for (auto& f : filter_v) {
f = gen(engine);
#ifdef TEST_CONV_IMAGE_ALL_1
f = 0.01;
#endif
}
SHADOW_LOG << "after gen input and filter ...";
......@@ -268,8 +291,7 @@ TEST(conv2d, compute_image2d_1x1) {
SHADOW_LOG << "filter_dim.production(): "
<< filter_dim.production();
SHADOW_LOG << "out_dim.production(): " << out_dim.production();
SHADOW_LOG << "bias_dim.production(): "
<< bias_dim.production();
SHADOW_LOG << "bias_dim.production(): " << bias_dim.production();
SHADOW_LOG << "4 * input_image_height * input_image_width: "
<< 4 * input_image_height * input_image_width;
SHADOW_LOG << "4 * filter_image_width * filter_image_height: "
......@@ -306,12 +328,20 @@ TEST(conv2d, compute_image2d_1x1) {
// assign filter as target arm
filter.Assign<float, lite::DDim, TARGET(kARM)>(filter_v.data(),
filter_dim);
SHADOW_LOG << " lite输入 input_v ..... ";
for (int i = 0; i < input_v.size(); i++) {
SHADOW_LOG << "(" << i << ")" << input_v[i];
}
SHADOW_LOG << " lite输入 input_image2d ..... ";
for (int i = 0; i < x_image_v.size(); i++) {
SHADOW_LOG << "(" << i << ")" << Half2Float(x_image_v[i]);
}
// auto* filter_image2d =
// filter.mutable_data<half_t, cl::Image2D>(
// filter.mutable_data<uint16_t, cl::Image2D>(
// filter_image_width,
// filter_image_height,
// filter_image_v.data());
SHADOW_LOG << "卷积核: ---- ";
SHADOW_LOG << "卷积核 : ---- ";
for (int i = 0; i < filter_v.size(); i++) {
SHADOW_LOG << "(" << i << ")" << filter_v[i];
}
......@@ -328,6 +358,11 @@ TEST(conv2d, compute_image2d_1x1) {
SHADOW_LOG << "(" << i << ")" << *filter_p2;
filter_p2++;
}
SHADOW_LOG << "卷积核 image : ---- ";
for (int i = 0; i < filter_image_v.size(); i++) {
SHADOW_LOG << "(" << i << ")" << Half2Float(filter_image_v[i]);
}
if (bias_flag) {
for (int i = 0; i < bias_dim.production(); ++i) {
bias_v[i] = static_cast<int>(gen(engine));
......@@ -389,6 +424,14 @@ TEST(conv2d, compute_image2d_1x1) {
output_v.data(),
out_image_shape,
output.dims());
SHADOW_LOG << " lite输出 out_image_v ..... ";
for (int i = 0; i < out_image_v.size(); i++) {
SHADOW_LOG << "(" << i << ")" << Half2Float(out_image_v[i]);
}
SHADOW_LOG << " lite输出 output_v ..... ";
for (int i = 0; i < output_v.size(); i++) {
SHADOW_LOG << "(" << i << ")" << output_v[i];
}
SHADOW_LOG << "mutable_data out_ref_data: ";
// run cpu ref
......@@ -428,8 +471,11 @@ TEST(conv2d, compute_image2d_1x1) {
for (int i = 0; i < out_dim.production(); i++) {
auto relative_diff =
COMPUTE_RELATIVE_DIFF(output_v[i], out_ref_data[i]);
EXPECT_LT(relative_diff, FP16_MAX_DIFF);
if (relative_diff > FP16_MAX_DIFF) {
auto abs_diff = COMPTUE_ABS_DIFF(output_v[i], out_ref_data[i]);
// EXPECT_LT(relative_diff, FP16_MAX_DIFF);
EXPECT_FALSE(relative_diff > FP16_MAX_DIFF &&
abs_diff > FP16_ABS_DIFF);
if (relative_diff > FP16_MAX_DIFF && abs_diff > FP16_ABS_DIFF) {
LOG(FATAL) << "error idx:" << i << "output_v[" << i
<< "]:" << output_v[i] << " "
"out_ref_data["
......@@ -443,14 +489,15 @@ TEST(conv2d, compute_image2d_1x1) {
}
}
}
}
#else
// nothing to do.
#endif
}
#undef LOOP_TEST
#undef PRINT_RESULT
#endif
#ifdef TEST_CONV_IMAGE_3x3
// #define PRINT_RESULT
// #define LOOP_TEST
TEST(conv2d, compute_image2d_3x3) {
......@@ -475,7 +522,7 @@ TEST(conv2d, compute_image2d_3x3) {
const int dilation = 1;
#if 0 // small scale with group, but result of cpu reference is wrong
const int stride = 2;
const int stride = 2;
const int group = 2;
const int batch_size = 1;
const int ic = 1;
......@@ -823,6 +870,10 @@ TEST(conv2d, compute_image2d_3x3) {
#undef LOOP_TEST
#undef PRINT_RESULT
#endif
#ifdef TEST_CONV_IMAGE_5x5
// #define PRINT_RESULT
// #define LOOP_TEST
TEST(conv2d, compute_image2d_5x5) {
......@@ -839,7 +890,7 @@ TEST(conv2d, compute_image2d_5x5) {
for (int oc = 1; oc < 10; oc += 1) { // oc
for (int ih = 5; ih < 9; ih += 1) { // ih
int iw = ih;
for (int ic = 1; ic < 10; ic += 1) { // ic
for (int ic = 2; ic < 10; ic += 1) { // ic
for (bool bias_flag : {true, false}) {
for (std::string relu_flag : {/*true,*/ "relu"}) {
#else
......@@ -847,7 +898,8 @@ TEST(conv2d, compute_image2d_5x5) {
const int oc = 1;
const int ih = 5;
const int iw = 5;
const int ic = 1;
// ic = 1 会进入depthwise的路由 .
const int ic = 2;
const bool bias_flag = true;
const std::string relu_flag = "relu";
#endif
......@@ -1139,8 +1191,10 @@ TEST(conv2d, compute_image2d_5x5) {
for (int i = 0; i < out_dim.production(); i++) {
auto relative_diff =
COMPUTE_RELATIVE_DIFF(output_v[i], out_ref_data[i]);
EXPECT_LT(relative_diff, FP16_MAX_DIFF);
if (relative_diff > FP16_MAX_DIFF) {
auto abs_diff = COMPTUE_ABS_DIFF(output_v[i], out_ref_data[i]);
EXPECT_FALSE(relative_diff > FP16_MAX_DIFF &&
abs_diff > FP16_ABS_DIFF);
if (relative_diff > FP16_MAX_DIFF && abs_diff > FP16_ABS_DIFF) {
LOG(FATAL) << "error idx:" << i << "output_v[" << i
<< "]:" << output_v[i] << " "
"out_ref_data["
......@@ -1161,13 +1215,16 @@ TEST(conv2d, compute_image2d_5x5) {
}
#undef LOOP_TEST
#undef PRINT_RESULT
#endif
#ifdef TEST_CONV_IMAGE_7x7
#undef FP16_ABS_DIFF
#define FP16_ABS_DIFF (1e0)
// #define LOOP_TEST
TEST(conv2d, compute_image2d_7x7) {
// conv infos
const int ksize = 7;
const int stride = 1;
const int pad = 2;
const int pad = 3;
const int group = 1;
const int dilation = 1;
// int loop_cnt = 0;
......@@ -1177,7 +1234,7 @@ TEST(conv2d, compute_image2d_7x7) {
for (int oc = 1; oc < 10; oc += 1) { // oc
for (int ih = 7; ih < 15; ih += 1) { // ih
int iw = ih;
for (int ic = 1; ic < 10; ic += 1) { // ic
for (int ic = 2; ic < 10; ic += 1) { // ic
for (bool bias_flag : {true, false}) {
for (std::string relu_flag : {"relu"}) {
#else
......@@ -1185,7 +1242,8 @@ TEST(conv2d, compute_image2d_7x7) {
const int oc = 1;
const int ih = 7;
const int iw = 7;
const int ic = 1;
// ic = 1会进入 depthwise路由
const int ic = 2;
const bool bias_flag = false;
const std::string relu_flag = "";
#endif
......@@ -1286,11 +1344,15 @@ TEST(conv2d, compute_image2d_7x7) {
SHADOW_LOG << "gen input and filter ...";
for (auto& i : input_v) {
i = gen(engine);
// i = 1;
#ifdef TEST_CONV_IMAGE_ALL_1
i = 1;
#endif
}
for (auto& f : filter_v) {
f = gen(engine);
// f = 1;
#ifdef TEST_CONV_IMAGE_ALL_1
f = 1;
#endif
}
LOG(INFO) << "bias: " << bias_flag;
LOG(INFO) << "relu: " << relu_flag;
......@@ -1340,7 +1402,7 @@ TEST(conv2d, compute_image2d_7x7) {
}
SHADOW_LOG << "输入image : ---- ";
for (int i = 0; i < x_image_v.size(); i++) {
SHADOW_LOG << "(" << i << ")" << x_image_v[i];
SHADOW_LOG << "(" << i << ")" << Half2Float(x_image_v[i]);
}
SHADOW_LOG << "set mapped filter ...";
CLImageConverterFolder folder_convertor;
......@@ -1353,7 +1415,7 @@ TEST(conv2d, compute_image2d_7x7) {
}
SHADOW_LOG << "卷积核image: ---- ";
for (int i = 0; i < filter_image_v.size(); i++) {
SHADOW_LOG << "(" << i << ")" << filter_image_v[i];
SHADOW_LOG << "(" << i << ")" << Half2Float(filter_image_v[i]);
}
auto* input_image2d = input.mutable_data<half_t, cl::Image2D>(
input_image_width, input_image_height, x_image_v.data());
......@@ -1437,7 +1499,7 @@ TEST(conv2d, compute_image2d_7x7) {
SHADOW_LOG << "输出image: ---- ";
for (int i = 0; i < out_image_v.size(); i++) {
SHADOW_LOG << "(" << i << ")" << out_image_v[i];
SHADOW_LOG << "(" << i << ")" << Half2Float(out_image_v[i]);
}
SHADOW_LOG << "mutable_data out_ref_data: ";
......@@ -1478,8 +1540,10 @@ TEST(conv2d, compute_image2d_7x7) {
for (int i = 0; i < out_dim.production(); i++) {
auto relative_diff =
COMPUTE_RELATIVE_DIFF(output_v[i], out_ref_data[i]);
EXPECT_LT(relative_diff, FP16_MAX_DIFF);
if (relative_diff > FP16_MAX_DIFF) {
auto abs_diff = COMPTUE_ABS_DIFF(output_v[i], out_ref_data[i]);
EXPECT_FALSE(relative_diff > FP16_MAX_DIFF &&
abs_diff > FP16_ABS_DIFF);
if (relative_diff > FP16_MAX_DIFF && abs_diff > FP16_ABS_DIFF) {
LOG(FATAL) << "error idx:" << i << "output_v[" << i
<< "]:" << output_v[i] << " "
"out_ref_data["
......@@ -1500,7 +1564,14 @@ TEST(conv2d, compute_image2d_7x7) {
}
#undef LOOP_TEST
#undef PRINT_RESULT
#endif
#undef SHADOW_LOG
#undef TEST_CONV_IMAGE_1x1
#undef TEST_CONV_IMAGE_3x3
#undef TEST_CONV_IMAGE_5x5
#undef TEST_CONV_IMAGE_7x7
#undef TEST_CONV_IMAGE_ALL_1
} // namespace lite
} // namespace paddle
......
Markdown is supported
0% 或 .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑！
想要评论请 注册 或