Skip to content

Commit cc8475b

Browse files
committed
[OPENCL] add opencl kernel for sin cos tan atan asin acos ,test=develop (PaddlePaddle#4799)
1 parent 2123035 commit cc8475b

File tree

5 files changed

+364
-123
lines changed

5 files changed

+364
-123
lines changed

lite/backends/opencl/cl_kernel/image/trigonometric_kernel.cl

Lines changed: 80 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -30,3 +30,83 @@ __kernel void trigonometric_sin(__read_only image2d_t input,
3030
in.w = native_sin(in.w);
3131
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, (int2)(x, y), in);
3232
}
33+
34+
__kernel void trigonometric_cos(__read_only image2d_t input,
35+
__write_only image2d_t output) {
36+
const int x = get_global_id(0); // image_width
37+
const int y = get_global_id(1); // image_height
38+
39+
const sampler_t sampler =
40+
CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
41+
42+
CL_DTYPE4 in = READ_IMG_TYPE(CL_DTYPE_CHAR, input, sampler, (int2)(x, y));
43+
in.x = native_cos(in.x);
44+
in.y = native_cos(in.y);
45+
in.z = native_cos(in.z);
46+
in.w = native_cos(in.w);
47+
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, (int2)(x, y), in);
48+
}
49+
50+
__kernel void trigonometric_tan(__read_only image2d_t input,
51+
__write_only image2d_t output) {
52+
const int x = get_global_id(0); // image_width
53+
const int y = get_global_id(1); // image_height
54+
55+
const sampler_t sampler =
56+
CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
57+
58+
CL_DTYPE4 in = READ_IMG_TYPE(CL_DTYPE_CHAR, input, sampler, (int2)(x, y));
59+
in.x = native_tan(in.x);
60+
in.y = native_tan(in.y);
61+
in.z = native_tan(in.z);
62+
in.w = native_tan(in.w);
63+
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, (int2)(x, y), in);
64+
}
65+
66+
__kernel void trigonometric_atan(__read_only image2d_t input,
67+
__write_only image2d_t output) {
68+
const int x = get_global_id(0); // image_width
69+
const int y = get_global_id(1); // image_height
70+
71+
const sampler_t sampler =
72+
CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
73+
74+
CL_DTYPE4 in = READ_IMG_TYPE(CL_DTYPE_CHAR, input, sampler, (int2)(x, y));
75+
in.x = atan(in.x);
76+
in.y = atan(in.y);
77+
in.z = atan(in.z);
78+
in.w = atan(in.w);
79+
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, (int2)(x, y), in);
80+
}
81+
82+
__kernel void trigonometric_asin(__read_only image2d_t input,
83+
__write_only image2d_t output) {
84+
const int x = get_global_id(0); // image_width
85+
const int y = get_global_id(1); // image_height
86+
87+
const sampler_t sampler =
88+
CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
89+
90+
CL_DTYPE4 in = READ_IMG_TYPE(CL_DTYPE_CHAR, input, sampler, (int2)(x, y));
91+
in.x = asin(in.x);
92+
in.y = asin(in.y);
93+
in.z = asin(in.z);
94+
in.w = asin(in.w);
95+
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, (int2)(x, y), in);
96+
}
97+
98+
__kernel void trigonometric_acos(__read_only image2d_t input,
99+
__write_only image2d_t output) {
100+
const int x = get_global_id(0); // image_width
101+
const int y = get_global_id(1); // image_height
102+
103+
const sampler_t sampler =
104+
CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
105+
106+
CL_DTYPE4 in = READ_IMG_TYPE(CL_DTYPE_CHAR, input, sampler, (int2)(x, y));
107+
in.x = acos(in.x);
108+
in.y = acos(in.y);
109+
in.z = acos(in.z);
110+
in.w = acos(in.w);
111+
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, (int2)(x, y), in);
112+
}

lite/kernels/opencl/CMakeLists.txt

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -116,7 +116,7 @@ lite_cc_test(test_pad2d_image_opencl SRCS pad2d_image_compute_test.cc
116116
lite_cc_test(test_box_coder_image_opencl SRCS box_coder_image_compute_test.cc
117117
DEPS box_coder_opencl_image op_registry program context)
118118

119-
lite_cc_test(test_sin_image_opencl SRCS sin_image_compute_test.cc
119+
lite_cc_test(test_trigonometric_image_opencl SRCS trigonometric_image_compute_test.cc
120120
DEPS trigonometric_opencl_image op_registry program context)
121121

122122
######################

lite/kernels/opencl/sin_image_compute_test.cc

Lines changed: 0 additions & 122 deletions
This file was deleted.

lite/kernels/opencl/trigonometric_image_compute.cc

Lines changed: 100 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -136,6 +136,26 @@ class TrigonometricComputeImage2D
136136
class SinComputeImage2D : public TrigonometricComputeImage2D {
137137
std::string KernelFunctionName() override { return "trigonometric_sin"; }
138138
};
139+
140+
class CosComputeImage2D : public TrigonometricComputeImage2D {
141+
std::string KernelFunctionName() override { return "trigonometric_cos"; }
142+
};
143+
144+
class TanComputeImage2D : public TrigonometricComputeImage2D {
145+
std::string KernelFunctionName() override { return "trigonometric_tan"; }
146+
};
147+
148+
class AtanComputeImage2D : public TrigonometricComputeImage2D {
149+
std::string KernelFunctionName() override { return "trigonometric_atan"; }
150+
};
151+
152+
class AsinComputeImage2D : public TrigonometricComputeImage2D {
153+
std::string KernelFunctionName() override { return "trigonometric_asin"; }
154+
};
155+
156+
class AcosComputeImage2D : public TrigonometricComputeImage2D {
157+
std::string KernelFunctionName() override { return "trigonometric_acos"; }
158+
};
139159
} // namespace opencl
140160
} // namespace kernels
141161
} // namespace lite
@@ -156,3 +176,83 @@ REGISTER_LITE_KERNEL(sin,
156176
PRECISION(kFP16),
157177
DATALAYOUT(kImageDefault))})
158178
.Finalize();
179+
180+
REGISTER_LITE_KERNEL(cos,
181+
kOpenCL,
182+
kFP16,
183+
kImageDefault,
184+
paddle::lite::kernels::opencl::CosComputeImage2D,
185+
image2d)
186+
.BindInput("X",
187+
{LiteType::GetTensorTy(TARGET(kOpenCL),
188+
PRECISION(kFP16),
189+
DATALAYOUT(kImageDefault))})
190+
.BindOutput("Out",
191+
{LiteType::GetTensorTy(TARGET(kOpenCL),
192+
PRECISION(kFP16),
193+
DATALAYOUT(kImageDefault))})
194+
.Finalize();
195+
196+
REGISTER_LITE_KERNEL(tan,
197+
kOpenCL,
198+
kFP16,
199+
kImageDefault,
200+
paddle::lite::kernels::opencl::TanComputeImage2D,
201+
image2d)
202+
.BindInput("X",
203+
{LiteType::GetTensorTy(TARGET(kOpenCL),
204+
PRECISION(kFP16),
205+
DATALAYOUT(kImageDefault))})
206+
.BindOutput("Out",
207+
{LiteType::GetTensorTy(TARGET(kOpenCL),
208+
PRECISION(kFP16),
209+
DATALAYOUT(kImageDefault))})
210+
.Finalize();
211+
212+
REGISTER_LITE_KERNEL(atan,
213+
kOpenCL,
214+
kFP16,
215+
kImageDefault,
216+
paddle::lite::kernels::opencl::AtanComputeImage2D,
217+
image2d)
218+
.BindInput("X",
219+
{LiteType::GetTensorTy(TARGET(kOpenCL),
220+
PRECISION(kFP16),
221+
DATALAYOUT(kImageDefault))})
222+
.BindOutput("Out",
223+
{LiteType::GetTensorTy(TARGET(kOpenCL),
224+
PRECISION(kFP16),
225+
DATALAYOUT(kImageDefault))})
226+
.Finalize();
227+
228+
REGISTER_LITE_KERNEL(asin,
229+
kOpenCL,
230+
kFP16,
231+
kImageDefault,
232+
paddle::lite::kernels::opencl::AsinComputeImage2D,
233+
image2d)
234+
.BindInput("X",
235+
{LiteType::GetTensorTy(TARGET(kOpenCL),
236+
PRECISION(kFP16),
237+
DATALAYOUT(kImageDefault))})
238+
.BindOutput("Out",
239+
{LiteType::GetTensorTy(TARGET(kOpenCL),
240+
PRECISION(kFP16),
241+
DATALAYOUT(kImageDefault))})
242+
.Finalize();
243+
244+
REGISTER_LITE_KERNEL(acos,
245+
kOpenCL,
246+
kFP16,
247+
kImageDefault,
248+
paddle::lite::kernels::opencl::AcosComputeImage2D,
249+
image2d)
250+
.BindInput("X",
251+
{LiteType::GetTensorTy(TARGET(kOpenCL),
252+
PRECISION(kFP16),
253+
DATALAYOUT(kImageDefault))})
254+
.BindOutput("Out",
255+
{LiteType::GetTensorTy(TARGET(kOpenCL),
256+
PRECISION(kFP16),
257+
DATALAYOUT(kImageDefault))})
258+
.Finalize();

0 commit comments

Comments
 (0)