6 Replies Latest reply on Jan 20, 2011 6:42 AM by dal3boy

    Image 2D and greyscale pictures

    dal3boy

      I'm researching about error finding on 2D image (Grajscale for now), and couldn't find anything about read pixel colour (not RGB only greyscale)

      Now I have started reading little more on at Khronos Group’s website and is possible to get only one number per pixel (for greyscale only 0-255)?

      Like on picture attached, every pixel have int or float number to represent pixel colour

      In specification exist Cl_R or CL_INTENSITY may do single floats or int at each index, maybe I'm wrong.....

      Much easier would be to work only with one colour insted od RGB(3 colour)

       

        • Image 2D and greyscale pictures
          nou

          just create CL_R image and use read_image[f|i|ui]() functions. it is same as reading RGBA image.

          and just ignore that it return vector and use first value. more in specification section 6.11.13

            • Image 2D and greyscale pictures
              Meteorhead

              When I create images of channel order CL_R, CL_INTENSITY or CL_LUMINANCE, should I store the "trash" data also into the image in the host code, or the extra data will only be returned by the sampler, but actually they are not stored inside the image?

                • Image 2D and greyscale pictures
                  nou

                  thez will be just returned from read_image(). and they have specified that they are zeros. look at the end of chapter 6 of openCl specification.

                    • Image 2D and greyscale pictures
                      dal3boy

                      for now I'm on good way, just idea for kernel  pixel scaning

                      code below, suggest me if You have something to add

                      __kernel void scanImage(__read_only image2d_t src, some matrix will be output(later to implement)) { sampler_t smplr = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST; float4 curPixel; int w = get_image_width(src); int h = get_image_height(src); int i = 0; int j = 0; int2 x_y_axis; // Access rows for ( i = 0; i < h; i++) { x_y_axis[1] = i; // Access columns for ( j = 0; j < w; j++) { x_y_axis[0] = j; curTexel = read_imagef(src, smplr, x_y_axis); //some code to implement } } }

                        • Image 2D and greyscale pictures
                          nou

                          you use it like this

                          read_imagef(src, smplr, x_y_axis).x;

                          and curTexel as float.

                            • Image 2D and greyscale pictures
                              dal3boy

                              I have started on example, this only copy image (for now)

                              Now to find error on image(big colour differences) I have to scan pixel by pixel, so i have look Sobel filter image and I think I can make similar

                              example sobel filter image

                              float4 Gx = (float4)(0);
                              float4 Gy = Gx;
                                 
                              if( coord.x >= 1 && coord.x < (get_global_size(0)-1) && coord.y >= 1 && coord.y < get_global_size(1) - 1)
                                  {
                                      float4 i00 = convert_float4(read_imageui(imageIn, imgSampler, (int2)(coord.x - 1, coord.y + 1)));
                                      float4 i10 = convert_float4(read_imageui(imageIn, imgSampler, (int2)(coord.x - 0, coord.y + 1)));
                                      float4 i20 = convert_float4(read_imageui(imageIn, imgSampler, (int2)(coord.x + 1, coord.y + 1)));
                                      float4 i01 = convert_float4(read_imageui(imageIn, imgSampler, (int2)(coord.x - 1, coord.y + 0)));
                                      float4 i11 = convert_float4(read_imageui(imageIn, imgSampler, (int2)(coord.x - 0, coord.y + 0)));
                                      float4 i21 = convert_float4(read_imageui(imageIn, imgSampler, (int2)(coord.x + 1, coord.y + 0)));
                                      float4 i02 = convert_float4(read_imageui(imageIn, imgSampler, (int2)(coord.x - 1, coord.y - 1)));
                                      float4 i12 = convert_float4(read_imageui(imageIn, imgSampler, (int2)(coord.x - 0, coord.y - 1)));
                                      float4 i22 = convert_float4(read_imageui(imageIn, imgSampler, (int2)(coord.x + 1, coord.y - 1)));

                                      Gx =   i00 + (float4)(2) * i10 + i20 - i02  - (float4)(2) * i12 - i22;

                                      Gy =   i00 - i20  + (float4)(2)*i01 - (float4)(2)*i21 + i02  -  i22;

                                      Gx = native_divide(native_sqrt(Gx * Gx + Gy * Gy), (float4)(2));

                                      write_imageui(outputImage, coord, convert_uint4(Gx));
                                  }

                              if I sum all 8 neighbouring pixels (something like up in Sobel) (RGBA) and find BIG differences between them and write out pixel (black or white to indicate start of error on picture)

                              If is that possible to do like this

                              i11 is our readed pixel

                              float4 black = (float4)(0);

                              sum =(i00+i10+i20+i01+i21+i02+i12+i22)/8;

                              let say colour may be tolerated by 20

                              if(i11<sum-20 || i>sum+20)

                              { write_imageui(imageOut, coord, black); }

                              write_imageui(imageOut, coord, pixel);

                               

                               

                              //imgKernel.cl __constant sampler_t imgSampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_LINEAR; __kernel void imgScan(__read_only image2d_t imageIn,__write_only image2d_t imageOut) { int2 coord = (int2)(get_global_id(0), get_global_id(1)); uint4 pixel; pixel=read_imageui(imageIn,imgSampler,coord); write_imageui (imageOut,coord,pixel); } /main.cpp #include <stdio.h> #include <stdlib.h> #include "CL/cl.h" #define BLOCK_SIZE 16 int main() { // OpenCL varijable cl_uint num_devs_returned; cl_context_properties properties[3]; cl_device_id device_id; cl_int err; cl_platform_id platform_id; cl_uint num_platforms_returned; cl_context context; cl_command_queue command_queue; cl_program program; cl_kernel kernel; size_t global[2]; size_t local[2]; size_t result=0; unsigned char header [54]; FILE *input = fopen("lena256.bmp", "rb"); result = fread(header,1,54,input); fseek (input, 54, SEEK_SET); unsigned char *tab; tab = (unsigned char*)malloc(256*256*4); result = fread(tab, 1, 3*256*256, input); fclose(input); void * image= (unsigned char *)tab; unsigned char *tab2; tab2 = (unsigned char*)malloc(256*256*4); void * image2= (unsigned char *)tab2; FILE * inputF = fopen("imgKernel.cl", "rb");; size_t size = 0; char* binary = NULL; if(inputF == NULL) { return false; } fseek(inputF, 0L, SEEK_END); size = ftell(inputF); rewind(inputF); binary = (char*)malloc(size); if(binary == NULL) { return false; } fread(binary, sizeof(char), size, inputF); fclose(inputF); //printf("kernel: %s",binary); // dobivanje platformID err = clGetPlatformIDs(1,&platform_id,&num_platforms_returned); if (err != CL_SUCCESS) { printf("clGetPlatformIDs error = %d\n",err); exit(1); } // GPU ili CPU err = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, &num_devs_returned); if (err != CL_SUCCESS) { printf("clGetDeviceIDs error = %d\n",err); exit(1); } // CL_CONTEXT_PLATFORM postavljanje, mora završiti s 0 properties[0]= CL_CONTEXT_PLATFORM; properties[1]= (cl_context_properties) platform_id; properties[2]= 0; // postavljanje parametara context = clCreateContext(properties, 1, &device_id, NULL, NULL, &err); if (err != CL_SUCCESS) { printf("clCreateContext error = %d\n",err); exit(1); } // kreiranje naredbe za izvršavanje kernela command_queue = clCreateCommandQueue(context,device_id, 0, &err); if (err != CL_SUCCESS) { printf("clCreateCommandQueue error = %d\n",err); exit(1); } // kreiranje programa iz gore navedenog kernel koda program = clCreateProgramWithSource(context, 1 ,(const char **) &binary, NULL, &err); if (err != CL_SUCCESS) { printf("clCreateProgramWithSource error = %d\n",err); exit(1); } //kompajliranje kernek programa err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL); if (err != CL_SUCCESS) { printf("clBuildProgram error = %d\n", err); size_t len; char buffer[4096]; // info clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len); printf("--- Info -- \n %s\n",buffer); exit(1); } //kreiranje kernela GPU kernel = clCreateKernel(program, "imgScan", &err); if (err != CL_SUCCESS) { printf("clCreateKernel error = %d\n",err); exit(1); } size_t width = 256; size_t height = 256; size_t rowpitch = 0; cl_image_format format; format.image_channel_order = CL_RGBA; format.image_channel_data_type = CL_UNSIGNED_INT8; cl_mem_flags flags; flags = CL_MEM_READ_ONLY; cl_mem Image = clCreateImage2D(context, flags, &format, width, height, rowpitch, 0, &err); if (err != CL_SUCCESS) { printf("clCreateImage2D error = %d\n",err); exit(1); } cl_mem_flags flags2; flags2 = CL_MEM_WRITE_ONLY; cl_mem Image2 = clCreateImage2D(context, flags2, &format, width, height, rowpitch, 0,&err); if (err != CL_SUCCESS) { printf("clCreateImage2D error = %d\n",err); exit(1); } // postavljanje kernel argumenata if ( clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&Image) || clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*)&Image2) != CL_SUCCESS) { printf("clSetKernelArg error = %d\n",err); exit(1); } size_t origin[3]; origin[0] = 0; origin[1] = 0; origin[2] = 0; size_t region[3]; region[0] = width; region[1] = height; region[2] = 1; err = clEnqueueWriteImage(command_queue, Image, CL_TRUE, origin, region, width*sizeof(char)*4, 0, image, 0, NULL,NULL); if (err != CL_SUCCESS) { printf("clEnqueueWriteImage error = %d\n",err); exit(1); } // postavljanje global & local radne velicine global[0]= 256; global[1]= 256; local[0]=BLOCK_SIZE; local[1]=BLOCK_SIZE; // slanje kernel objekta na izvodjenje s prije navedenim parametrima err = clEnqueueNDRangeKernel(command_queue, kernel, 2, NULL, global,local, 0, NULL, NULL); if (err != CL_SUCCESS) { printf("clEnqueueNDRangeKernel error = %d\n",err); exit(1); } // dekanje izvrsenje kernel naredbe clFinish(command_queue); err = clEnqueueReadImage (command_queue, Image2, CL_TRUE, origin, region, 0, 0, image2, 0, NULL, NULL); if (err != CL_SUCCESS) { printf("clEnqueueReadImage error = %d\n",err); exit(1); } FILE *nk = fopen("outImg.bmp", "wb"); fwrite(header,1,54,nk); fwrite(image2, 1, (3*256*256), nk); fclose(nk); // oslobodjenje memorije clReleaseProgram(program); clReleaseKernel(kernel); clReleaseCommandQueue(command_queue); clReleaseContext(context); return 0; }