OpenCL OpFuse erro:Build program failed, err:-11 !

[mi-tao]

平台(如果交叉编译请再附上交叉编译目标平台):

Platform:

ubuntu opencl

Github版本:

Github Version:

2.62

Erro log

Build program failed, err:-11 ! 
kernelName.c_str()=s kernel_0 in buildKernelFromSource, 524 
CL ERROR CODE : -45, info:getKernel 
Program build log: BuildGroup_0:8:12: error: initializing '__private half4' (vector of 4 'half' values) with an expression of incompatible type 'float __attribute__((ext_vector_type(4)))' (vector of 4 'float' values)
    FLOAT4 var_0=read_imagef(input_0, SAMPLER, offset);
           ^     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
BuildGroup_0:13:5: error: no matching function for call to 'write_imagef'
    write_imagef(output_0, offset, var_2);
    ^~~~~~~~~~~~
BuildGroup_0:13:5: note: candidate function not viable: no known conversion from '__private half4' (vector of 4 'half' values) to 'float __attribute__((ext_vector_type(4)))' (vector of 4 'float' values) for 3rd argument
    write_imagef(output_0, offset, var_2);
    ^                              ~~~~~
BuildGroup_0:13:5: note: candidate function not viable: no known conversion from '__private __write_only image2d_t' to '__write_only image1d_array_t' for 1st argument
    write_imagef(output_0, offset, var_2);
    ^            ~~~~~~~~
BuildGroup_0:13:5: note: candidate function not viable: no known conversion from '__private __write_only image2d_t' to '__write_only image2d_depth_t' for 1st argument
    write_imagef(output_0, offset, var_2);
    ^            ~~~~~~~~
BuildGroup_0:13:5: note: candidate function not viable: no known conversion from '__private __write_only image2d_t' to '__read_write image2d_t' for 1st argument
    write_imagef(output_0, offset, var_2);
    ^            ~~~~~~~~
BuildGroup_0:13:5: note: candidate function not viable: no known conversion from '__private __write_only image2d_t' to '__read_write image1d_array_t' for 1st argument
    write_imagef(output_0, offset, var_2);
    ^            ~~~~~~~~
BuildGroup_0:13:5: note: candidate function not viable: no known conversion from '__private __write_only image2d_t' to '__read_write image2d_depth_t' for 1st argument
    write_imagef(output_0, offset, var_2);
    ^            ~~~~~~~~
BuildGroup_0:13:5: note: candidate function not viable: no known conversion from '__private __write_only image2d_t' to '__write_only image2d_array_t' for 1st argument
    write_imagef(output_0, offset, var_2);
    ^            ~~~~~~~~
BuildGroup_0:13:5: note: candidate function not viable: no known conversion from '__private __write_only image2d_t' to '__write_only image1d_t' for 1st argument
    write_imagef(output_0, offset, var_2);
    ^            ~~~~~~~~
BuildGroup_0:13:5: note: candidate function not viable: no known conversion from '__private __write_only image2d_t' to '__write_only image3d_t' for 1st argument
    write_imagef(output_0, offset, var_2);
    ^            ~~~~~~~~
BuildGroup_0:13:5: note: candidate function not viable: no known conversion from '__private __write_only image2d_t' to '__write_only image2d_array_depth_t' for 1st argument
    write_imagef(output_0, offset, var_2);
    ^            ~~~~~~~~
BuildGroup_0:13:5: note: candidate function not viable: no known conversion from '__private __write_only image2d_t' to '__write_only image1d_buffer_t' for 1st argument
    write_imagef(output_0, offset, var_2);
    ^            ~~~~~~~~
BuildGroup_0:13:5: note: candidate function not viable: no known conversion from '__private __write_only image2d_t' to '__read_write image2d_array_t' for 1st argument
    write_imagef(output_0, offset, var_2);
    ^            ~~~~~~~~
BuildGroup_0:13:5: note: candidate function not viable: no known conversion from '__private __write_only image2d_t' to '__read_write image1d_t' for 1st argument
    write_imagef(output_0, offset, var_2);
    ^            ~~~~~~~~
BuildGroup_0:13:5: note: candidate function not viable: no known conversion from '__private __write_only image2d_t' to '__read_write image1d_buffer_t' for 1st argument
    write_imagef(output_0, offset, var_2);
    ^            ~~~~~~~~
BuildGroup_0:13:5: note: candidate function not viable: no known conversion from '__private __write_only image2d_t' to '__read_write image3d_t' for 1st argument
    write_imagef(output_0, offset, var_2);
    ^            ~~~~~~~~
BuildGroup_0:13:5: note: candidate function not viable: no known conversion from '__private __write_only image2d_t' to '__read_write image2d_array_depth_t' for 1st argument
    write_imagef(output_0, offset, var_2);
    ^            ~~~~~~~~
BuildGroup_0:13:5: note: candidate function not viable: requires 4 arguments, but 3 were provided
BuildGroup_0:13:5: note: candidate function not viable: requires 4 arguments, but 3 were provided
BuildGroup_0:13:5: note: candidate function not viable: requires 4 arguments, but 3 were provided
BuildGroup_0:13:5: note: candidate function not viable: requires 4 arguments, but 3 were provided
BuildGroup_0:13:5: note: candidate function not viable: requires 4 arguments, but 3 were provided
BuildGroup_0:13:5: note: candidate function not viable: requires 4 arguments, but 3 were provided
BuildGroup_0:13:5: note: candidate function not viable: requires 4 arguments, but 3 were provided
BuildGroup_0:13:5: note: candidate function not viable: requires 4 arguments, but 3 were provided
BuildGroup_0:13:5: note: candidate function not viable: requires 4 arguments, but 3 were provided
BuildGroup_0:13:5: note: candidate function not viable: requires 4 arguments, but 3 were provided
BuildGroup_0:13:5: note: candidate function not viable: requires 4 arguments, but 3 were provided
BuildGroup_0:13:5: note: candidate function not viable: requires 4 arguments, but 3 were provided
BuildGroup_0:13:5: note: candidate function not viable: requires 4 arguments, but 3 were provided
BuildGroup_0:13:5: note: candidate function not viable: requires 4 arguments, but 3 were provided

[mi-tao]

将其中一个cl文件打印出来没有发现什么问题,为什么会在构建program报上述问题呢，cl如下

#define OFFSET_CHECK\
	const int c = get_global_id(0), w = get_global_id(1), hb = get_global_id(2);\
	if (c >= global_size_dim0 || w >= global_size_dim1 || hb >= global_size_dim2) { return; }\
	const int2 offset = (int2)(mad24(c, global_size_dim1, w), hb);
	__constant sampler_t SAMPLER = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
__kernel void kernel_0(__read_only image2d_t input_0, __write_only image2d_t output_0, __private const int global_size_dim0, __private const int global_size_dim1, __private const int global_size_dim2){
    OFFSET_CHECK;
    FLOAT4 var_0=read_imagef(input_0, SAMPLER, offset);
    FLOAT4 var_1;
    var_1=native_recip((float4)1+native_exp(convert_float4(-var_0)));
    FLOAT4 var_2;
    var_2=(var_0*var_1);
    write_imagef(output_0, offset, var_2);
}

[mi-tao]

当硬件配置设置成backend_config.precision = MNN::BackendConfig::Precision_Low;就会出现这个问题，原因在于配置为low的时候，数据类型为half，而在codegen部分codegen/opencl/OpenCLTarget.cpp文件中，写的数据类型为float，如 case UnaryOpOperation_SIGMOID: ss << inpName << "=native_recip((float4)1+native_exp(convert_float4(-" << operand << ")))"; break;，期待问题早日修复

[bitxsw93]

感谢反馈，修复中。

[jxt1234]

已经修正