I have already posted this question to the Khronos Forums as well as Stack Overflow to no avail. For a small program I wrote, the use of image2d_t memory objects as opposed to regular buffers would be beneficial (I think I could save on logic and compute on the ALU/FPUs). For computations I read pgm grayscale images to matrices stored in row-major order and copy those to buffers.
If I understand it correctly, for image2d_t memory objects I would have to blow up those matrices to matrices with four times as many entries, specifying for OpenCL what data is put in what channel. I tried doing something like this:
std::vector<float> A(256, 0.0);
for(int i=0; i<256; i++)
{
A[i] = 255.0f - i;
std::cout << A[i] << " ";
}
std::cout << std::endl;
// Blow up to float4 array
std::vector<float> A_img(1024, 0.0);
for(int i=0; i<256; i++)
{
A_img[4*i] = A[i];
A_img[(4*i)+3] = 1;
}
Maybe due to my device being little endian, the order of entries in the float4-vectors of the image2d_t would be reversed; I am not sure on that.
However, I can't seem to get this working. Attached you find a minimal working example of my efforts. The code is supposed to create a 16x16 grayscale image with a white pixel in the top left corner, a black pixel in the bottom right corner and any grayscale in between, "blow it up to a float4-vector", copy it to a device, read it, write it to another image and read back the copy.
If i run the code posted below on my Macbook, the "copied image" by the kernel contains only zeroes.
Any input on what I missed or am doing wrong would be greatly appreciated. Even a hint on literature or an open source project where I might find information on this would be really helpful, I couln't find anything.
#define CL_HPP_TARGET_OPENCL_VERSION 120
#define CL_HPP_MINIMUM_OPENCL_VERSION 120
#if defined(__APPLE__)
#include <OpenCL/cl2.hpp>
#else
#include <CL/cl2.hpp>
#endif
#include <iostream>
#include <string>
#include <vector>
#include <fstream>
int main(void)
{
// Set up platform, device and context
std::vector<cl::Platform> platforms;
std::vector<cl::Device> devices;
cl::Device default_device;
cl::Platform::get(&platforms);
if (platforms.size() == 0)
{
std::cout << "No OpenCL platform found, check installation!" << std::endl;
exit(-1);
}
platforms[0].getDevices(CL_DEVICE_TYPE_ALL, &devices);
if (devices.size() == 0)
{
std::cout << "No devices found in platform, check installation!" << std::endl;
exit(-1);
}
default_device = devices[0];
cl::Context context(default_device);
std::ifstream program_file("read_write_image.cl");
std::string program_string(std::istreambuf_iterator<char>(program_file), (std::istreambuf_iterator<char>()));
cl::Program::Sources source { program_string };
cl::Program dummy_program(context, source);
if (dummy_program.build()!=CL_SUCCESS)
{
std::cout << "Error building: " << dummy_program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(default_device)<< std::endl;
exit(-1);
}
cl::Kernel kernel(dummy_program, "read_write_image");
cl::CommandQueue queue(context, default_device);
// Set up dummy grayscale image
std::vector<float> A(256, 0.0);
for(int i=0; i<256; i++)
{
A[i] = 255.0f - i;
std::cout << A[i] << " ";
}
std::cout << std::endl;
// Blow up to float4 array
std::vector<float> A_img(1024, 0.0);
for(int i=0; i<256; i++)
{
A_img[4*i] = A[i];
A_img[(4*i)+3] = 1;
}
std::vector<float> B_img(1024, 0.0);
cl::ImageFormat grayscale(CL_R, CL_FLOAT);
cl::Image2D Input_Image(context, CL_MEM_READ_ONLY, grayscale, 16, 16);
cl::Image2D Output_Image(context, CL_MEM_WRITE_ONLY, grayscale, 16, 16);
std::array<cl::size_type, 3> origin {0,0,0};
std::array<cl::size_type, 3> region {16, 16, 1};
queue.enqueueWriteImage(Input_Image, CL_TRUE, origin, region, 0, 0, &A_img[0]);
kernel.setArg(0, Input_Image);
kernel.setArg(1, Output_Image);
queue.enqueueNDRangeKernel(kernel, cl::NullRange, cl::NDRange(16,16), cl::NullRange, NULL);
queue.enqueueReadImage(Output_Image, CL_TRUE, origin, region, 0, 0, &B_img[0]);
for(int i=0; i<1024; i++)
{
std::cout << B_img[i] << " ";
}
std::cout << std::endl;
return EXIT_SUCCESS;
}
Now the kernel code:
__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE|CLK_ADDRESS_CLAMP|CLK_FILTER_LINEAR;
__kernel void read_write_image(read_only image2d_t input_image, write_only image2d_t output_image)
{
int i = get_global_id(0);
int j = get_global_id(1);
float tmp = read_imagef(input_image, sampler, (int2) (i,j)).x;
write_imagef(output_image, (int2) (i,j), (float4) (tmp,0,0,1));
}
