Простое ядро OpenCL выдает вертикальные полосы вместо сплошного цвета в выходном изображении. Зачем?
Я пытаюсь написать приложение для обработки изображений OpenCL, но моя проблема в том, что любая попытка изменить входное изображение приводит к появлению артефактов, которые выглядят как вертикальные полосы. Этого не происходит, если я копирую пиксели изображения, не изменяя их. Так, например, эта строка производит артефакты:
pixel = (uint4)(image1_pixel.x,
image1_pixel.y,
image1_pixel.z,
255);
…но этот работает как положено:
pixel = (uint4)(image1_pixel.x,
image1_pixel.y,
image1_pixel.z,
image1_pixel.w);
вход является непрозрачным 32-битным изображением PNG, поэтому я ожидаю, что обе строки кода приведут к одному и тому же результату. В действительности, однако, только вторая линия работает, как ожидалось. Первая строка дает вывод с артефактами.
Вот мое ядро:
__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
CLK_ADDRESS_CLAMP |
CLK_FILTER_NEAREST;
__kernel void test(__read_only image2d_t image1,
__write_only image2d_t out) {
const int2 pos = (int2)(get_global_id(0), get_global_id(1) );
uint4 image1_pixel = read_imageui(image1, sampler, pos);
uint4 pixel = (uint4)(image1_pixel.x,
image1_pixel.y,
image1_pixel.z,
255);
write_imageui(out, pos, pixel);
}
Вот соответствующая часть кода main.cpp:
CImg<unsigned char> image1("../input.png");
...
Image2D clImage1 = Image2D(context,
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
ImageFormat(CL_RGBA, CL_UNSIGNED_INT8),
image1.width(), image1.height(), 0, image1.data() );
Image2D clResult = Image2D(context, CL_MEM_WRITE_ONLY,
ImageFormat(CL_RGBA, CL_UNSIGNED_INT8),
image1.width(), image1.height(), 0, NULL);
Kernel test = Kernel(program, "test");
test.setArg(0, clImage1); test.setArg(1, clResult);
Event kernel_event, read_event;
queue.enqueueNDRangeKernel(test, NullRange,
NDRange(image1.width(), image1.height() ),
NullRange, NULL, &kernel_event);
cl::size_t<3> origin;
origin.push_back(0); origin.push_back(0); origin.push_back(0);
cl::size_t<3> region;
region.push_back(image1.width() );
region.push_back(image1.height() ); region.push_back(1);
queue.enqueueReadImage(clResult, CL_TRUE,
origin, region, 0, 0,
image1.data(), NULL, NULL);
kernel_event.wait();
image1.save("../output.png");
Вот можно загрузить полный исходный код для моего тестового приложения (он содержит короткий main.cpp до 30 строк, CMakeLists.txt, readme.txt, объясняющий, как его скомпилировать и запустить, входное изображение и ядро). Я использую библиотеку CImg для загрузки и сохранения изображений. Я дважды проверил, что вход открывается как 32-битное изображение RGBA. Я попытался запустить ядро с AMD или NVidia SDK и получил тот же результат.
Есть идеи, почему я получаю неожиданный результат?
Решение
CImg использует планарный формат, а OpenCL image2d_t ожидает чередования
(например, planar: R1R2R3R4R5R6 … G1G2G3G4G5G6 … B1B2B3B4B5B6 … и чередование: R1G1B1R2G2B2R3G3B3 …)
Когда вы копируете пиксели, не изменяя их, планарный формат не теряется, и он работает, но если вы начнете изменять один компонент изображения, вы на самом деле будете изменять все компоненты пикселя и, следовательно, темноту из-за этого плоского Чередованная проблема.
Если вы хотите придерживаться CImg, вы можете мультиплексировать изображение перед отправкой на устройство:
CImg<unsigned int> muxRGBA(const CImg<unsigned int>& pic)
{
assert( pic.spectrum() == 3 );
CImg<float> mux(pic.width()*4,pic.height(),pic.depth(),1);
cimg_forXYZ(pic,x,y,z){
for(int k=0; k<3; k++) {
mux(x*4+k,y,z,0) = pic(x,y,z,k);
}
mux(x*4+3,y,z,0) = 1.0f; // alpha channel
}
return mux;
}
а затем демультиплексируйте его после обработки, так что вы все равно можете отобразить / сохранить его с помощью CImg:
CImg<unsigned int> demuxRGBA(const CImg<unsigned int>& pic)
{
assert( pic.spectrum()==1 );
assert( pic.width()%4 == 0 );
CImg<float> demux(pic.width()/4,pic.height(),pic.depth(),3);
cimg_forXYZ(demux,x,y,z) {
for(int k=0; k<3; k++) {
demux(x,y,z,k) = pic(x*4+k,y,z,0);
}
}
return demux;
}
Другие решения
Ядро правильное, но CImg, кажется, делает что-то очень неправильное, поэтому в этом случае не стоит использовать его. Вместо этого я использовал Magick ++. Это не так просто, как CImg, но он более надежен и работает.
Вот как я читаю изображение из файла и конвертирую его в формат RGBA, чтобы использовать его с cl :: Image2D:
Magick::Image image1;
image1.read("input1/0.png");
long image1_size = 4 * image1.rows() * image1.columns();
uint8_t *image1_pixels = new uint8_t[image1_size];
image1.write(0, 0, image1.columns(), image1.rows(),
"RGBA", CharPixel, image1_pixels);
...
Image2D clImage1 = Image2D(context,
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
ImageFormat(CL_RGBA, CL_UNSIGNED_INT8),
image1.columns(), image1.rows(), 0, image1_pixels);
И вот как я записываю результат в файл:
queue.enqueueReadImage(clResult, CL_TRUE,
origin, region, 0, 0,
image1_pixels, NULL, NULL);
image1.read(image1.columns(), image1.rows(),
"RGBA", CharPixel, image1_pixels);
image1.write("../output.png");
Вот можно скачать полный исходный код.