loongarch 平台下，跑 PaddleLite-generic-demo.tar.gz ，显示feed无法找到 about paddle-lite

[csy0225]

valid_places.push_back(paddle::lite_api::Place{TARGET(kHost), PRECISION(kFloat)}); 试试？

[Sherry-yanqi]

valid_places.push_back(paddle::lite_api::Place{TARGET(kHost), PRECISION(kFloat)}); 试试？

感谢，成功奏效了，又产生了新的问题：
/usr/include/clc/image/image.h:19:24: note: passing argument to parameter 'image' here
input.cl:480:31: error: passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'
input.cl:38:28: note: expanded from macro 'WRITE_IMG_TYPE'
input.cl:36:24: note: expanded from macro '
_WRITE_IMG_TYPE'
/usr/include/clc/image/image.h:19:24: note: passing argument to parameter 'image' here
input.cl:483:31: error: passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'
input.cl:38:28: note: expanded from macro 'WRITE_IMG_TYPE'
input.cl:36:24: note: expanded from macro '
_WRITE_IMG_TYPE'
/usr/include/clc/image/image.h:19:24: note: passing argument to parameter 'image' here
input.cl:486:31: error: passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'
input.cl:38:28: note: expanded from macro 'WRITE_IMG_TYPE'
input.cl:36:24: note: expanded from macro '
_WRITE_IMG_TYPE'
/usr/include/clc/image/image.h:19:24: note: passing argument to parameter 'image' here
input.cl:550:31: error: passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'
input.cl:38:28: note: expanded from macro 'WRITE_IMG_TYPE'
input.cl:36:24: note: expanded from macro '
_WRITE_IMG_TYPE'

[csy0225]

valid_places.emplace_back(
Place{TARGET(kOpenCL), PRECISION(kFP16), DATALAYOUT(kImageDefault)});
valid_places.emplace_back(
Place{TARGET(kOpenCL), PRECISION(kFP16), DATALAYOUT(kImageFolder)});
valid_places.emplace_back(
Place{TARGET(kOpenCL), PRECISION(kFloat), DATALAYOUT(kNCHW)});
valid_places.emplace_back(
Place{TARGET(kOpenCL), PRECISION(kAny), DATALAYOUT(kImageDefault)});
valid_places.emplace_back(
Place{TARGET(kOpenCL), PRECISION(kAny), DATALAYOUT(kImageFolder)});
valid_places.emplace_back(
Place{TARGET(kOpenCL), PRECISION(kAny), DATALAYOUT(kNCHW)});
valid_places.emplace_back(
Place{TARGET(kOpenCL), PRECISION(kInt32), DATALAYOUT(kNCHW)});
valid_places.emplace_back(Place{TARGET(kARM)});

[csy0225]

你设置一下这些 Place，替换掉 demo 中现有的，跑下试试

[Sherry-yanqi]

替换掉后依然出现如下错误。WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, (int2)(pos_x, pos_y), out);
}E0325 09:45:10.467700 2910 cl_runtime.cc:351] OpenCL error with code CL_BUILD_PROGRAM_FAILURE happened in file /home/lixianrui/Paddle-Lite/lite/backends/opencl/cl_runtime.cc at line 351. Exiting.
F0325 09:45:10.467792 2910 cl_runtime.cc:357] Program build error: input.cl:1:9: warning: #pragma once in main file
input.cl:249:39: error: passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'
input.cl:38:28: note: expanded from macro 'WRITE_IMG_TYPE'
input.cl:36:24: note: expanded from macro '
_WRITE_IMG_TYPE'
/usr/include/clc/image/image.h:19:24: note: passing argument to parameter 'image' here
input.cl:389:39: error: passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'
input.cl:38:28: note: expanded from macro 'WRITE_IMG_TYPE'
input.cl:36:24: note: expanded from macro '
_WRITE_IMG_TYPE'
/usr/include/clc/image/image.h:19:24: note: passing argument to parameter 'image' here
input.cl:460:31: error: passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'
input.cl:38:28: note: expanded from macro 'WRITE_IMG_TYPE'
input.cl:36:24: note: expanded from macro '
_WRITE_IMG_TYPE'
/usr/include/clc/image/image.h:19:24: note: passing argument to parameter 'image' here
input.cl:478:31: error: passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'
input.cl:38:28: note: expanded from macro 'WRITE_IMG_TYPE'
input.cl:36:24: note: expanded from macro '
_WRITE_IMG_TYPE'
/usr/include/clc/image/image.h:19:24: note: passing argument to parameter 'image' here
input.cl:480:31: error: passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'
input.cl:38:28: note: expanded from macro 'WRITE_IMG_TYPE'
input.cl:36:24: note: expanded from macro '
_WRITE_IMG_TYPE'
/usr/include/clc/image/image.h:19:24: note: passing argument to parameter 'image' here
input.cl:483:31: error: passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'
input.cl:38:28: note: expanded from macro 'WRITE_IMG_TYPE'
input.cl:36:24: note: expanded from macro '
_WRITE_IMG_TYPE'
/usr/include/clc/image/image.h:19:24: note: passing argument to parameter 'image' here
input.cl:486:31: error: passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'
input.cl:38:28: note: expanded from macro 'WRITE_IMG_TYPE'
input.cl:36:24: note: expanded from macro '
_WRITE_IMG_TYPE'
/usr/include/clc/image/image.h:19:24: note: passing argument to parameter 'image' here
input.cl:550:31: error: passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'
input.cl:38:28: note: expanded from macro 'WRITE_IMG_TYPE'
input.cl:36:24: note: expanded from macro '
_WRITE_IMG_TYPE'
/usr/include/clc/image/image.h:19:24: note: passing argument to parameter 'image' here

[Sherry-yanqi]

替换掉后依然出现如下错误。WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, (int2)(pos_x, pos_y), out); }E0325 09:45:10.467700 2910 cl_runtime.cc:351] OpenCL error with code CL_BUILD_PROGRAM_FAILURE happened in file /home/lixianrui/Paddle-Lite/lite/backends/opencl/cl_runtime.cc at line 351. Exiting. F0325 09:45:10.467792 2910 cl_runtime.cc:357] Program build error: input.cl:1:9: warning: #pragma once in main file input.cl:249:39: error: passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t' input.cl:38:28: note: expanded from macro 'WRITE_IMG_TYPE' input.cl:36:24: note: expanded from macro ' _WRITE_IMG_TYPE' /usr/include/clc/image/image.h:19:24: note: passing argument to parameter 'image' here input.cl:389:39: error: passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t' input.cl:38:28: note: expanded from macro 'WRITE_IMG_TYPE' input.cl:36:24: note: expanded from macro ' _WRITE_IMG_TYPE' /usr/include/clc/image/image.h:19:24: note: passing argument to parameter 'image' here input.cl:460:31: error: passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t' input.cl:38:28: note: expanded from macro 'WRITE_IMG_TYPE' input.cl:36:24: note: expanded from macro ' _WRITE_IMG_TYPE' /usr/include/clc/image/image.h:19:24: note: passing argument to parameter 'image' here input.cl:478:31: error: passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t' input.cl:38:28: note: expanded from macro 'WRITE_IMG_TYPE' input.cl:36:24: note: expanded from macro ' _WRITE_IMG_TYPE' /usr/include/clc/image/image.h:19:24: note: passing argument to parameter 'image' here input.cl:480:31: error: passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t' input.cl:38:28: note: expanded from macro 'WRITE_IMG_TYPE' input.cl:36:24: note: expanded from macro ' _WRITE_IMG_TYPE' /usr/include/clc/image/image.h:19:24: note: passing argument to parameter 'image' here input.cl:483:31: error: passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t' input.cl:38:28: note: expanded from macro 'WRITE_IMG_TYPE' input.cl:36:24: note: expanded from macro ' _WRITE_IMG_TYPE' /usr/include/clc/image/image.h:19:24: note: passing argument to parameter 'image' here input.cl:486:31: error: passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t' input.cl:38:28: note: expanded from macro 'WRITE_IMG_TYPE' input.cl:36:24: note: expanded from macro ' _WRITE_IMG_TYPE' /usr/include/clc/image/image.h:19:24: note: passing argument to parameter 'image' here input.cl:550:31: error: passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t' input.cl:38:28: note: expanded from macro 'WRITE_IMG_TYPE' input.cl:36:24: note: expanded from macro ' _WRITE_IMG_TYPE' /usr/include/clc/image/image.h:19:24: note: passing argument to parameter 'image' here

#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);
#endif
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) {
out[index + size_ch] = convert_uchar_sat(in.y);
}
if (C - 4 * in_c >= 3) {
out[index + size_ch * 2] = convert_uchar_sat(in.z);
}
if (C - 4 * in_c >= 4) {
out[index + size_ch * 3] = convert_uchar_sat(in.w);
}
}
__kernel void image2d_default_to_image2d_folder(__read_only image2d_t input,
__write_only image2d_t output,
__private const int in_img_w,
__private const int in_img_h) {
const int pos_x = get_global_id(0);
const int pos_y = get_global_id(1);
CL_DTYPE4 in =
READ_IMG_TYPE(CL_DTYPE_CHAR, input, SAMPLER, (int2)(pos_x, pos_y));
CL_DTYPE4 in0 = 0.f;
CL_DTYPE4 in1 = 0.f;
CL_DTYPE4 in2 = 0.f;
CL_DTYPE4 in3 = 0.f;
in0 = READ_IMG_TYPE(CL_DTYPE_CHAR, input, SAMPLER, (int2)(pos_x * 4, pos_y));
if (pos_x * 4 + 1 < in_img_w) {
in1 = READ_IMG_TYPE(
CL_DTYPE_CHAR, input, SAMPLER, (int2)(pos_x * 4 + 1, pos_y));
}
if (pos_x * 4 + 2 < in_img_w) {
in2 = READ_IMG_TYPE(
CL_DTYPE_CHAR, input, SAMPLER, (int2)(pos_x * 4 + 2, pos_y));
}
if (pos_x * 4 + 3 < in_img_w) {
in3 = READ_IMG_TYPE(
CL_DTYPE_CHAR, input, SAMPLER, (int2)(pos_x * 4 + 3, pos_y));
}
CL_DTYPE4 out = (CL_DTYPE4)(in0.x, in1.x, in2.x, in3.x);
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, (int2)(pos_x, pos_y), out);
}
__kernel void image2d_folder_to_image2d_default(__read_only image2d_t input,
__write_only image2d_t output,
__private const int out_img_w,
__private const int out_img_h) {
const int pos_x = get_global_id(0);
const int pos_y = get_global_id(1);
CL_DTYPE4 in =
READ_IMG_TYPE(CL_DTYPE_CHAR, input, SAMPLER, (int2)(pos_x, pos_y));
CL_DTYPE4 out0 = 0.f;
CL_DTYPE4 out1 = 0.f;
CL_DTYPE4 out2 = 0.f;
CL_DTYPE4 out3 = 0.f;
out0.x = in.x;
out1.x = in.y;
out2.x = in.z;
out3.x = in.w;
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, (int2)(pos_x * 4, pos_y), out0);
if (pos_x * 4 + 1 < out_img_w) {
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, (int2)(pos_x * 4 + 1, pos_y), out1);
}
if (pos_x * 4 + 2 < out_img_w) {
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, (int2)(pos_x * 4 + 2, pos_y), out2);
}
if (pos_x * 4 + 3 < out_img_w) {
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, (int2)(pos_x * 4 + 3, pos_y), out3);
}
}
__kernel void image2d_folder_to_buffer(__read_only image2d_t input,
__global MUTABLE_TYPE* output,
__private const int out_h,
__private const int out_w) {
const int pos_x = get_global_id(0);
const int pos_y = get_global_id(1);
CL_DTYPE4 in =
READ_IMG_TYPE(CL_DTYPE_CHAR, input, SAMPLER, (int2)(pos_x, pos_y));
CL_DTYPE4 out = in;
int outpos_base = out_w * pos_y + pos_x * 4;
output[outpos_base] = CONVERT_TYPE_TO(out.x, MUTABLE_TYPE);
if (pos_x * 4 + 1 < out_w) {
output[outpos_base + 1] = CONVERT_TYPE_TO(out.y, MUTABLE_TYPE);
}
if (pos_x * 4 + 2 < out_w) {
output[outpos_base + 2] = CONVERT_TYPE_TO(out.z, MUTABLE_TYPE);
}
if (pos_x * 4 + 3 < out_w) {
output[outpos_base + 3] = CONVERT_TYPE_TO(out.w, MUTABLE_TYPE);
}
}
__kernel void image2d_folder_to_buffer_half2float(__read_only image2d_t input,
__global float* output,
__private const int out_h,
__private const int out_w) {
const int pos_x = get_global_id(0);
const int pos_y = get_global_id(1);
CL_DTYPE4 in =
READ_IMG_TYPE(CL_DTYPE_CHAR, input, SAMPLER, (int2)(pos_x, pos_y));
float4 out = convert_float4(in);
int outpos_base = out_w * pos_y + pos_x * 4;
output[outpos_base] = out.x;
if (pos_x * 4 + 1 < out_w) {
output[outpos_base + 1] = out.y;
}
if (pos_x * 4 + 2 < out_w) {
output[outpos_base + 2] = out.z;
}
if (pos_x * 4 + 3 < out_w) {
output[outpos_base + 3] = out.w;
}
}
__kernel void buffer_to_image2d_folder(__global const MUTABLE_TYPE* input,
__write_only image2d_t output,
__private const int out_h,
__private const int out_w,
__private const int length) {
const int pos_x = get_global_id(0);
const int pos_y = get_global_id(1);
int inpos_base = out_w * pos_y + pos_x * 4;
CL_COMPUTE_DTYPE4 out = (CL_COMPUTE_DTYPE4)(0.f, 0.f, 0.f, 0.f);
out.x = CONVERT_TYPE_TO(input[inpos_base], CL_COMPUTE_DTYPE);
if (inpos_base + 1 < length) {
out.y = CONVERT_TYPE_TO(input[inpos_base + 1], CL_COMPUTE_DTYPE);
}
if (inpos_base + 2 < length) {
out.z = CONVERT_TYPE_TO(input[inpos_base + 2], CL_COMPUTE_DTYPE);
}
if (inpos_base + 3 < length) {
out.w = CONVERT_TYPE_TO(input[inpos_base + 3], CL_COMPUTE_DTYPE);
}
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, (int2)(pos_x, pos_y), out);
}E0325 09:45:10.467700 2910 cl_runtime.cc:351] OpenCL error with code CL_BUILD_PROGRAM_FAILURE happened in file /home/lixianrui/Paddle-Lite/lite/backends/opencl/cl_runtime.cc at line 351. Exiting.
F0325 09:45:10.467792 2910 cl_runtime.cc:357] Program build error: input.cl:1:9: warning: #pragma once in main file
input.cl:249:39: error: passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'

[csy0225]

这个问题你能整理一下代码提个 PR 看下么

[Sherry-yanqi]

这个问题你能整理一下代码提个 PR 看下么

#10477
已提交PR

[sprouteer]

看起来是驱动问题导致的运行时编译失败，无法解析 __write_only image2d_t 和 __read_only image2d_t。目前的解决方案是采用opencl buffer版本的code，但是buffer支持的算子不多，不确定能否跑通这个模型。

[Sherry-yanqi]

看起来是驱动问题导致的运行时编译失败，无法解析 __write_only image2d_t 和 __read_only image2d_t。目前的解决方案是采用opencl buffer版本的code，但是buffer支持的算子不多，不确定能否跑通这个模型。

您好，具体指的是opencl的哪个版本及以上那

[Sherry-yanqi]

看起来是驱动问题导致的运行时编译失败，无法解析 __write_only image2d_t 和 __read_only image2d_t。目前的解决方案是采用opencl buffer版本的code，但是buffer支持的算子不多，不确定能否跑通这个模型。

更新了mesa版本，跑通了，感谢