diff --git a/src/operators/kernel/cl/cl_kernel/cl_common.h b/src/operators/kernel/cl/cl_kernel/cl_common.h index 95f0ab820883aeb0a467fa0612fe13de7319d2df..d718ea48aee5c38498f3fd1b8b3a7ea4b1b8b6dc 100644 --- a/src/operators/kernel/cl/cl_kernel/cl_common.h +++ b/src/operators/kernel/cl/cl_kernel/cl_common.h @@ -12,11 +12,9 @@ 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; - #pragma OPENCL EXTENSION cl_khr_fp16 : enable -inline hafl4 activation(half4 in +inline half4 activation(half4 in #ifdef PRELU , half4 prelu_alpha @@ -28,7 +26,7 @@ inline hafl4 activation(half4 in #endif #ifdef RELU - fmax(in, 0.0); + output = fmax(in, (half4)(0.0f)); #endif return output; } diff --git a/src/operators/kernel/cl/cl_kernel/conv_add_bn_relu_kernel.cl b/src/operators/kernel/cl/cl_kernel/conv_add_bn_relu_kernel.cl index 074280a23522efbf3220de8fde396d24e2165d30..75b14767a1332972605070f175f08187fac2ed65 100644 --- a/src/operators/kernel/cl/cl_kernel/conv_add_bn_relu_kernel.cl +++ b/src/operators/kernel/cl/cl_kernel/conv_add_bn_relu_kernel.cl @@ -12,10 +12,12 @@ 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 OPENCL EXTENSION cl_khr_fp16 : enable #define BIASE #define BATCH_NORM +#define RELU + +#include "cl_kernel/cl_common.h" __kernel void conv_3x3(__private const int global_size_dim0, __private const int global_size_dim1,