Link Copied
Hi Vitor,
Thanks for providing the source code and the logs, we looked into it.
Could you please bring buffers out of the queue.submit scope and try running the code?
Also, I'm provided with the edited version of the code which we tested and executing without any errors, just in case.
// bufferlibclass.cpp
#define STB_IMAGE_IMPLEMENTATION
#define STB_IMAGE_WRITE_IMPLEMENTATION
#include"bufferlibclass.h"
//#include"../stb/stb_image.h"
//#include"../stb/stb_image_write.h"
using namespace std;
size_t cc::height()
{
return this->h;
}
size_t cc::width()
{
return this->w;
}
int cc::image()
// int cc::image(size_t m, size_t n, uint8_t* data, uint8_t* target)
{
float pi = 0.0;
// vector<uint8_t> arr(m*n);
cl::sycl::default_selector selector;
//cl::sycl::device dev = selector.select_device();
auto dev = selector.select_device();
// Device information
cout<<"Name: "<<dev.get_info<cl::sycl::info::device::name>()<<endl;
cout<<"Version: "<<dev.get_info<cl::sycl::info::device::version>()<<endl;
cout<<"Vendor: "<<dev.get_info<cl::sycl::info::device::vendor>()<<endl;
cout<<"Driver version: "<<dev.get_info<cl::sycl::info::device::driver_version>()<<endl;
// cout<<"Image support: "<<dev.get_info<cl::sycl::info::device::image_support>()<<endl;
// cout<<"Max mem alloc size: "<<dev.get_info<cl::sycl::info::device::max_mem_alloc_size>()<<endl;
// cout<<"Vendor ID: "<<dev.get_info<cl::sycl::info::device::vendor_id>()<<endl;
try
{
cl::sycl::queue queue(selector);
cl::sycl::buffer<uint8_t, 2> buf(data, cl::sycl::range<2>{w, h}); //EDITED
cl::sycl::buffer<uint8_t, 2> res(result, cl::sycl::range<2>{w, h}); //EDITED
queue.submit([&](cl::sycl::handler &hand)
{
// Try to create buffer from data
// cl::sycl::buffer<uint8_t, 2> buf(data, cl::sycl::range<2>{w, h});
// cl::sycl::buffer<uint8_t, 2> res(result, cl::sycl::range<2>{w, h});
auto edge = cl::sycl::range<2>{w-2, h-2};
// Os accessors têm que estar dentro do submit (e referenciando o handler)
auto bu = buf.get_access<cl::sycl::access::mode::read>(hand);
auto re = res.get_access<cl::sycl::access::mode::write>(hand);
// Usa essa implementação do parallel_for. Nela tem o parametro (segundo)
// que permite colocar um offSet. Exatamente o que precisa pra não usar as bordas
//
// template <typename KernelName,
// typename KernelType, int dimensions>
// void parallel_for(range<dimensions> numWorkItems,
// id<dimensions> workItemOffset, KernelType kernelFunc);
hand.parallel_for<class edge>
(cl::sycl::range<2>(w-2, h-2),
cl::sycl::id<2>(1, 1),
[=](cl::sycl::id<2> idx)
{
// idx[0] varia m
// idx[1] varia n
// if(bu[idx]>0) re[idx] = 1; //bu[idx];
// re[idx] = bu[idx];
//cout<<"("<<idx[0]<<")("<<idx[1]<<") ";
// Seguindo o Cel%{erity fazer um laplaciano
float sumf = (bu[idx[0]-1][idx[1]-1] +
bu[idx[0]][idx[1]-1] +
bu[idx[0]+1][idx[1]-1] +
bu[idx[0]-1][idx[1]] +
bu[idx[0]][idx[1]] +
bu[idx[0]+1][idx[1]] +
bu[idx[0]-1][idx[1]+1] +
bu[idx[0]][idx[1]+1] +
bu[idx[0]+1][idx[1]+1]
)/9;
int sum = static_cast<uint8_t>(sumf);
re[idx] = static_cast<uint8_t>(std::min(std::max(0, sum - (9 * bu[idx])), 255));
});
});
queue.wait();
} // end try
catch (std::exception& ex)
{
std::cerr<<"exception caught: "<< ex.what() << std::endl;
return 1;
}
return 0;
}
<output>
Please let us know if this resolves your issue.
Have a Good day!
Thanks & Regards
Goutham
