- Mark as New
- Bookmark
- Subscribe
- Mute
- Subscribe to RSS Feed
- Permalink
- Report Inappropriate Content
I got wrong results on my computer i7-6700 (HD530), win10 and vs2017. Here is my code,
#ifdef __APPLE__ #include <OpenCL/opencl.h> #else #include <CL/cl.h> #endif #include<memory> #include<iostream> using namespace std; const char* kernelSource = R"( __constant sampler_t samplerf = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR; __kernel void toy(__read_only image2d_t img, __write_only image2d_t rst) { int x = get_global_id(0); int y = get_global_id(1); float4 src = read_imagef(img, samplerf, (int2)(x,y)); write_imagef(rst, (int2)(x, y), src); } )"; int main() { cl_platform_id platform_id; cl_device_id device_id; cl_context context; cl_command_queue command_queue; cl_program program; cl_kernel kernel; clGetPlatformIDs(1, &platform_id, NULL); clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL); context = clCreateContext(NULL, 1, &device_id, NULL, NULL, NULL); command_queue = clCreateCommandQueueWithProperties(context, device_id, 0, NULL); program = clCreateProgramWithSource(context, 1, (const char**)&kernelSource, NULL, NULL); clBuildProgram(program, 1, &device_id, NULL, NULL, NULL); kernel = clCreateKernel(program, "toy", NULL); if (kernel == NULL) { cout << "failed" << endl; } float input[2 * 3 * 4], output[2 * 3 * 4]; for (int i = 0; i < 24; i++) input = i; cl_image_desc desc; memset((void*)&desc, 0, sizeof(cl_image_desc)); desc.image_width = 2; desc.image_height = 3; desc.image_type = CL_MEM_OBJECT_IMAGE2D; cl_image_format fmt; memset((void*)&fmt, 0, sizeof(cl_image_format)); fmt.image_channel_data_type = CL_FLOAT; fmt.image_channel_order = CL_RGBA; cl_mem src = clCreateImage(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, &fmt, &desc, input, NULL); cl_mem dst = clCreateImage(context, CL_MEM_WRITE_ONLY, &fmt, &desc, NULL, NULL); clSetKernelArg(kernel, 0, sizeof(cl_mem), &src); clSetKernelArg(kernel, 1, sizeof(cl_mem), &dst); const size_t orign[3] = { 0,0,0 }; const size_t region[3] = { 2,3,1 }; const size_t global_work[2] = { 2,3 }; clEnqueueNDRangeKernel(command_queue, kernel, 2, NULL, global_work, NULL, 0, NULL, NULL); clEnqueueReadImage(command_queue, dst, true, orign, region, 0, 0, output, 0, NULL, NULL); clReleaseMemObject(src); clReleaseMemObject(dst); for (int i = 0; i < 24; i++) { cout << output << " "; } cout << endl; return 0; }
I got results " 0 1 2 3 2 3 4 5 4 5 6 7 6 7 8 9 12 13 14 15 14 15 16 17 ", not expected "0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23". I am not sure the reason.
Thanks very much.
- Mark as New
- Bookmark
- Subscribe
- Mute
- Subscribe to RSS Feed
- Permalink
- Report Inappropriate Content
Hi XiYu,
Thanks for the question... the semantics of the sampler can be challenging. This question can probably help a lot of users figuring out the details of image sampling in OpenCL.
Using the CLK_FILTER_LINEAR sampler gives undefined results per the spec in this case... :
Please see Page 115 of the OpenCL-C 2.0 spec from Khronos:
The read_imagef calls that take integer coordinates must use a sampler with filter mode set to CLK_FILTER_NEAREST, normalized coordinates set to CLK_NORMALIZED_COORDS_FALSE and addressing mode set to CLK_ADDRESS_CLAMP_TO_EDGE, CLK_ADDRESS_CLAMP or CLK_ADDRESS_NONE; otherwise the values returned are undefined.
I think what's key here is to ask what you intend to interpolate between... If you intend to provide integer based unnormalized coordinates.... do you want to interpolate between anything at all? What would you expect to be linearly interpolating between?
Consider if you want to sample with normalized float coordinates (0.0 <-> 1.0) instead?
Thanks for asking such a productive question.
-MichaelC
Link Copied
- Mark as New
- Bookmark
- Subscribe
- Mute
- Subscribe to RSS Feed
- Permalink
- Report Inappropriate Content
In addition, I found that it get correct results when sampler is set as "CL_FILTER_NEAREST" .
- Mark as New
- Bookmark
- Subscribe
- Mute
- Subscribe to RSS Feed
- Permalink
- Report Inappropriate Content
Hi XiYu,
Thanks for the question... the semantics of the sampler can be challenging. This question can probably help a lot of users figuring out the details of image sampling in OpenCL.
Using the CLK_FILTER_LINEAR sampler gives undefined results per the spec in this case... :
Please see Page 115 of the OpenCL-C 2.0 spec from Khronos:
The read_imagef calls that take integer coordinates must use a sampler with filter mode set to CLK_FILTER_NEAREST, normalized coordinates set to CLK_NORMALIZED_COORDS_FALSE and addressing mode set to CLK_ADDRESS_CLAMP_TO_EDGE, CLK_ADDRESS_CLAMP or CLK_ADDRESS_NONE; otherwise the values returned are undefined.
I think what's key here is to ask what you intend to interpolate between... If you intend to provide integer based unnormalized coordinates.... do you want to interpolate between anything at all? What would you expect to be linearly interpolating between?
Consider if you want to sample with normalized float coordinates (0.0 <-> 1.0) instead?
Thanks for asking such a productive question.
-MichaelC
- Subscribe to RSS Feed
- Mark Topic as New
- Mark Topic as Read
- Float this Topic for Current User
- Bookmark
- Subscribe
- Printer Friendly Page