Skip to content

P
Paddle-Lite

PaddlePaddle / Paddle-Lite

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

[LITE][OPENCL]Fix opencl (#3433) 

* [lite][opencl] remove event with clfinish, add strict check for cl warning. add conv 3x3opt fallback opt layout cast ,test=develop

* [LITE][OPENCL]rm event in element_add_buffer_compute test=develop

* [LITE][OPENCL]suite cl_functions_test.cc test=develop

* [LITE][OPENCL] suite cl_common.sh lint check test=develop

* [LITE][OPENCL] suite conv_image_compute.cc lint check test=develop

* [LITE][OPENCL] suite cl_wait_list() lint check test=develop
上级 48ff11f9
展开全部 隐藏空白更改
内联 并排
Showing with 1074 addition and 1207 deletion
lite/backends/opencl/cl_functions_test.cc
浏览文件 @ 44fc847d
......@@ -100,16 +100,18 @@ TEST(cl_test, kernel_test) {
size_t width = in_image.ImageWidth();
size_t height = in_image.ImageHeight();
auto global_work_size = cl::NDRange{width, height};
cl::Event event;
status = context->GetCommandQueue().enqueueNDRangeKernel(
kernel, cl::NullRange, global_work_size, cl::NullRange, nullptr, &event);
kernel, cl::NullRange, global_work_size, cl::NullRange, nullptr, nullptr);
CL_CHECK_FATAL(status);
status = context->GetCommandQueue().finish();
CL_CHECK_FATAL(status);
#if 0
double start_nanos = event.getProfilingInfo<CL_PROFILING_COMMAND_START>();
double stop_nanos = event.getProfilingInfo<CL_PROFILING_COMMAND_END>();
double elapsed_micros = (stop_nanos - start_nanos) / 1000.0;
LOG(INFO) << "Kernel Run Cost Time: " << elapsed_micros << " us.";
#endif
LOG(INFO) << out_image;
}
......
lite/backends/opencl/cl_image_converter.cc
浏览文件 @ 44fc847d
......@@ -73,7 +73,7 @@ void CLImageConverterDefault::NCHWToImage(float *nchw,
i2 += 4;
p++;
} else {
image[i2] = 0.0;
image[i2] = Float2Half(0.f);
i2 += 4;
}
}
......@@ -261,7 +261,7 @@ void CLImageConverterNWBlock::NCHWToImage(float *tensor,
image[index] = Float2Half(*p);
p++;
} else {
image[index] = 0.0;
image[index] = Float2Half(0.f);
}
if (index >= (width * height * 4)) {
LOG(INFO) << " index out of range ";
......
lite/backends/opencl/cl_kernel/cl_common.h
浏览文件 @ 44fc847d
......@@ -11,7 +11,6 @@ distributed under the License is distributed on an "AS IS" BASIS,
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and
limitations under the License. */
#pragma once
/////////////////////////////////
......@@ -108,7 +107,8 @@ inline CL_DTYPE4 activation_type4(CL_DTYPE4 in
#endif
#ifdef RELU6
output = clamp(in, (CL_DTYPE4)0, (CL_DTYPE4)6);
in = fmax((CL_DTYPE4)(0.0f, 0.0f, 0.0f, 0.0f), in);
output = fmin((CL_DTYPE4)(6.0f, 6.0f, 6.0f, 6.0f), in);
#endif
return output;
}
lite/backends/opencl/cl_kernel/image/activation_kernel.cl
浏览文件 @ 44fc847d
......@@ -14,36 +14,30 @@ limitations under the License. */
#include <cl_common.h>
__kernel void relu(__read_only image2d_t input,
__write_only image2d_t output,
__private const float threshold,
__private const float scale) {
const int x = get_global_id(0); // image_width
const int y = get_global_id(1); // image_height
const int x = get_global_id(0); // image_width
const int y = get_global_id(1); // image_height
const sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE |
CLK_ADDRESS_CLAMP |
CLK_FILTER_NEAREST;
const sampler_t sampler =
CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
CL_DTYPE4 in = READ_IMG_TYPE(CL_DTYPE_CHAR, input, sampler, (int2)(x, y));
in = max((CL_DTYPE4)(0.0f), in);
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, (int2)(x, y), in);
}
__kernel void relu6(__read_only image2d_t input,
__write_only image2d_t output,
__private const float threshold,
__private const float scale){
__private const float scale) {
const int x = get_global_id(0);
const int y = get_global_id(1);
const sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE |
CLK_ADDRESS_CLAMP |
CLK_FILTER_NEAREST;
const sampler_t sampler =
CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
CL_DTYPE4 in = READ_IMG_TYPE(CL_DTYPE_CHAR, input, sampler, (int2)(x, y));
in = max((CL_DTYPE4)(0.0f, 0.0f, 0.0f, 0.0f), in);
......@@ -51,7 +45,6 @@ __kernel void relu6(__read_only image2d_t input,
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, (int2)(x, y), in);
}
__kernel void sigmoid(__read_only image2d_t input,
__write_only image2d_t output,
__private const float threshold,
......@@ -64,70 +57,66 @@ __kernel void sigmoid(__read_only image2d_t input,
CL_DTYPE4 in = READ_IMG_TYPE(CL_DTYPE_CHAR, input, sampler, (int2)(x, y));
CL_DTYPE4 out;
out.x = 1.0 / (1.0 + pow(2.71828182, -1.0 * (float)(in.x)));
out.y = 1.0 / (1.0 + pow(2.71828182, -1.0 * (float)(in.y)));
out.z = 1.0 / (1.0 + pow(2.71828182, -1.0 * (float)(in.z)));
out.w = 1.0 / (1.0 + pow(2.71828182, -1.0 * (float)(in.w)));
out.x = (CL_DTYPE)(1.0f / (1.0f + pow(2.71828182f, -1.0f * (float)(in.x))));
out.y = (CL_DTYPE)(1.0f / (1.0f + pow(2.71828182f, -1.0f * (float)(in.y))));
out.z = (CL_DTYPE)(1.0f / (1.0f + pow(2.71828182f, -1.0f * (float)(in.z))));
out.w = (CL_DTYPE)(1.0f / (1.0f + pow(2.71828182f, -1.0f * (float)(in.w))));
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, (int2)(x, y), out);
}
__kernel void leaky_relu(__read_only image2d_t input,
__write_only image2d_t output,
__private const float threshold,
__private const float scale) {
__write_only image2d_t output,
__private const float threshold,
__private const float scale) {
const int x = get_global_id(0);
const int y = get_global_id(1);
const sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE |
CLK_ADDRESS_CLAMP |
CLK_FILTER_NEAREST;
const sampler_t sampler =
CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
CL_DTYPE4 in = READ_IMG_TYPE(CL_DTYPE_CHAR, input, sampler, (int2)(x, y));
CL_DTYPE4 s_val = CONVERT_TYPE_TO(scale, CL_DTYPE) * in;
if (in.x < 0.0f){
if (in.x < 0.0f) {
in.x = s_val.x;
}
if (in.y < 0.0f){
if (in.y < 0.0f) {
in.y = s_val.y;
}
if (in.z < 0.0f){
if (in.z < 0.0f) {
in.z = s_val.z;
}
if (in.w < 0.0f){
if (in.w < 0.0f) {
in.w = s_val.w;
}
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, (int2)(x, y), in);
}
__kernel void tanh_act(__read_only image2d_t input,
__write_only image2d_t output,
__private const float threshold,
__private const float scale) {
const int x = get_global_id(0); // image_width
const int y = get_global_id(1); // image_height
__write_only image2d_t output,
__private const float threshold,
__private const float scale) {
const int x = get_global_id(0); // image_width
const int y = get_global_id(1); // image_height
const sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE |
CLK_ADDRESS_CLAMP |
CLK_FILTER_NEAREST;
const sampler_t sampler =
CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
CL_DTYPE4 in = READ_IMG_TYPE(CL_DTYPE_CHAR, input, sampler, (int2)(x, y));
CL_DTYPE4 out= (exp(in) - exp(-in))/ (exp(in) + exp(-in));
CL_DTYPE4 out = (exp(in) - exp(-in)) / (exp(in) + exp(-in));
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, (int2)(x, y), out);
}
__kernel void exp_act(__read_only image2d_t input,
__write_only image2d_t output,
__private const float threshold,
__private const float scale) {
const int x = get_global_id(0); // image_width
const int y = get_global_id(1); // image_height
__private const float scale) {
const int x = get_global_id(0); // image_width
const int y = get_global_id(1); // image_height
const sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE |
CLK_ADDRESS_CLAMP |
CLK_FILTER_NEAREST;
const sampler_t sampler =
CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
CL_DTYPE4 in = READ_IMG_TYPE(CL_DTYPE_CHAR, input, sampler, (int2)(x, y));
CL_DTYPE4 out = exp(in);
......@@ -135,19 +124,16 @@ __kernel void exp_act(__read_only image2d_t input,
}
__kernel void swish(__read_only image2d_t input,
__write_only image2d_t output,
__private const float threshold,
__private const float scale) {
const int x = get_global_id(0); // image_width
const int y = get_global_id(1); // image_height
__write_only image2d_t output,
__private const float threshold,
__private const float scale) {
const int x = get_global_id(0); // image_width
const int y = get_global_id(1); // image_height
const sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE |
CLK_ADDRESS_CLAMP |
CLK_FILTER_NEAREST;
const sampler_t sampler =
CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
CL_DTYPE4 in = READ_IMG_TYPE(CL_DTYPE_CHAR, input, sampler, (int2)(x, y));
CL_DTYPE4 out = in / (1 + exp(-(CL_DTYPE)scale * in));
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, (int2)(x, y), out);
}
lite/backends/opencl/cl_kernel/image/conv2d_1x1_opt_kernel.cl
浏览文件 @ 44fc847d
#include <cl_common.h>
__kernel void conv2d_1x1_opt(__private const int global_size_dim0,
__private const int global_size_dim1,
__private const int global_size_dim2,
__read_only image2d_t input_image,
__read_only image2d_t filter,
__kernel void conv2d_1x1_opt(
__private const int global_size_dim0,
__private const int global_size_dim1,
__private const int global_size_dim2,
__read_only image2d_t input_image,
__read_only image2d_t filter,
#if defined(BIASE_CH) || defined(BIASE_ELE)
__read_only image2d_t bias,
#endif
#ifdef BATCH_NORM
__read_only image2d_t new_scale,
__read_only image2d_t new_biase,
__read_only image2d_t new_scale,
__read_only image2d_t new_biase,
#endif
__write_only image2d_t output_image,
__private const int stride,
__private const int offset,
__private const int input_c_block,
__private const int input_c_origin,
__private const int dilation,
__private const int input_width, /* of one block */
__private const int input_height, /* of one block */
__private const int output_width,
__private const int output_height,
__private const int old_w) {
__write_only image2d_t output_image,
__private const int stride,
__private const int offset,
__private const int input_c_block,
__private const int input_c_origin,
__private const int dilation,
__private const int input_width, /* of one block */
__private const int input_height, /* of one block */
__private const int output_width,
__private const int output_height,
__private const int old_w) {
const int out_c = get_global_id(0);
const int out_w = get_global_id(1);
const int out_nh = get_global_id(2);
if (out_c >= global_size_dim0 || out_w >= global_size_dim1 ||
out_nh >= global_size_dim2) {
return;
}
int out_w0 = out_w;
int out_w1 = out_w + global_size_dim1;
int out_w2 = out_w + global_size_dim1 * 2;
......@@ -73,10 +76,10 @@ __read_only image2d_t new_scale,
CL_DTYPE4 output3 = output0;
#else
CL_DTYPE4 output0 = 0.0f;
CL_DTYPE4 output1 = 0.0f;
CL_DTYPE4 output2 = 0.0f;
CL_DTYPE4 output3 = 0.0f;
CL_DTYPE4 output0 = (CL_DTYPE4)(0.0f, 0.0f, 0.0f, 0.0f);
CL_DTYPE4 output1 = (CL_DTYPE4)(0.0f, 0.0f, 0.0f, 0.0f);
CL_DTYPE4 output2 = (CL_DTYPE4)(0.0f, 0.0f, 0.0f, 0.0f);
CL_DTYPE4 output3 = (CL_DTYPE4)(0.0f, 0.0f, 0.0f, 0.0f);
#endif
int max_w_bound = input_c_block * input_width;
......@@ -85,6 +88,14 @@ __read_only image2d_t new_scale,
// ------------0---------------
int2 pos_in = (int2)(i * input_width + in_pos_in_one_block0.x,
in_pos_in_one_block0.y);
pos_in.x = select(
pos_in.x,
-1,
(pos_in.x < i * input_width + in_pos_in_one_block0.x ||
pos_in.x >= i * input_width + in_pos_in_one_block0.x + input_width));
pos_in.y =
select(pos_in.y, -1, (pos_in.y < 0 || pos_in.y >= global_size_dim2));
CL_DTYPE4 input0 =
READ_IMG_TYPE(CL_DTYPE_CHAR, input_image, sampler, pos_in);
......@@ -131,6 +142,14 @@ __read_only image2d_t new_scale,
// -------------1--------------
pos_in = (int2)(i * input_width + in_pos_in_one_block1.x,
in_pos_in_one_block1.y);
pos_in.x = select(
pos_in.x,
-1,
(pos_in.x < i * input_width + in_pos_in_one_block0.x ||
pos_in.x >= i * input_width + in_pos_in_one_block0.x + input_width));
pos_in.y =
select(pos_in.y, -1, (pos_in.y < 0 || pos_in.y >= global_size_dim2));
CL_DTYPE4 input1 =
READ_IMG_TYPE(CL_DTYPE_CHAR, input_image, sampler, pos_in);
......@@ -167,6 +186,14 @@ __read_only image2d_t new_scale,
// -------------2--------------
pos_in = (int2)(i * input_width + in_pos_in_one_block2.x,
in_pos_in_one_block2.y);
pos_in.x = select(
pos_in.x,
-1,
(pos_in.x < i * input_width + in_pos_in_one_block0.x ||
pos_in.x >= i * input_width + in_pos_in_one_block0.x + input_width));
pos_in.y =
select(pos_in.y, -1, (pos_in.y < 0 || pos_in.y >= global_size_dim2));
CL_DTYPE4 input2 =
READ_IMG_TYPE(CL_DTYPE_CHAR, input_image, sampler, pos_in);
......@@ -203,6 +230,14 @@ __read_only image2d_t new_scale,
// -------------3--------------
pos_in = (int2)(i * input_width + in_pos_in_one_block3.x,
in_pos_in_one_block3.y);
pos_in.x = select(
pos_in.x,
-1,
(pos_in.x < i * input_width + in_pos_in_one_block0.x ||
pos_in.x >= i * input_width + in_pos_in_one_block0.x + input_width));
pos_in.y =
select(pos_in.y, -1, (pos_in.y < 0 || pos_in.y >= global_size_dim2));
CL_DTYPE4 input3 =
READ_IMG_TYPE(CL_DTYPE_CHAR, input_image, sampler, pos_in);
......@@ -287,7 +322,7 @@ __kernel void conv2d_1x1_simple(
__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,
......@@ -304,7 +339,10 @@ __read_only image2d_t new_scale,
const int out_c = get_global_id(0);
const int out_w = get_global_id(1);
const int out_nh = get_global_id(2);
if (out_c >= global_size_dim0 || out_w >= global_size_dim1 ||
out_nh >= global_size_dim2) {
return;
}
int out_w0 = out_w;
int out_w1 = out_w + global_size_dim1;
int out_w2 = out_w + global_size_dim1 * 2;
......@@ -350,16 +388,25 @@ __read_only image2d_t new_scale,
CL_DTYPE4 output3 = output0;
#else
CL_DTYPE4 output0 = 0.0f;
CL_DTYPE4 output1 = 0.0f;
CL_DTYPE4 output2 = 0.0f;
CL_DTYPE4 output3 = 0.0f;
CL_DTYPE4 output0 = (CL_DTYPE4)(0.0f, 0.0f, 0.0f, 0.0f);
CL_DTYPE4 output1 = (CL_DTYPE4)(0.0f, 0.0f, 0.0f, 0.0f);
CL_DTYPE4 output2 = (CL_DTYPE4)(0.0f, 0.0f, 0.0f, 0.0f);
CL_DTYPE4 output3 = (CL_DTYPE4)(0.0f, 0.0f, 0.0f, 0.0f);
#endif
for (int i = 0; i < input_c; ++i) {
// ------------0---------------
int2 pos_in = (int2)(i * input_width + in_pos_in_one_block0.x,
in_pos_in_one_block0.y);
pos_in.x = select(
pos_in.x,
-1,
(pos_in.x < i * input_width + in_pos_in_one_block0.x ||
pos_in.x >= i * input_width + in_pos_in_one_block0.x + input_width));
pos_in.y =
select(pos_in.y, -1, (pos_in.y < 0 || pos_in.y >= global_size_dim2));
CL_DTYPE4 input0 =
READ_IMG_TYPE(CL_DTYPE_CHAR, input_image, sampler, pos_in);
......@@ -379,6 +426,15 @@ __read_only image2d_t new_scale,
pos_in = (int2)(i * input_width + in_pos_in_one_block1.x,
in_pos_in_one_block1.y);
pos_in.x = select(
pos_in.x,
-1,
(pos_in.x < i * input_width + in_pos_in_one_block0.x ||
pos_in.x >= i * input_width + in_pos_in_one_block0.x + input_width));
pos_in.y =
select(pos_in.y, -1, (pos_in.y < 0 || pos_in.y >= global_size_dim2));
CL_DTYPE4 input1 =
READ_IMG_TYPE(CL_DTYPE_CHAR, input_image, sampler, pos_in);
output1 = mad(input1.x, weight0, output1);
......@@ -388,6 +444,14 @@ __read_only image2d_t new_scale,
pos_in = (int2)(i * input_width + in_pos_in_one_block2.x,
in_pos_in_one_block2.y);
pos_in.x = select(
pos_in.x,
-1,
(pos_in.x < i * input_width + in_pos_in_one_block0.x ||
pos_in.x >= i * input_width + in_pos_in_one_block0.x + input_width));
pos_in.y =
select(pos_in.y, -1, (pos_in.y < 0 || pos_in.y >= global_size_dim2));
CL_DTYPE4 input2 =
READ_IMG_TYPE(CL_DTYPE_CHAR, input_image, sampler, pos_in);
output2 = mad(input2.x, weight0, output2);
......@@ -397,6 +461,16 @@ __read_only image2d_t new_scale,
pos_in = (int2)(i * input_width + in_pos_in_one_block3.x,
in_pos_in_one_block3.y);
pos_in.x = select(
pos_in.x,
-1,
(pos_in.x < i * input_width + in_pos_in_one_block0.x ||
pos_in.x >= i * input_width + in_pos_in_one_block0.x + input_width));
pos_in.y =
select(pos_in.y, -1, (pos_in.y < 0 || pos_in.y >= global_size_dim2));
CL_DTYPE4 input3 =
READ_IMG_TYPE(CL_DTYPE_CHAR, input_image, sampler, pos_in);
output3 = mad(input3.x, weight0, output3);
......@@ -428,6 +502,16 @@ __read_only image2d_t new_scale,
output2 = activation_type4(output2);
output3 = activation_type4(output3);
// const int debug_pos = 0;
// int2 pos_test = (int2)(debug_pos, debug_pos);
// if (input_height == 112 && input_width == 112 && output_width == 112 &&
// output_height == 112) {
// output0 = READ_IMG_TYPE(CL_DTYPE_CHAR, input_image, sampler, pos_test);
// output1 = output0;
// output2 = output1;
// output3 = output2;
// }
if (out_w0 < old_w) {
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output_image, output_pos0, output0);
}
......
lite/backends/opencl/cl_kernel/image/conv2d_3x3_opt_kernel.cl
浏览文件 @ 44fc847d
......@@ -13,7 +13,6 @@ See the License for the specific language governing permissions and
limitations under the License. */
#include <cl_common.h>
__kernel void conv2d_3x3_opt(__private const int item_ch,
__private const int item_w,
__private const int item_h,
......
lite/backends/opencl/cl_kernel/image/layout_kernel.cl
浏览文件 @ 44fc847d
......@@ -18,7 +18,7 @@ limitations under the License. */
////////////////////////////////////////////////////////
// buffer -> image2d
////////////////////////////////////////////////////////
__kernel void buffer_to_image2d(__global CL_DTYPE *in,
__kernel void buffer_to_image2d(__global CL_DTYPE* in,
__write_only image2d_t output_image,
__private const int out_H,
__private const int out_W,
......@@ -26,11 +26,14 @@ __kernel void buffer_to_image2d(__global CL_DTYPE *in,
__private const int Stride0,
__private const int Stride1,
__private const int Stride2) {
const int out_c = get_global_id(0);
const int out_w = get_global_id(1);
const int out_nh = get_global_id(2);
if (out_c >= out_C || out_w >= out_W || out_nh >= out_H) {
return;
}
const int out_n = out_nh / out_H;
const int out_h = out_nh % out_H;
......@@ -56,26 +59,41 @@ __kernel void buffer_to_image2d(__global CL_DTYPE *in,
if (out_C - 4 * out_c >= 2) {
output.y = CONVERT_TYPE_TO(in[input_pos1], CL_COMPUTE_DTYPE);
} else {
output.y = CONVERT_TYPE_TO(0.f, CL_COMPUTE_DTYPE);
}
if (out_C - 4 * out_c >= 3) {
output.z = CONVERT_TYPE_TO(in[input_pos2], CL_COMPUTE_DTYPE);
} else {
output.z = CONVERT_TYPE_TO(0.f, CL_COMPUTE_DTYPE);
}
if (out_C - 4 * out_c >= 4) {
output.w = CONVERT_TYPE_TO(in[input_pos3], CL_COMPUTE_DTYPE);
} else {
output.w = CONVERT_TYPE_TO(0.f, CL_COMPUTE_DTYPE);
}
#ifdef DEBUG
if (out_w > 2045) {
printf("out_w:%d, out_C - 4 * out_c:%d, input[pos0~pos3]:%.2f %.2f %.2f %.2f\n",
out_w,
out_C - 4 * out_c,
(float)(in[input_pos0]),
(float)(in[input_pos1]),
(float)(in[input_pos2]),
(float)(in[input_pos3]));
printf("buffer2image ===> %d,%d,%d, out(%d,%d): %.2f %.2f %.2f %.2f \n", out_c, out_w, out_nh,
output_pos.x, output_pos.y,
(float)(output.x), (float)(output.y), (float)(output.z), (float)(output.w));
printf(
"out_w:%d, out_C - 4 * out_c:%d, input[pos0~pos3]:%.2f %.2f %.2f "
"%.2f\n",
out_w,
out_C - 4 * out_c,
(float)(in[input_pos0]),
(float)(in[input_pos1]),
(float)(in[input_pos2]),
(float)(in[input_pos3]));
printf("buffer2image ===> %d,%d,%d, out(%d,%d): %.2f %.2f %.2f %.2f \n",
out_c,
out_w,
out_nh,
output_pos.x,
output_pos.y,
(float)(output.x),
(float)(output.y),
(float)(output.z),
(float)(output.w));
}
#endif
......@@ -101,34 +119,44 @@ __kernel void image2d_to_buffer(__read_only image2d_t input,
const int in_h = in_nh % in_height;
const sampler_t sampler =
CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
const int pos_x = mad24(in_c, in_width, in_w);
CL_COMPUTE_DTYPE4 in = READ_IMG_TYPE(CL_COMPUTE_DTYPE_CHAR, input, sampler, (int2)(pos_x, in_nh));
CL_COMPUTE_DTYPE4 in = READ_IMG_TYPE(
CL_COMPUTE_DTYPE_CHAR, input, sampler, (int2)(pos_x, in_nh));
#ifdef DEBUG
if (in_w > 2045) {
printf("image2buffer ===> %d,%d,%d, in(%d,%d): %.2f %.2f %.2f %.2f \n", in_c, in_w, in_nh,
pos_x, in_nh,
(float)(in.x), (float)(in.y), (float)(in.z), (float)(in.w));
printf("image2buffer ===> %d,%d,%d, in(%d,%d): %.2f %.2f %.2f %.2f \n",
in_c,
in_w,
in_nh,
pos_x,
in_nh,
(float)(in.x),
(float)(in.y),
(float)(in.z),
(float)(in.w));
}
#endif
const int index = in_n * size_batch + in_c * size_block + in_h * in_width + in_w;
const int index =
in_n * size_batch + in_c * size_block + in_h * in_width + in_w;
out[index] = CONVERT_TYPE_TO(in.x, CL_DTYPE);
if (C - 4 * in_c >= 2) {
out[index + size_ch] = CONVERT_TYPE_TO(in.y, CL_DTYPE);
}
if(C - 4 * in_c >= 3) {
if (C - 4 * in_c >= 3) {
out[index + size_ch * 2] = CONVERT_TYPE_TO(in.z, CL_DTYPE);
}
if(C - 4 * in_c >= 4) {
if (C - 4 * in_c >= 4) {
out[index + size_ch * 3] = CONVERT_TYPE_TO(in.w, CL_DTYPE);
}
}
#if 0 // NOTE(ysh329): keep, un-used from paddle-mobile
#if 0 // NOTE(ysh329): keep, un-used from paddle-mobile
////////////////////////////////////////////////////////
// buffer -> image2d_nw
////////////////////////////////////////////////////////
......@@ -182,8 +210,7 @@ __kernel void buffer_to_image2d_nw(__global CL_DTYPE* in,
}
#endif
#if 0 // NOTE(ysh329): keep, un-used from paddle-mobile
#if 0 // NOTE(ysh329): keep, un-used from paddle-mobile
// image2d -> buffer
__kernel void image2d_to_buffer_2d(__private const int in_height,
__private const int in_width,
......@@ -208,15 +235,14 @@ __kernel void image2d_to_buffer_2d(__private const int in_height,
////////////////////////////////////////////////////////
// buffer -> image2d (divide by 255 to normalize)
////////////////////////////////////////////////////////
__kernel void buffer_to_image2d_with_pre255(__global uchar *in,
__kernel void buffer_to_image2d_with_pre255(__global uchar* in,
__write_only image2d_t output_image,
__private const int out_H,
__private const int out_W,
__private const int out_C,
__private const int Stride0,
__private const int Stride1,
__private const int Stride2){
__private const int Stride2) {
const int out_c = get_global_id(0);
const int out_w = get_global_id(1);
const int out_nh = get_global_id(2);
......@@ -231,7 +257,6 @@ __kernel void buffer_to_image2d_with_pre255(__global uchar *in,
const int in_h = out_h;
const int in_w = out_w;
int input_pos0 = in_n * Stride2 + in_c0 * Stride1 + in_h * Stride0 + in_w;
int input_pos1 = in_n * Stride2 + in_c1 * Stride1 + in_h * Stride0 + in_w;
int input_pos2 = in_n * Stride2 + in_c2 * Stride1 + in_h * Stride0 + in_w;
......@@ -243,30 +268,29 @@ __kernel void buffer_to_image2d_with_pre255(__global uchar *in,
CL_COMPUTE_DTYPE4 output = (CL_COMPUTE_DTYPE4)0.0f;
output.x = CONVERT_TYPE_TO(in[input_pos0], CL_COMPUTE_DTYPE) / 255;
if(out_C - 4 * out_c>=2){
output.y = CONVERT_TYPE_TO(in[input_pos1], CL_COMPUTE_DTYPE) / 255;
if (out_C - 4 * out_c >= 2) {
output.y = CONVERT_TYPE_TO(in[input_pos1], CL_COMPUTE_DTYPE) / 255;
}
if(out_C - 4 * out_c>=3){
output.z = CONVERT_TYPE_TO(in[input_pos2], CL_COMPUTE_DTYPE) / 255;
if (out_C - 4 * out_c >= 3) {
output.z = CONVERT_TYPE_TO(in[input_pos2], CL_COMPUTE_DTYPE) / 255;
}
if(out_C - 4 * out_c>=4){
output.w = CONVERT_TYPE_TO(in[input_pos3], CL_COMPUTE_DTYPE) / 255;
if (out_C - 4 * out_c >= 4) {
output.w = CONVERT_TYPE_TO(in[input_pos3], CL_COMPUTE_DTYPE) / 255;
}
WRITE_IMG_TYPE(CL_COMPUTE_DTYPE_CHAR, output_image, output_pos, output);
}
////////////////////////////////////////////////////////
// image2d -> buffer (multiply by 255 to de-normalize)
////////////////////////////////////////////////////////
__kernel void image2d_to_buffer_with_post255(__read_only image2d_t input,
__private const int in_width,
__private const int in_height,
__global uchar* out,
__private const int size_ch,
__private const int size_block,
__private const int size_batch,
__private const int C) {
__private const int in_width,
__private const int in_height,
__global uchar* out,
__private const int size_ch,
__private const int size_block,
__private const int size_batch,
__private const int C) {
const int in_c = get_global_id(0);
const int in_w = get_global_id(1);
const int in_nh = get_global_id(2);
......@@ -277,22 +301,34 @@ __kernel void image2d_to_buffer_with_post255(__read_only image2d_t input,
CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
const int pos_x = mad24(in_c, in_width, in_w);
CL_COMPUTE_DTYPE4 in = READ_IMG_TYPE(CL_COMPUTE_DTYPE_CHAR, input, sampler, (int2)(pos_x, in_nh)) * 255;
CL_COMPUTE_DTYPE4 in =
READ_IMG_TYPE(
CL_COMPUTE_DTYPE_CHAR, input, sampler, (int2)(pos_x, in_nh)) *
255;
#ifdef DEBUG
printf("in_c:%d, in_w:%d, in_nh:%d ===> in(%d,%d): %.2f %.2f %.2f %.2f\n",
in_c, in_w, in_nh, pos_x, in_nh, in.x, in.y, in.z, in.w);
in_c,
in_w,
in_nh,
pos_x,
in_nh,
in.x,
in.y,
in.z,
in.w);
#endif
const int index = in_n * size_batch + in_c * size_block + in_h * in_width + in_w;
const int index =
in_n * size_batch + in_c * size_block + in_h * in_width + in_w;
out[index] = convert_uchar_sat(in.x);
if(C - 4 * in_c>=2){
if (C - 4 * in_c >= 2) {
out[index + size_ch] = convert_uchar_sat(in.y);
}
if(C - 4 * in_c>=3){
if (C - 4 * in_c >= 3) {
out[index + size_ch * 2] = convert_uchar_sat(in.z);
}
if(C - 4 * in_c>=4){
if (C - 4 * in_c >= 4) {
out[index + size_ch * 3] = convert_uchar_sat(in.w);
}
}
lite/backends/opencl/cl_runtime.cc
浏览文件 @ 44fc847d
......@@ -45,6 +45,9 @@ bool CLRuntime::Init() {
bool is_device_init = InitializeDevice();
is_init_success_ = is_platform_init && is_device_init;
initialized_ = true;
context_ = CreateContext();
command_queue_ = CreateCommandQueue(context());
return initialized_;
}
......@@ -55,7 +58,7 @@ cl::Platform& CLRuntime::platform() {
cl::Context& CLRuntime::context() {
if (context_ == nullptr) {
context_ = CreateContext();
LOG(FATAL) << "context_ create failed. ";
}
return *context_;
}
......@@ -67,7 +70,7 @@ cl::Device& CLRuntime::device() {
cl::CommandQueue& CLRuntime::command_queue() {
if (command_queue_ == nullptr) {
command_queue_ = CreateCommandQueue(context());
LOG(FATAL) << "command_queue_ create failed. ";
}
return *command_queue_;
}
......@@ -96,7 +99,7 @@ std::unique_ptr<cl::UserEvent> CLRuntime::CreateEvent(
bool CLRuntime::BuildProgram(cl::Program* program, const std::string& options) {
/* -I +CLRuntime::Global()->cl_path() + "/cl_kernel"*/
std::string build_option = options + " -cl-fast-relaxed-math ";
std::string build_option = options + " -cl-fast-relaxed-math -cl-mad-enable";
VLOG(4) << "OpenCL build_option: " << build_option;
status_ = program->build({*device_}, build_option.c_str());
CL_CHECK_ERROR(status_);
......
lite/backends/opencl/target_wrapper.cc
浏览文件 @ 44fc847d
......@@ -66,7 +66,8 @@ void *TargetWrapperCL::MallocImage<float>(const size_t cl_image2d_width,
cl_int status;
cl::Image2D *cl_image =
new cl::Image2D(CLRuntime::Global()->context(),
CL_MEM_READ_WRITE | (host_ptr ? CL_MEM_COPY_HOST_PTR : 0),
CL_MEM_READ_WRITE | (host_ptr ? CL_MEM_COPY_HOST_PTR
: CL_MEM_ALLOC_HOST_PTR),
img_format,
cl_image2d_width,
cl_image2d_height,
......@@ -89,7 +90,8 @@ void *TargetWrapperCL::MallocImage<uint16_t>(const size_t cl_image2d_width,
cl_int status;
cl::Image2D *cl_image =
new cl::Image2D(CLRuntime::Global()->context(),
CL_MEM_READ_WRITE | (host_ptr ? CL_MEM_COPY_HOST_PTR : 0),
CL_MEM_READ_WRITE | (host_ptr ? CL_MEM_USE_HOST_PTR
: CL_MEM_ALLOC_HOST_PTR),
img_format,
cl_image2d_width,
cl_image2d_height,
......@@ -112,7 +114,8 @@ void *TargetWrapperCL::MallocImage<int32_t>(const size_t cl_image2d_width,
cl_int status;
cl::Image2D *cl_image =
new cl::Image2D(CLRuntime::Global()->context(),
CL_MEM_READ_WRITE | (host_ptr ? CL_MEM_COPY_HOST_PTR : 0),
CL_MEM_READ_WRITE | (host_ptr ? CL_MEM_COPY_HOST_PTR
: CL_MEM_ALLOC_HOST_PTR),
img_format,
cl_image2d_width,
cl_image2d_height,
......@@ -192,7 +195,6 @@ void TargetWrapperCL::MemcpySync(void *dst,
size_t size,
IoDirection dir) {
cl_int status;
cl::Event event;
auto stream = CLRuntime::Global()->command_queue();
switch (dir) {
case IoDirection::DtoD:
......@@ -202,9 +204,9 @@ void TargetWrapperCL::MemcpySync(void *dst,
0,
size,
nullptr,
&event);
nullptr);
CL_CHECK_FATAL(status);
event.wait();
CLRuntime::Global()->command_queue().finish();
break;
case IoDirection::HtoD:
status = stream.enqueueWriteBuffer(*static_cast<cl::Buffer *>(dst),
......@@ -283,7 +285,6 @@ void TargetWrapperCL::ImgcpySync(void *dst,
cl::array<size_t, 3> origin = {0, 0, 0};
cl::array<size_t, 3> region = {cl_image2d_width, cl_image2d_height, 1};
cl_int status;
cl::Event event;
auto stream = CLRuntime::Global()->command_queue();
switch (dir) {
case IoDirection::DtoD:
......@@ -293,9 +294,9 @@ void TargetWrapperCL::ImgcpySync(void *dst,
origin,
region,
nullptr,
&event);
nullptr);
CL_CHECK_FATAL(status);
event.wait();
CLRuntime::Global()->command_queue().finish();
break;
case IoDirection::HtoD:
status = stream.enqueueWriteImage(*static_cast<cl::Image2D *>(dst),
......
lite/core/context.h
浏览文件 @ 44fc847d
......@@ -340,27 +340,17 @@ class Context<TargetType::kX86> {
template <>
class Context<TargetType::kOpenCL> {
std::shared_ptr<CLContext> cl_context_;
using WaitListType =
std::unordered_map<decltype(static_cast<const void*>(nullptr)),
std::shared_ptr<cl::Event>>;
std::shared_ptr<WaitListType> cl_wait_list_;
public:
CLContext* cl_context() { return cl_context_.get(); }
WaitListType* cl_wait_list() { return cl_wait_list_.get(); }
void InitOnce() {
// Init cl runtime.
CHECK(CLRuntime::Global()->IsInitSuccess()) << "OpenCL runtime init failed";
cl_context_ = std::make_shared<CLContext>();
cl_wait_list_ = std::make_shared<WaitListType>();
}
void CopySharedTo(OpenCLContext* ctx) {
ctx->cl_context_ = cl_context_;
ctx->cl_wait_list_ = cl_wait_list_;
}
void CopySharedTo(OpenCLContext* ctx) { ctx->cl_context_ = cl_context_; }
};
#endif
......
lite/kernels/opencl/activation_buffer_compute.cc
浏览文件 @ 44fc847d
......@@ -62,23 +62,21 @@ class ReluCompute
CL_CHECK_FATAL(status);
auto global_work_size = cl::NDRange{count};
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_buf, event_);
}
private:
std::string kernel_func_name_{"relu"};
std::string build_options_{"-DCL_DTYPE_float -DRELU"};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
};
class SigmoidCompute
......@@ -121,23 +119,21 @@ class SigmoidCompute
CL_CHECK_FATAL(status);
auto global_work_size = cl::NDRange{count};
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_buf, event_);
}
private:
std::string kernel_func_name_{"sigmoid"};
std::string build_options_{"-DCL_DTYPE_float -DSIGMOID"};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
};
} // namespace opencl
......
lite/kernels/opencl/activation_buffer_compute_test.cc
浏览文件 @ 44fc847d
......@@ -85,16 +85,9 @@ TEST(opencl_relu_buffer, compute) {
kernel->Launch();
auto *wait_list = context->As<OpenCLContext>().cl_wait_list();
auto *out_ptr = param.Out->data<float, cl::Buffer>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl tensor. ---";
auto &event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the target cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
// run compute ref and check
std::unique_ptr<float[]> out_ref(new float[x_dim.production()]);
......@@ -145,16 +138,9 @@ TEST(opencl_sigmoid_buffer, compute) {
kernel->Launch();
auto *wait_list = context->As<OpenCLContext>().cl_wait_list();
auto *out_ptr = param.Out->data<float, cl::Buffer>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl tensor. ---";
auto &event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the target cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
// run compute ref and check
std::unique_ptr<float[]> out_ref(new float[x_dim.production()]);
......
lite/kernels/opencl/activation_image_compute.cc
浏览文件 @ 44fc847d
......@@ -147,16 +147,15 @@ class ActivationComputeImageDefault
auto& context = ctx_->As<OpenCLContext>();
CHECK(context.cl_context() != nullptr);
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size_,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_img, event_);
}
private:
......@@ -175,7 +174,6 @@ class ActivationComputeImageDefault
static_cast<size_t>(1), static_cast<size_t>(1), static_cast<size_t>(1)};
std::string build_options_{"-DCL_DTYPE_half"};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
};
} // namespace opencl
} // namespace kernels
......
lite/kernels/opencl/activation_image_compute_test.cc
浏览文件 @ 44fc847d
......@@ -234,19 +234,9 @@ TEST(act_image2d_fp16, compute) {
img_to_buf_kernel->Launch();
// wait for opencl
auto *wait_list = context->As<OpenCLContext>().cl_wait_list();
auto *out_ptr = ImageToBufferParam.y->data<float, cl::Buffer>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl "
"tensor. ---";
auto &event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the target "
"cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
// compute ref cpu
act_compute_ref<float>(
......
lite/kernels/opencl/bilinear_interp_image_compute.cc
浏览文件 @ 44fc847d
......@@ -142,16 +142,14 @@ class BilinearInterpImageCompute
static_cast<cl::size_type>(default_work_size[1]),
static_cast<cl::size_type>(default_work_size[2])};
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_img, event_);
#ifndef LITE_SHUTDOWN_LOG
VLOG(4) << "global_work_size:[2D]:" << global_work_size[0] << " "
<< global_work_size[1] << " " << global_work_size[2];
......@@ -163,7 +161,6 @@ class BilinearInterpImageCompute
std::string kernel_func_name_{"bilinear_interp"};
std::string build_options_{"-DCL_DTYPE_half"};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
};
} // namespace opencl
......
lite/kernels/opencl/bilinear_interp_image_compute_test.cc
浏览文件 @ 44fc847d
......@@ -187,18 +187,7 @@ TEST(bilinear_interp_image2d, compute) {
// LOG(INFO) << "out_image:" << out_image;
kernel->Launch();
auto* wait_list = context->As<OpenCLContext>().cl_wait_list();
auto* out_ptr = param.Out->data<half_t, cl::Image2D>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl "
"tensor. ---";
auto& event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the "
"target cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
std::unique_ptr<float[]> out_ref(
new float[out_dim.production()]);
......
lite/kernels/opencl/box_coder_image_compute.cc
浏览文件 @ 44fc847d
......@@ -47,8 +47,10 @@ class BoxCoderComputeImage : public KernelLite<TARGET(kOpenCL),
}
CHECK(context.cl_context() != nullptr);
VLOG(1) << "kernel_func_name_:" << kernel_func_name_;
context.cl_context()->AddKernel(
kernel_func_name_, "image/box_coder_kernel.cl", build_options_);
context.cl_context()->AddKernel(kernel_func_name_,
"image/box_coder_kernel.cl",
build_options_,
time_stamp_);
}
void Run() override {
......@@ -81,7 +83,7 @@ class BoxCoderComputeImage : public KernelLite<TARGET(kOpenCL),
auto& context = ctx_->As<OpenCLContext>();
CHECK(context.cl_context() != nullptr);
STL::stringstream kernel_key;
kernel_key << kernel_func_name_ << build_options_;
kernel_key << kernel_func_name_ << build_options_ << time_stamp_;
auto kernel = context.cl_context()->GetKernel(kernel_key.str());
auto default_work_size =
......@@ -120,16 +122,14 @@ class BoxCoderComputeImage : public KernelLite<TARGET(kOpenCL),
cl::NDRange{static_cast<cl::size_type>(default_work_size[0]),
static_cast<cl::size_type>(default_work_size[2])};
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_buf, event_);
#ifndef LITE_SHUTDOWN_LOG
VLOG(4) << "global_work_size:[2D]:" << global_work_size[0] << " "
......@@ -142,7 +142,7 @@ class BoxCoderComputeImage : public KernelLite<TARGET(kOpenCL),
param_t* boxcoder_param_{nullptr};
std::string kernel_func_name_{};
std::string build_options_{" -DCL_DTYPE_half"};
std::shared_ptr<cl::Event> event_{nullptr};
std::string time_stamp_{GetTimeStamp()};
};
} // namespace opencl
......
lite/kernels/opencl/box_coder_image_compute_test.cc
浏览文件 @ 44fc847d
......@@ -216,18 +216,7 @@ TEST(box_coder_image2d, compute) {
out_image_shape[0], out_image_shape[1]);
kernel->Launch();
auto* wait_list = context->As<OpenCLContext>().cl_wait_list();
auto* out_ptr = param.proposals->data<half_t, cl::Image2D>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl "
"tensor. ---";
auto& event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the "
"target cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
lite::Tensor out_ref_tensor;
out_ref_tensor.Resize(out_dim);
......
lite/kernels/opencl/concat_buffer_compute.cc
浏览文件 @ 44fc847d
......@@ -123,16 +123,15 @@ class ConcatCompute : public KernelLite<TARGET(kOpenCL),
CL_CHECK_FATAL(status);
status = kernel.setArg(++arg_idx, total1);
CL_CHECK_FATAL(status);
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_buf, event_);
} else {
auto start = 0;
for (int i = 0; i < inputs.size(); i++) {
......@@ -157,16 +156,15 @@ class ConcatCompute : public KernelLite<TARGET(kOpenCL),
CL_CHECK_FATAL(status);
status = kernel.setArg(++arg_idx, total0);
CL_CHECK_FATAL(status);
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_buf, event_);
start += size;
}
}
......@@ -182,7 +180,6 @@ class ConcatCompute : public KernelLite<TARGET(kOpenCL),
std::string kernel_func_name_{};
std::string build_options_{"-DCL_DTYPE_float"};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
};
} // namespace opencl
......
lite/kernels/opencl/concat_buffer_compute_test.cc
浏览文件 @ 44fc847d
......@@ -142,16 +142,7 @@ TEST(opencl_concat_buffer, compute) {
kernel->SetContext(std::move(concat_context));
kernel->Launch();
auto *wait_list = context->As<OpenCLContext>().cl_wait_list();
auto *out_ptr = param.output->data<float, cl::Buffer>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl tensor. ---";
auto &event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the target cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
// run compute ref and check
auto *out_ref_data = out_ref.mutable_data<float>(TARGET(kARM));
......
lite/kernels/opencl/concat_image_compute.cc
浏览文件 @ 44fc847d
......@@ -187,16 +187,15 @@ class ConcatComputeImage : public KernelLite<TARGET(kOpenCL),
CL_CHECK_FATAL(status);
status = kernel.setArg(++arg_idx, width_);
CL_CHECK_FATAL(status);
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_buf, event_);
} else {
auto start = 0;
for (int i = 0; i < inputs.size(); i++) {
......@@ -231,16 +230,15 @@ class ConcatComputeImage : public KernelLite<TARGET(kOpenCL),
status = kernel.setArg(++arg_idx, width_);
CL_CHECK_FATAL(status);
CL_CHECK_FATAL(status);
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_buf, event_);
start += inputs[i]->dims()[axis_];
}
}
......@@ -256,7 +254,6 @@ class ConcatComputeImage : public KernelLite<TARGET(kOpenCL),
std::string kernel_func_name_{};
std::string build_options_{" -DCL_DTYPE_half"};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
};
} // namespace opencl
......
lite/kernels/opencl/concat_image_compute_test.cc
浏览文件 @ 44fc847d
......@@ -245,20 +245,7 @@ TEST(concat_image2d, compute) {
LOG(INFO) << "run kernel: img_to_buf_kernel";
img_to_buf_kernel->Launch();
// wait for opencl
auto *wait_list = context->As<OpenCLContext>().cl_wait_list();
auto *out_ptr = ImageToBufferParam.y->data<float, cl::Buffer>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl "
"tensor. ---";
auto &event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the target "
"cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
// compute ref cp_u
std::vector<const float *> ins_ptr;
......
lite/kernels/opencl/conv_buffer_compute.cc
浏览文件 @ 44fc847d
......@@ -205,7 +205,7 @@ void ConvCompute::GemmlikeConv2d() {
CL_CHECK_FATAL(status);
auto global_work_size = cl::NDRange{static_cast<size_t>(out_stride)};
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
img2col_kernel,
cl::NullRange,
......@@ -301,17 +301,14 @@ void ConvCompute::GemmBatched(cl::Kernel& kernel,
status = kernel.setArg(++arg_idx, batch_size);
CL_CHECK_FATAL(status);
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
local_work_size,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(output_d, event_);
}
void ConvCompute::Run() { (this->*impl_)(); }
......
lite/kernels/opencl/conv_buffer_compute.h
浏览文件 @ 44fc847d
......@@ -57,7 +57,6 @@ class ConvCompute
std::vector<std::string> kernel_func_paths_{};
std::vector<std::string> build_options_{};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
};
} // namespace opencl
......
lite/kernels/opencl/conv_buffer_compute_test.cc
浏览文件 @ 44fc847d
......@@ -304,25 +304,14 @@ TEST(conv2d, compute_conv2d_1x1) {
// run opencl kernel
kernel->Launch();
auto* wait_list = context->As<OpenCLContext>().cl_wait_list();
auto* out_ptr = param.output->data<float, cl::Buffer>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl "
"tensor. ---";
auto& event = *(it->second);
event.wait();
double start_nanos =
event.getProfilingInfo<CL_PROFILING_COMMAND_START>();
double stop_nanos =
event.getProfilingInfo<CL_PROFILING_COMMAND_END>();
double elapsed_micros = (stop_nanos - start_nanos) / 1000.0;
LOG(INFO) << "Kernel Run Cost Time: " << elapsed_micros
<< " us.";
} else {
LOG(FATAL) << "Could not find the sync event for the target "
"cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
// double start_nanos =
// event.getProfilingInfo<CL_PROFILING_COMMAND_START>();
// double stop_nanos =
// event.getProfilingInfo<CL_PROFILING_COMMAND_END>();
// double elapsed_micros = (stop_nanos - start_nanos) / 1000.0;
// LOG(INFO) << "Kernel Run Cost Time: " << elapsed_micros
// << " us.";
// run cpu ref
auto* out_ref_data = out_ref.mutable_data<float>(TARGET(kARM));
......@@ -536,25 +525,15 @@ TEST(conv2d, compute_conv2d_gemm) {
// run opencl kernel
kernel->Launch();
auto* wait_list = context->As<OpenCLContext>().cl_wait_list();
auto* out_ptr = param.output->data<float, cl::Buffer>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl "
"tensor. ---";
auto& event = *(it->second);
event.wait();
double start_nanos =
event.getProfilingInfo<CL_PROFILING_COMMAND_START>();
double stop_nanos =
event.getProfilingInfo<CL_PROFILING_COMMAND_END>();
double elapsed_micros = (stop_nanos - start_nanos) / 1000.0;
LOG(INFO) << "Kernel Run Cost Time: " << elapsed_micros
<< " us.";
} else {
LOG(FATAL) << "Could not find the sync event for the target "
"cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
// double start_nanos =
// event.getProfilingInfo<CL_PROFILING_COMMAND_START>();
// double stop_nanos =
// event.getProfilingInfo<CL_PROFILING_COMMAND_END>();
// double elapsed_micros = (stop_nanos - start_nanos) /
// 1000.0;
// LOG(INFO) << "Kernel Run Cost Time: " << elapsed_micros
// << " us.";
// run cpu ref
auto* out_ref_data = out_ref.mutable_data<float>(TARGET(kARM));
......
lite/kernels/opencl/conv_image_compute.h
浏览文件 @ 44fc847d
......@@ -58,9 +58,11 @@ class ConvImageCompute : public KernelLite<TARGET(kOpenCL),
std::vector<std::string> kernel_func_paths_{};
std::vector<std::string> build_options_{};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
std::unique_ptr<Tensor> filter_gpu_image_{nullptr};
std::unique_ptr<Tensor> bias_gpu_image_{nullptr};
std::unique_ptr<Tensor> tensor_hold_filter_image_{nullptr};
std::unique_ptr<Tensor> tensor_hold_bias_image_{nullptr};
cl::NDRange global_work_size_ = cl::NDRange{
static_cast<size_t>(1), static_cast<size_t>(1), static_cast<size_t>(1)};
int c_blk_ = 1;
......
lite/kernels/opencl/conv_image_compute_test.cc
浏览文件 @ 44fc847d
......@@ -395,19 +395,7 @@ TEST(conv2d, compute_image2d_1x1) {
auto* output_image2d = output.mutable_data<half_t, cl::Image2D>(
out_image_width, out_image_height);
auto* wait_list = context->As<OpenCLContext>().cl_wait_list();
auto* out_ptr = param.output->data<half_t, cl::Image2D>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
SHADOW_LOG << "--- Find the sync event for the target cl "
"tensor. ---";
auto& event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the target"
"cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
TargetWrapperCL::ImgcpySync(out_image_v.data(),
output.data<half_t, cl::Image2D>(),
......@@ -530,11 +518,11 @@ const int stride = 2;
const int iw = 3;
const int oc = 2;
#else // big scale with group
const int stride = 1;
const int group = 32 / 1;
const int batch_size = 2;
const int ic = 32 / 1;
const int ih = 112 / 1;
const int stride = 2;
const int group = 1;
const int batch_size = 1;
const int ic = 3 / 1;
const int ih = 224 / 1;
const int iw = 112 / 1;
const int oc = 32 / 1;
#endif
......@@ -652,10 +640,10 @@ const int stride = 2;
SHADOW_LOG << "gen input and filter ...";
for (int i = 0; i < input_v.size(); ++i) {
input_v[i] = i * 0.001; // gen(engine);
input_v[i] = gen(engine);
}
for (int i = 0; i < filter_v.size(); ++i) {
filter_v[i] = 1 * 0.001; // gen(engine);
filter_v[i] = gen(engine);
}
SHADOW_LOG << "after gen input and filter ...";
......@@ -763,20 +751,7 @@ const int stride = 2;
auto* output_image2d = output.mutable_data<half_t, cl::Image2D>(
out_image_width, out_image_height);
auto* wait_list = context->As<OpenCLContext>().cl_wait_list();
auto* out_ptr = param.output->data<half_t, cl::Image2D>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
SHADOW_LOG << "--- Find the sync event for the target cl "
"tensor. ---";
auto& event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the target "
"cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
TargetWrapperCL::ImgcpySync(out_image_v.data(),
output.data<half_t, cl::Image2D>(),
out_image_width,
......@@ -848,8 +823,13 @@ const int stride = 2;
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 = COMPUTE_ABS_DIFF(output_v[i], out_ref_data[i]);
// EXPECT_LT(relative_diff, FP16_MAX_DIFF);
// EXPECT_LT(abs_diff, FP16_ABS_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["
......@@ -1115,19 +1095,7 @@ TEST(conv2d, compute_image2d_5x5) {
auto* output_image2d = output.mutable_data<half_t, cl::Image2D>(
out_image_width, out_image_height);
auto* wait_list = context->As<OpenCLContext>().cl_wait_list();
auto* out_ptr = param.output->data<half_t, cl::Image2D>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
SHADOW_LOG << "--- Find the sync event for the target cl "
"tensor. ---";
auto& event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the target "
"cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
TargetWrapperCL::ImgcpySync(out_image_v.data(),
output.data<half_t, cl::Image2D>(),
......@@ -1468,19 +1436,7 @@ TEST(conv2d, compute_image2d_7x7) {
auto* output_image2d = output.mutable_data<half_t, cl::Image2D>(
out_image_width, out_image_height);
auto* wait_list = context->As<OpenCLContext>().cl_wait_list();
auto* out_ptr = param.output->data<half_t, cl::Image2D>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
SHADOW_LOG << "--- Find the sync event for the target cl "
"tensor. ---";
auto& event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the target "
"cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
TargetWrapperCL::ImgcpySync(out_image_v.data(),
output.data<half_t, cl::Image2D>(),
......
lite/kernels/opencl/depthwise_conv2d_buffer_compute.cc
浏览文件 @ 44fc847d
......@@ -108,23 +108,21 @@ class DepthwiseConv2dCompute
status = kernel.setArg(++arg_idx, *bias_buf);
CL_CHECK_FATAL(status);
auto global_work_size = cl::NDRange(static_cast<size_t>(numel));
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(output_buf, event_);
}
private:
std::string kernel_func_name_{"depthwise_conv2d"};
std::string build_options_{"-DCL_DTYPE_float"};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
};
} // namespace opencl
......
lite/kernels/opencl/depthwise_conv2d_buffer_compute_test.cc
浏览文件 @ 44fc847d
......@@ -137,16 +137,7 @@ TEST(depthwise_conv2d_buffer_fp32, compute) {
output.Resize({4, 32, 110, 110});
kernel->Launch();
auto* wait_list = context->As<OpenCLContext>().cl_wait_list();
auto* out_ptr = param.output->data<float, cl::Buffer>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl tensor. ---";
auto& event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the target cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
lite::Tensor output_ref;
output_ref.Resize({4, 32, 110, 110});
......
lite/kernels/opencl/depthwise_conv2d_image_compute_test.cc
浏览文件 @ 44fc847d
......@@ -312,19 +312,7 @@ TEST(depthwise_conv2d, compute_basic) {
auto* output_image2d = output.mutable_data<half_t, cl::Image2D>(
out_image_width, out_image_height);
auto* wait_list = context->As<OpenCLContext>().cl_wait_list();
auto* out_ptr = param.output->data<half_t, cl::Image2D>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl "
"tensor. ---";
auto& event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the target "
"cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
TargetWrapperCL::ImgcpySync(out_image_v.data(),
output.data<half_t, cl::Image2D>(),
......@@ -503,20 +491,7 @@ TEST(depthwise_conv2d, compute_image2d_3x3) {
kernel->Launch();
auto* wait_list = context->As<OpenCLContext>().cl_wait_list();
auto* out_ptr = param.output->data<half_t, cl::Image2D>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl tensor. ---";
LOG(INFO) << "--- Find the sync event for the target cl tensor. ---";
auto& event = *(it->second);
event.wait();
} else {
LOG(FATAL)
<< "Could not find the sync event for the target cl tensor.";
LOG(INFO)
<< "Could not find the sync event for the target cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
lite::Tensor out_ref;
out_ref.Resize(output_dim);
......
lite/kernels/opencl/dropout_image_compute.cc
浏览文件 @ 44fc847d
......@@ -89,23 +89,20 @@ class DropoutComputeImage2D : public KernelLite<TARGET(kOpenCL),
static_cast<cl::size_type>(default_work_size.data()[1]),
static_cast<cl::size_type>(default_work_size.data()[2])};
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_img, event_);
}
private:
std::string kernel_func_name_{"dropout"};
std::string build_options_{"-DCL_DTYPE_half"};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
};
} // namespace opencl
......
lite/kernels/opencl/dropout_image_compute_test.cc
浏览文件 @ 44fc847d
......@@ -86,16 +86,7 @@ TEST(dropout_image2d_fp16, compute) {
LOG(INFO) << "out_image:" << out_image;
kernel->Launch();
auto* wait_list = context->As<OpenCLContext>().cl_wait_list();
auto* out_ptr = param.output->data<half_t, cl::Image2D>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl tensor. ---";
auto& event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the target cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
std::unique_ptr<float[]> out_ref(new float[out_dim.production()]);
dropout(input_v.data(), in_dim, out_ref.get(), 0.6);
......
lite/kernels/opencl/elementwise_add_buffer_compute.cc
浏览文件 @ 44fc847d
......@@ -63,16 +63,10 @@ void ElementwiseAddCompute::Run() {
CL_CHECK_FATAL(status);
auto global_work_size = cl::NDRange{channels_, batch_};
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
kernel, cl::NullRange, global_work_size, cl::NullRange, nullptr, nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_buf, event_);
}
void ElementwiseAddCompute::UpdateParams() {
......
lite/kernels/opencl/elementwise_add_buffer_compute.h
浏览文件 @ 44fc847d
......@@ -48,7 +48,6 @@ class ElementwiseAddCompute
std::string kernel_func_name_{"elementwise_add"};
std::string build_options_{"-DCL_DTYPE_float"};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
};
} // namespace opencl
......
lite/kernels/opencl/elementwise_add_buffer_compute_test.cc
浏览文件 @ 44fc847d
......@@ -144,16 +144,7 @@ TEST(elementwise_add_buffer, compute) {
kernel->Launch();
auto *wait_list = context->As<OpenCLContext>().cl_wait_list();
auto *out_ptr = param.Out->data<float, cl::Buffer>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl tensor. ---";
auto &event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the target cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
std::unique_ptr<float[]> out_ref(new float[out_dim.production()]);
elementwise_compute_ref<float>(
......@@ -225,16 +216,7 @@ TEST(fusion_elementwise_add_activation_buffer, compute) {
kernel->Launch();
auto *wait_list = context->As<OpenCLContext>().cl_wait_list();
auto *out_ptr = param.Out->data<float, cl::Buffer>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl tensor. ---";
auto &event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the target cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
std::unique_ptr<float[]> out_ref(new float[out_dim.production()]);
elementwise_compute_ref<float>(
......
lite/kernels/opencl/elementwise_add_image_compute.cc
浏览文件 @ 44fc847d
......@@ -153,16 +153,15 @@ void ElementwiseAddImageCompute::Run() {
auto& context = ctx_->As<OpenCLContext>();
CHECK(context.cl_context() != nullptr);
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size_,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_img, event_);
}
} // namespace opencl
......
lite/kernels/opencl/elementwise_add_image_compute.h
浏览文件 @ 44fc847d
......@@ -63,7 +63,6 @@ class ElementwiseAddImageCompute
cl::Kernel kernel_;
cl::NDRange global_work_size_ = cl::NDRange{
static_cast<size_t>(1), static_cast<size_t>(1), static_cast<size_t>(1)};
std::shared_ptr<cl::Event> event_{nullptr};
};
} // namespace opencl
......
lite/kernels/opencl/elementwise_mul_compute.cc
浏览文件 @ 44fc847d
......@@ -50,8 +50,10 @@ void ElementwiseMulFloatImageCompute::PrepareForRun() {
VLOG(4) << "y_dims.size():" << y_dims.size();
auto& context = ctx_->As<OpenCLContext>();
context.cl_context()->AddKernel(
kernel_func_name_, "image/elementwise_mul_kernel.cl", build_options_);
context.cl_context()->AddKernel(kernel_func_name_,
"image/elementwise_mul_kernel.cl",
build_options_,
time_stamp_);
}
void ElementwiseMulFloatImageCompute::Run() {
......@@ -88,7 +90,7 @@ void ElementwiseMulFloatImageCompute::Run() {
<< out_img_shape[1];
STL::stringstream kernel_key;
kernel_key << kernel_func_name_ << build_options_;
kernel_key << kernel_func_name_ << build_options_ << time_stamp_;
auto kernel = context.cl_context()->GetKernel(kernel_key.str());
int arg_idx = 0;
......@@ -150,16 +152,16 @@ void ElementwiseMulFloatImageCompute::Run() {
auto global_work_size = cl::NDRange{static_cast<cl::size_type>(x_img_width),
static_cast<cl::size_type>(x_img_height)};
event_ = std::shared_ptr<cl::Event>(new cl::Event);
auto status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_img, event_);
std::string time_stamp_{GetTimeStamp()};
VLOG(4) << "global_work_size:[2D]:" << x_img_width << " " << x_img_height;
}
......
lite/kernels/opencl/elementwise_mul_image_compute.cc
浏览文件 @ 44fc847d
......@@ -185,16 +185,15 @@ class ElementwiseMulImageCompute
auto global_work_size =
cl::NDRange{static_cast<cl::size_type>(x_img_width),
static_cast<cl::size_type>(x_img_height)};
event_ = std::shared_ptr<cl::Event>(new cl::Event);
auto status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_img, event_);
#ifndef LITE_SHUTDOWN_LOG
VLOG(4) << "global_work_size:[2D]:" << x_img_width << " " << x_img_height;
#endif
......@@ -205,7 +204,6 @@ class ElementwiseMulImageCompute
std::string kernel_func_name_{"elementwise_mul"};
std::string build_options_{"-DCL_DTYPE_half"};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
};
} // namespace opencl
......
lite/kernels/opencl/elementwise_sub_image_compute.cc
浏览文件 @ 44fc847d
......@@ -138,16 +138,9 @@ void ElementwiseSubImageCompute::Run() {
VLOG(4) << "global_work_size:[2D]:" << x_img_width << " " << x_img_height;
#endif
event_ = std::shared_ptr<cl::Event>(new cl::Event);
auto status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
kernel, cl::NullRange, global_work_size, cl::NullRange, nullptr, nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_img, event_);
}
} // namespace opencl
......
lite/kernels/opencl/elementwise_sub_image_compute.h
浏览文件 @ 44fc847d
......@@ -46,7 +46,6 @@ class ElementwiseSubImageCompute
std::string kernel_func_name_{"elementwise_sub"};
std::string build_options_{"-DCL_DTYPE_half"};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
};
} // namespace opencl
......
lite/kernels/opencl/fc_buffer_compute.cc
浏览文件 @ 44fc847d
......@@ -123,16 +123,15 @@ class FcCompute
auto& context = ctx_->As<OpenCLContext>();
CHECK(context.cl_context() != nullptr);
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size_,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_buf, event_);
}
private:
......@@ -145,7 +144,6 @@ class FcCompute
DDim last_x_dims_;
cl::NDRange global_work_size_;
cl::Kernel kernel_;
std::shared_ptr<cl::Event> event_{nullptr};
};
} // namespace opencl
......
lite/kernels/opencl/fc_buffer_compute_test.cc
浏览文件 @ 44fc847d
......@@ -162,17 +162,8 @@ TEST(fc, compute) {
// run opencl kernel
kernel->Launch();
// kernel->Launch();
auto* wait_list = context->As<OpenCLContext>().cl_wait_list();
auto* out_ptr = param.output->data<float, cl::Buffer>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl tensor. ---";
auto& event = *(it->second);
event.wait();
auto command_queue = CLRuntime::Global()->command_queue();
command_queue.finish();
CLRuntime::Global()->command_queue().finish();
#if 0
double start_nanos =
event.getProfilingInfo<CL_PROFILING_COMMAND_START>();
......@@ -181,10 +172,6 @@ TEST(fc, compute) {
double elapsed_micros = (stop_nanos - start_nanos) / 1000.0;
LOG(INFO) << "Kernel Run Cost Time: " << elapsed_micros << " us.";
#endif
} else {
LOG(FATAL)
<< "Could not find the sync event for the target cl tensor.";
}
std::vector<float> out_data_from_gpu(out_dim.production());
TargetWrapperCL::MemcpySync(
......
lite/kernels/opencl/grid_sampler_image_compute.cc
浏览文件 @ 44fc847d
......@@ -130,16 +130,15 @@ class GridSamplerImageCompute : public KernelLite<TARGET(kOpenCL),
auto& context = ctx_->As<OpenCLContext>();
CHECK(context.cl_context() != nullptr);
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size_,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_img, event_);
}
protected:
......@@ -154,7 +153,6 @@ class GridSamplerImageCompute : public KernelLite<TARGET(kOpenCL),
static_cast<size_t>(1), static_cast<size_t>(1), static_cast<size_t>(1)};
std::string build_options_{"-DCL_DTYPE_half"};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
};
} // namespace opencl
......
lite/kernels/opencl/grid_sampler_image_compute_test.cc
浏览文件 @ 44fc847d
......@@ -191,17 +191,7 @@ TEST(grid_samler_image2d, compute) {
// LOG(INFO) << "out_image:" << out_image;
kernel->Launch();
auto* wait_list = context->As<OpenCLContext>().cl_wait_list();
auto* out_ptr = param.out->data<half_t, cl::Image2D>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl tensor. ---";
auto& event = *(it->second);
event.wait();
} else {
LOG(FATAL)
<< "Could not find the sync event for the target cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
std::unique_ptr<float[]> out_ref(new float[out_dim.production()]);
gird_sampler_ref(
......
lite/kernels/opencl/instance_norm_image_compute.cc
浏览文件 @ 44fc847d
......@@ -137,16 +137,14 @@ class InstanceNormImageCompute : public KernelLite<TARGET(kOpenCL),
status = kernel.setArg(7, *out_img);
CL_CHECK_FATAL(status);
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
local_work_size,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_img, event_);
}
#else // paddle version
......@@ -260,16 +258,14 @@ class InstanceNormImageCompute : public KernelLite<TARGET(kOpenCL),
status = kernel.setArg(arg_idx++, in_w);
CL_CHECK_FATAL(status);
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
local_work_size,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_img, event_);
}
#endif
......@@ -278,7 +274,7 @@ class InstanceNormImageCompute : public KernelLite<TARGET(kOpenCL),
std::string kernel_func_name_{"instance_norm_onnx"};
std::string build_options_{"-DCL_DTYPE_half"};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
Tensor scale_image_;
Tensor bias_image_;
};
......
lite/kernels/opencl/io_copy_buffer_compute.cc
浏览文件 @ 44fc847d
......@@ -105,20 +105,11 @@ class IoCopykOpenCLToHostCompute
}
auto& context = ctx_->As<OpenCLContext>();
auto* wait_list = context.cl_wait_list();
auto it = wait_list->find(x_ptr);
if (it != wait_list->end()) {
#ifndef LITE_SHUTDOWN_LOG
VLOG(2) << "--- Find the sync event for the target cl tensor. ---";
VLOG(2) << "--- Find the sync event for the target cl tensor. ---";
#endif
auto& event = *(it->second);
event.wait();
auto command_queue = CLRuntime::Global()->command_queue();
command_queue.finish();
} else {
LOG(FATAL) << "Could not find the sync event for the target cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
CopyToHostSync(data, param.x->raw_data(), mem_size);
}
......
lite/kernels/opencl/io_copy_buffer_compute_test.cc
浏览文件 @ 44fc847d
......@@ -66,7 +66,6 @@ TEST(io_copy, compute) {
h2d_kernel->Launch();
auto* event_key = d_y.data<float, cl::Buffer>();
std::shared_ptr<cl::Event> event(new cl::Event);
context->As<OpenCLContext>().cl_wait_list()->emplace(event_key, event);
d2h_kernel->Launch();
auto* h_y_data = h_y.data<float>();
......
lite/kernels/opencl/layout_image_compute.cc
浏览文件 @ 44fc847d
......@@ -44,8 +44,10 @@ class LayoutComputeBufferChwToImageDefault
}
VLOG(1) << "kernel_func_name_:" << kernel_func_name_;
auto& context = ctx_->As<OpenCLContext>();
context.cl_context()->AddKernel(
kernel_func_name_, "image/layout_kernel.cl", build_options_);
context.cl_context()->AddKernel(kernel_func_name_,
"image/layout_kernel.cl",
build_options_,
time_stamp_);
}
void Run() override {
......@@ -95,7 +97,7 @@ class LayoutComputeBufferChwToImageDefault
auto& context = ctx_->As<OpenCLContext>();
CHECK(context.cl_context() != nullptr);
STL::stringstream kernel_key;
kernel_key << kernel_func_name_ << build_options_;
kernel_key << kernel_func_name_ << build_options_ << time_stamp_;
auto kernel = context.cl_context()->GetKernel(kernel_key.str());
int arg_idx = 0;
......@@ -122,16 +124,15 @@ class LayoutComputeBufferChwToImageDefault
cl::NDRange{static_cast<cl::size_type>((new_dims[1] + 3) / 4),
static_cast<cl::size_type>(new_dims[3]),
static_cast<cl::size_type>(new_dims[0] * new_dims[2])};
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(y_data, event_);
}
std::string doc() const override {
......@@ -140,9 +141,9 @@ class LayoutComputeBufferChwToImageDefault
}
private:
std::string time_stamp_{GetTimeStamp()};
std::string kernel_func_name_{"buffer_to_image2d"};
std::string build_options_{"-DCL_DTYPE_float"};
std::shared_ptr<cl::Event> event_{nullptr};
};
// [ImageDefault] -> [NCHW]
......@@ -158,8 +159,10 @@ class LayoutComputeImageDefaultToBufferChw
}
VLOG(1) << "kernel_func_name_:" << kernel_func_name_;
auto& context = ctx_->As<OpenCLContext>();
context.cl_context()->AddKernel(
kernel_func_name_, "image/layout_kernel.cl", build_options_);
context.cl_context()->AddKernel(kernel_func_name_,
"image/layout_kernel.cl",
build_options_,
time_stamp_);
}
void Run() override {
......@@ -202,7 +205,7 @@ class LayoutComputeImageDefaultToBufferChw
auto& context = ctx_->As<OpenCLContext>();
CHECK(context.cl_context() != nullptr);
STL::stringstream kernel_key;
kernel_key << kernel_func_name_ << build_options_;
kernel_key << kernel_func_name_ << build_options_ << time_stamp_;
auto kernel = context.cl_context()->GetKernel(kernel_key.str());
int arg_idx = 0;
......@@ -230,16 +233,15 @@ class LayoutComputeImageDefaultToBufferChw
cl::NDRange{static_cast<cl::size_type>((new_dims[1] + 3) / 4),
static_cast<cl::size_type>(new_dims[3]),
static_cast<cl::size_type>(new_dims[0] * new_dims[2])};
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(y_data, event_);
}
std::string doc() const override {
......@@ -248,9 +250,9 @@ class LayoutComputeImageDefaultToBufferChw
}
private:
std::string time_stamp_{GetTimeStamp()};
std::string kernel_func_name_{"image2d_to_buffer"};
std::string build_options_{"-DCL_DTYPE_float"};
std::shared_ptr<cl::Event> event_{nullptr};
};
// [NCHW] -> [ImageDW]
......@@ -263,8 +265,10 @@ class LayoutComputeBufferChwToImage2DNw
void PrepareForRun() override {
auto& context = ctx_->As<OpenCLContext>();
context.cl_context()->AddKernel(
kernel_func_name_, "buffer/layout_kernel.cl", build_options_);
context.cl_context()->AddKernel(kernel_func_name_,
"buffer/layout_kernel.cl",
build_options_,
time_stamp_);
}
void Run() override {
......@@ -298,7 +302,7 @@ class LayoutComputeBufferChwToImage2DNw
auto& context = ctx_->As<OpenCLContext>();
CHECK(context.cl_context() != nullptr);
STL::stringstream kernel_key;
kernel_key << kernel_func_name_ << build_options_;
kernel_key << kernel_func_name_ << build_options_ << time_stamp_;
auto kernel = context.cl_context()->GetKernel(kernel_key.str());
int arg_idx = 0;
......@@ -325,16 +329,15 @@ class LayoutComputeBufferChwToImage2DNw
cl::NDRange{static_cast<cl::size_type>((out_N + 3) / 4), // N blocks
static_cast<cl::size_type>(out_W), // w
static_cast<cl::size_type>(out_C * out_H)}; // ch
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(y_data, event_);
}
std::string doc() const override {
......@@ -342,9 +345,10 @@ class LayoutComputeBufferChwToImage2DNw
}
private:
std::string time_stamp_{GetTimeStamp()};
std::string kernel_func_name_{"buffer_to_image2d_nw"};
std::string build_options_{"-DCL_DTYPE_float "};
std::shared_ptr<cl::Event> event_{nullptr};
};
} // namespace opencl
......
lite/kernels/opencl/layout_image_compute_test.cc
浏览文件 @ 44fc847d
......@@ -246,20 +246,7 @@ TEST(layout_ImageDefault_With_Pre_Post, compute) {
LOG(INFO) << "run kernel: image2d_to_buffer_with_post255";
img_to_buf_kernel->Launch();
// wait for opencl
auto* wait_list = context->As<OpenCLContext>().cl_wait_list();
auto* out_ptr = ImageToBufferParam.y->data<float, cl::Buffer>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl "
"tensor. ---";
auto& event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the target "
"cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
// result
#ifdef PRINT_RESULT
......
lite/kernels/opencl/lrn_image_compute.cc
浏览文件 @ 44fc847d
......@@ -128,16 +128,14 @@ class LrnImageCompute : public KernelLite<TARGET(kOpenCL),
static_cast<cl::size_type>(default_work_size[1]),
static_cast<cl::size_type>(default_work_size[2])};
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_img, event_);
#ifndef LITE_SHUTDOWN_LOG
VLOG(4) << "global_work_size:[2D]:" << global_work_size[0] << " "
<< global_work_size[1] << " " << global_work_size[2];
......@@ -154,7 +152,6 @@ class LrnImageCompute : public KernelLite<TARGET(kOpenCL),
std::string kernel_func_name_{"lrn"};
std::string build_options_{"-DCL_DTYPE_half"};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
};
} // namespace opencl
......
lite/kernels/opencl/lrn_image_compute_test.cc
浏览文件 @ 44fc847d
......@@ -181,19 +181,7 @@ TEST(lrn_image2d, compute) {
// LOG(INFO) << "out_image:" << out_image;
kernel->Launch();
auto* wait_list =
context->As<OpenCLContext>().cl_wait_list();
auto* out_ptr = param.Out->data<half_t, cl::Image2D>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl "
"tensor. ---";
auto& event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the "
"target cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
std::unique_ptr<float[]> out_ref(
new float[out_dim.production()]);
......
lite/kernels/opencl/mul_buffer_compute.cc
浏览文件 @ 44fc847d
......@@ -91,16 +91,15 @@ class MulCompute
auto global_work_size = cl::NDRange{static_cast<size_t>((m_ + 3) / 4),
static_cast<size_t>((n_ + 3) / 4)};
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_buf, event_);
}
private:
......@@ -108,7 +107,6 @@ class MulCompute
std::string kernel_func_name_{"mat_mul"};
std::string build_options_{"-DCL_DTYPE_float"};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
};
} // namespace opencl
......
lite/kernels/opencl/mul_buffer_compute_test.cc
浏览文件 @ 44fc847d
......@@ -123,17 +123,7 @@ TEST(mul, compute) {
// run opencl kernel
kernel->Launch();
auto* wait_list = context->As<OpenCLContext>().cl_wait_list();
auto* out_ptr = param.output->data<float, cl::Buffer>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl tensor. ---";
auto& event = *(it->second);
event.wait();
} else {
LOG(FATAL)
<< "Could not find the sync event for the target cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
// run cpu ref
auto* out_ref_data = out_ref.mutable_data<float>(TARGET(kARM));
......
lite/kernels/opencl/nearest_interp_image_compute.cc
浏览文件 @ 44fc847d
......@@ -109,23 +109,21 @@ class NearestInterpComputeImageDefault
cl::NDRange{static_cast<cl::size_type>(default_work_size.data()[0]),
static_cast<cl::size_type>(default_work_size.data()[1]),
static_cast<cl::size_type>(default_work_size.data()[2])};
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_img, event_);
}
private:
std::string kernel_func_name_{"nearest_interp"};
std::string build_options_{" -DCL_DTYPE_half"};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
};
} // namespace opencl
......
lite/kernels/opencl/nearest_interp_image_compute_test.cc
浏览文件 @ 44fc847d
......@@ -208,20 +208,7 @@ TEST(nearest_interp_image2d, compute) {
LOG(INFO) << "run kernel: img_to_buf_kernel";
img_to_buf_kernel->Launch();
// wait for opencl
auto *wait_list = context->As<OpenCLContext>().cl_wait_list();
auto *out_ptr = ImageToBufferParam.y->data<float, cl::Buffer>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl "
"tensor. ---";
auto &event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the target "
"cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
// compute ref cpu
for (int nid = 0; nid < x_dim[0]; ++nid) {
......
lite/kernels/opencl/pad2d_image_compute.cc
浏览文件 @ 44fc847d
......@@ -142,16 +142,14 @@ class Pad2dCompute : public KernelLite<TARGET(kOpenCL),
static_cast<cl::size_type>(default_work_size[1]),
static_cast<cl::size_type>(default_work_size[2])};
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_img, event_);
#ifndef LITE_SHUTDOWN_LOG
VLOG(4) << "global_work_size:[2D]:" << global_work_size[0] << " "
<< global_work_size[1] << " " << global_work_size[2];
......@@ -163,7 +161,6 @@ class Pad2dCompute : public KernelLite<TARGET(kOpenCL),
std::string kernel_func_name_{};
std::string build_options_{"-DCL_DTYPE_half"};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
};
} // namespace opencl
......
lite/kernels/opencl/pad2d_image_compute_test.cc
浏览文件 @ 44fc847d
......@@ -262,22 +262,8 @@ TEST(pad2d_image2d, compute) {
img_to_buf_kernel->Launch();
// wait for opencl
auto *wait_list =
context->As<OpenCLContext>().cl_wait_list();
auto *out_ptr =
ImageToBufferParam.y->data<float, cl::Buffer>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl "
"tensor. ---";
auto &event = *(it->second);
event.wait();
} else {
LOG(FATAL)
<< "Could not find the sync event for the target "
"cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
// compute ref cpu
pad2d_ref(mapped_x,
......
lite/kernels/opencl/pool_buffer_compute.cc
浏览文件 @ 44fc847d
......@@ -105,23 +105,21 @@ class PoolCompute
status = kernel.setArg(++arg_idx, *output_buf);
CL_CHECK_FATAL(status);
auto global_work_size = cl::NDRange(static_cast<size_t>(numel));
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(output_buf, event_);
}
private:
std::string kernel_func_name_{"pool_"};
std::string build_options_{"-DCL_DTYPE_float"};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
};
} // namespace opencl
......
lite/kernels/opencl/pool_buffer_compute_test.cc
浏览文件 @ 44fc847d
......@@ -119,16 +119,7 @@ TEST(pool2d_buffer_fp32, compute) {
kernel->Launch();
auto* wait_list = context->As<OpenCLContext>().cl_wait_list();
auto* out_ptr = param.output->data<float, cl::Buffer>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl tensor. ---";
auto& event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the target cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
std::unique_ptr<float[]> out_ref(new float[out_dim.production()]);
pool_avg(0, 0, 1, 1, 7, 7, mapped_x, in_dim, out_ref.get(), out_dim);
......
lite/kernels/opencl/pool_image_compute.cc
浏览文件 @ 44fc847d
......@@ -150,23 +150,20 @@ class PoolComputeImage2D : public KernelLite<TARGET(kOpenCL),
status = kernel.setArg(++arg_idx, static_cast<const int>(paddings[0]));
CL_CHECK_FATAL(status);
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_img, event_);
}
private:
std::string kernel_func_name_{"pool_"};
std::string build_options_{"-DCL_DTYPE_half"};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
};
} // namespace opencl
......
lite/kernels/opencl/pool_image_compute_test.cc
浏览文件 @ 44fc847d
......@@ -137,16 +137,7 @@ TEST(pool2d_image2d, compute) {
LOG(INFO) << "out_image:" << out_image;
kernel->Launch();
auto* wait_list = context->As<OpenCLContext>().cl_wait_list();
auto* out_ptr = param.output->data<half_t, cl::Image2D>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl tensor. ---";
auto& event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the target cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
std::unique_ptr<float[]> out_ref(new float[out_dim.production()]);
pool_avg(0, 0, 1, 1, 7, 7, input_v.data(), in_dim, out_ref.get(), out_dim);
......
lite/kernels/opencl/reshape_image_compute.cc
浏览文件 @ 44fc847d
......@@ -154,23 +154,20 @@ class ReshapeComputeFloatImage : public KernelLite<TARGET(kOpenCL),
static_cast<size_t>(default_work_size.data()[1]),
static_cast<size_t>(default_work_size.data()[2])};
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_image, event_);
}
private:
std::string kernel_func_name_{"reshape"};
std::string build_options_{"-DCL_DTYPE_half"};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
};
} // namespace opencl
......
lite/kernels/opencl/reshape_image_compute_test.cc
浏览文件 @ 44fc847d
......@@ -184,17 +184,7 @@ TEST(reshape_opencl, compute) {
LOG(INFO) << "kernel launch ...";
kernel->Launch();
auto* wait_list = context->As<OpenCLContext>().cl_wait_list();
auto* out_ptr = param.output->data<half_t, cl::Image2D>();
auto it = wait_list->find(out_image);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl tensor. ---";
auto& event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the target cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
half_t* out_image_data = new half_t[out_image_shape.production() * 4];
TargetWrapperCL::ImgcpySync(out_image_data,
......
lite/kernels/opencl/scale_image_compute.cc
浏览文件 @ 44fc847d
......@@ -93,23 +93,20 @@ class ScaleComputeImage2D : public KernelLite<TARGET(kOpenCL),
status = kernel.setArg(3, bias);
CL_CHECK_FATAL(status);
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size_,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_img, event_);
}
private:
std::string kernel_func_name_{"scale"};
std::string build_options_{"-DCL_DTYPE_half"};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
param_t* scale_param_{nullptr};
cl::Kernel kernel_;
......
lite/kernels/opencl/scale_image_compute_test.cc
浏览文件 @ 44fc847d
......@@ -88,16 +88,7 @@ TEST(scale_image2d_fp32, compute) {
LOG(INFO) << "out_image:" << out_image;
kernel->Launch();
auto* wait_list = context->As<OpenCLContext>().cl_wait_list();
auto* out_ptr = param.output->data<half_t, cl::Image2D>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl tensor. ---";
auto& event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the target cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
std::unique_ptr<float[]> out_ref(new float[out_dim.production()]);
scale(input_v.data(), in_dim, out_ref.get(), 1.5f, 0.3f);
......
lite/kernels/opencl/slice_image_compute.cc
浏览文件 @ 44fc847d
......@@ -96,23 +96,20 @@ class SliceComputeImage2D : public KernelLite<TARGET(kOpenCL),
static_cast<cl::size_type>(default_work_size.data()[1]),
static_cast<cl::size_type>(default_work_size.data()[2])};
event_ = std::shared_ptr<cl::Event>(new cl::Event);
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
nullptr);
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_img, event_);
}
private:
std::string kernel_func_name_{"slice"};
std::string build_options_{"-DCL_DTYPE_half"};
std::string time_stamp_{GetTimeStamp()};
std::shared_ptr<cl::Event> event_{nullptr};
};
} // namespace opencl
......
lite/kernels/opencl/slice_image_compute_test.cc
浏览文件 @ 44fc847d
......@@ -98,16 +98,7 @@ TEST(slice_image2d_fp16, compute) {
LOG(INFO) << "out_image:" << out_image;
kernel->Launch();
auto* wait_list = context->As<OpenCLContext>().cl_wait_list();
auto* out_ptr = param.Out->data<half_t, cl::Image2D>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl tensor. ---";
auto& event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the target cl tensor.";
}
CLRuntime::Global()->command_queue().finish();
std::unique_ptr<float[]> out_ref(new float[out_dim.production()]);
slice_channel(input_v.data(), in_dim, out_ref.get(), 2, 5);
......
Markdown is supported
0% 或 .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑！
想要评论请 注册 或