void triangularSmooth(unsigned char *grayImage, unsigned char *smoothImage, const int width, const int height, const float *filter, NSTimer &timer) { NSTimer kernelTime = NSTimer("kernelTime", false, false); kernelTime.start(); // Kernel for ( int y = 0; y < height; y++ ) { for ( int x = 0; x < width; x++ ) { unsigned int filterItem = 0; float filterSum = 0.0f; float smoothPix = 0.0f; for ( int fy = y - 2; fy < y + 3; fy++ ) { for ( int fx = x - 2; fx < x + 3; fx++ ) { if ( ((fy < 0) || (fy >= height)) || ((fx < 0) || (fx >= width)) ) { filterItem++; continue; } smoothPix += grayImage[(fy * width) + fx] * filter[filterItem]; //if (grayImage[(fy*width) + fx] == 0) //printf("Yesssssssssss\n"); filterSum += filter[filterItem]; filterItem++; } } smoothPix /= filterSum; smoothImage[(y * width) + x] = static_cast< unsigned char >(smoothPix); } } // /Kernel kernelTime.stop(); cout << fixed << setprecision(6); cout << "triangularSmooth (kernel): \t" << kernelTime.getElapsed() << " seconds." << endl; }
void rgb2gray(unsigned char *inputImage, unsigned char *grayImage, const int width, const int height, NSTimer &timer) { NSTimer kernelTime = NSTimer("kernelTime", false, false); kernelTime.start(); // Kernel for ( int y = 0; y < height; y++ ) { for ( int x = 0; x < width; x++ ) { float grayPix = 0.0f; float r = static_cast< float >(inputImage[(y * width) + x]); float g = static_cast< float >(inputImage[(width * height) + (y * width) + x]); float b = static_cast< float >(inputImage[(2 * width * height) + (y * width) + x]); grayPix = (0.3f * r) + (0.59f * g) + (0.11f * b); grayImage[(y * width) + x] = static_cast< unsigned char >(grayPix); } } // /Kernel kernelTime.stop(); cout << fixed << setprecision(6); cout << "rgb2gray (kernel): \t\t" << kernelTime.getElapsed() << " seconds." << endl; }
int main(int argc, char *argv[]) { NSTimer total = NSTimer("total", false, false); if ( argc != 2 ) { cerr << "Usage: " << argv[0] << " <filename>" << endl; return 1; } // Load the input image CImg< unsigned char > inputImage = CImg< unsigned char >(argv[1]); if ( displayImages ) { inputImage.display("Input Image"); } if ( inputImage.spectrum() != 3 ) { cerr << "The input must be a color image." << endl; return 1; } // Image is loaded, start timing total.start(); // Convert the input image to grayscale CImg< unsigned char > grayImage = CImg< unsigned char >(inputImage.width(), inputImage.height(), 1, 1); rgb2gray(inputImage.data(), grayImage.data(), inputImage.width(), inputImage.height(), total); total.stop(); if ( displayImages ) { grayImage.display("Grayscale Image"); } if ( saveAllImages ) { grayImage.save("./grayscale.bmp"); } total.start(); // Compute 1D histogram CImg< unsigned char > histogramImage = CImg< unsigned char >(BAR_WIDTH * HISTOGRAM_SIZE, HISTOGRAM_SIZE, 1, 1); unsigned int *histogram = new unsigned int [HISTOGRAM_SIZE]; histogram1D(grayImage.data(), histogramImage.data(), grayImage.width(), grayImage.height(), histogram, HISTOGRAM_SIZE, BAR_WIDTH, total); total.stop(); if ( displayImages ) { histogramImage.display("Histogram"); } if ( saveAllImages ) { histogramImage.save("./histogram.bmp"); } total.start(); // Contrast enhancement contrast1D(grayImage.data(), grayImage.width(), grayImage.height(), histogram, HISTOGRAM_SIZE, CONTRAST_THRESHOLD, total); total.stop(); if ( displayImages ) { grayImage.display("Contrast Enhanced Image"); } if ( saveAllImages ) { grayImage.save("./contrast.bmp"); } total.start(); delete [] histogram; // Triangular smooth (convolution) CImg< unsigned char > smoothImage = CImg< unsigned char >(grayImage.width(), grayImage.height(), 1, 1); triangularSmooth(grayImage.data(), smoothImage.data(), grayImage.width(), grayImage.height(), filter, total); // Job done, stop timing total.stop(); if ( displayImages ) { smoothImage.display("Smooth Image"); } smoothImage.save("./smooth.bmp"); // Wrap up cout << fixed << setprecision(6) << endl; cout << "Execution time: \t\t" << total.getElapsed() << " seconds." << endl << endl; return 0; }
int main(int argc, char *argv[]) { int err; cl_device_id device_id; // compute device id cl_context context; // compute context cl_command_queue commands; // compute command queue cl_program program; // compute program cl_kernel ko_vadd; // compute kernel //char *kernel_source = getKernelSource("filters.cl"); //buffers cl_mem d_img; // device memory used for the input a vector cl_mem d_gry; FILE *fp; char fileName[] = "filters.cl"; char *source_str; size_t source_size; size_t global; /* Load the source code containing the kernel*/ fp = fopen(fileName, "r"); if (!fp) { fprintf(stderr, "Failed to load kernel.\n"); exit(1); } source_str = (char*)malloc(MAX_SOURCE_SIZE); source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp); fclose(fp); //printf("Start 1 \n"); //===================openCL seetup======================================== // Set up platform and GPU device cl_uint numPlatforms; // Find number of platforms err = clGetPlatformIDs(0, NULL, &numPlatforms); checkError(err, "Finding platforms"); if (numPlatforms == 0) { printf("Found 0 platforms!\n"); return EXIT_FAILURE; } // Get all platforms cl_platform_id Platform[numPlatforms]; err = clGetPlatformIDs(numPlatforms, Platform, NULL); checkError(err, "Getting platforms"); // Secure a GPU for (int i = 0; i < numPlatforms; i++) { err = clGetDeviceIDs(Platform[i], DEVICE, 1, &device_id, NULL); if (err == CL_SUCCESS) { break; } } if (device_id == NULL) checkError(err, "Finding a device"); err = output_device_info(device_id); checkError(err, "Printing device output"); // Create a compute context context = clCreateContext(0, 1, &device_id, NULL, NULL, &err); checkError(err, "Creating context"); // Create a command queue commands = clCreateCommandQueue(context, device_id, CL_QUEUE_PROFILING_ENABLE, &err); checkError(err, "Creating command queue"); // Create the compute program from the source buffer //program = clCreateProgramWithSource(context, 1, (const char **) & KernelSource, NULL, &err); program = clCreateProgramWithSource(context, 1, (const char **)&source_str,(const size_t *)&source_size, &err); checkError(err, "Creating program"); //printf("Start 2 \n"); // Build the program err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL); if (err != CL_SUCCESS) { size_t len; char buffer[2048]; printf("Error: Failed to build program executable!\n%s\n", err_code(err)); clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len); printf("%s\n", buffer); return EXIT_FAILURE; } // Create the compute kernel from the program ko_vadd = clCreateKernel(program, "rgbnaargrey", &err); checkError(err, "Creating kernel"); //printf("Start 3 \n"); //===================openCL seetup END======================================== if ( argc != 2 ) { cerr << "Usage: " << argv[0] << " <filename>" << endl; return 1; } // Load the input image CImg< unsigned char > inputImage = CImg< unsigned char >(argv[1]); if ( displayImages ) { inputImage.display("Input Image"); } if ( inputImage.spectrum() != 3 ) { cerr << "The input must be a color image." << endl; return 1; } // Convert the input image to grayscale CImg< unsigned char > grayImage = CImg< unsigned char >(inputImage.width(), inputImage.height(), 1, 1); //===================openCL buffers======================================== int w=inputImage.width(); int h=inputImage.height(); //d_img=clCreateImage(context,CL_MEM_READ_ONLY,_cl_image_format(CL_RGB ,CL_UNORM_SHORT_555),); printf("\n Width=%d Height=%d \n", w,h); d_img= clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(unsigned char) *w*h*3, NULL, &err); d_gry= clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(unsigned char) *w*h, NULL, &err); //printf("Start 4 \n"); err = clEnqueueWriteBuffer(commands, d_img, CL_TRUE, 0, sizeof(unsigned char) *w*h*3, inputImage, 0, NULL, NULL); checkError(err, "Copying h_a to device at d_a"); //printf("Start 4 .1 \n"); // Set the arguments to our compute kernel err = clSetKernelArg(ko_vadd, 0, sizeof(cl_mem), &d_img); //printf("Start 4 .2 \n"); err |= clSetKernelArg(ko_vadd, 1, sizeof(cl_mem), &d_gry); //printf("Start 4 .3 \n"); err |= clSetKernelArg(ko_vadd, 2, sizeof(unsigned int), &w); //printf("Start 4 .4 \n"); err |= clSetKernelArg(ko_vadd, 3, sizeof(unsigned int), &h); checkError(err, "Setting kernel arguments"); //printf("Start 5 \n"); //===================openCL kernel ======================================== global = w*h; NSTimer kernelTime = NSTimer("kernelTime", false, false); cl_event event0; clEnqueueNDRangeKernel(commands, ko_vadd, 1, NULL, &global, NULL, 0, NULL, &event0); clWaitForEvents (1, &event0); /* unsigned long start = 0; unsigned long end = 0; clGetEventProfilingInfo(event0,CL_PROFILING_COMMAND_START, sizeof(cl_ulong),&start,NULL); clGetEventProfilingInfo(event0,CL_PROFILING_COMMAND_END, sizeof(cl_ulong),&end,NULL); */ cl_ulong time_start, time_end; double total_time; clGetEventProfilingInfo(event0, CL_PROFILING_COMMAND_START, sizeof(time_start), &time_start, NULL); clGetEventProfilingInfo(event0, CL_PROFILING_COMMAND_END, sizeof(time_end), &time_end, NULL); total_time = time_end - time_start; printf("\nExecution time in milliseconds = %0.3f ms\n", (total_time / 1000000.0) ); // Compute the duration in nanoseconds //unsigned long duration = end - start; /*printf("rgb2gray (gpu):%lf \n",start); printf("rgb2gray (gpu):%lf \n",end); printf("rgb2gray (gpu):%lf \n",duration); */ //checkError(err, "Enqueueing kernel"); //kernelTime.start(); // Wait for the commands to complete before stopping the timer err = clFinish(commands); checkError(err, "Waiting for kernel to finish"); //kernelTime.stop(); //cout << fixed << setprecision(6); //cout << "rgb2gray (gpu): \t\t" << kernelTime.getElapsed() << " seconds." << endl; //printf("Start 6 \n"); err = clEnqueueReadBuffer( commands, d_gry, CL_TRUE, 0, sizeof(unsigned char) *w*h, grayImage, 0, NULL, NULL ); if (err != CL_SUCCESS) { printf("Error: Failed to read output array!\n%s\n", err_code(err)); exit(1); } //printf("Start 7 \n"); //===================openCL kernel END======================================== rgb2gray(inputImage.data(), grayImage.data(), inputImage.width(), inputImage.height()); //rgb2grayCuda if ( displayImages ) { grayImage.display("Grayscale Image"); } if ( saveAllImages ) { grayImage.save("./grayscale.bmp"); } // Compute 1D histogram CImg< unsigned char > histogramImage = CImg< unsigned char >(BAR_WIDTH * HISTOGRAM_SIZE, HISTOGRAM_SIZE, 1, 1); unsigned int *histogram = new unsigned int [HISTOGRAM_SIZE]; //histogram1D(grayImage.data(), histogramImage.data(), grayImage.width(), grayImage.height(), histogram, HISTOGRAM_SIZE, BAR_WIDTH); //histogram1DCuda if ( displayImages ) { histogramImage.display("Histogram"); } if ( saveAllImages ) { histogramImage.save("./histogram.bmp"); } // Contrast enhancement //contrast1D(grayImage.data(), grayImage.width(), grayImage.height(), histogram, HISTOGRAM_SIZE, CONTRAST_THRESHOLD); //contrast1DCuda if ( displayImages ) { grayImage.display("Contrast Enhanced Image"); } if ( saveAllImages ) { grayImage.save("./contrast.bmp"); } delete [] histogram; // Triangular smooth (convolution) CImg< unsigned char > smoothImage = CImg< unsigned char >(grayImage.width(), grayImage.height(), 1, 1); //triangularSmooth(grayImage.data(), smoothImage.data(), grayImage.width(), grayImage.height(), filter); //triangularSmoothCuda if ( displayImages ) { smoothImage.display("Smooth Image"); } if ( saveAllImages ) { smoothImage.save("./smooth.bmp"); } clReleaseMemObject(d_img); clReleaseMemObject(d_gry); //clReleaseMemObject(d_c); clReleaseEvent(event0); clReleaseProgram(program); clReleaseKernel(ko_vadd); clReleaseCommandQueue(commands); clReleaseContext(context); /* free(inputImage.data()); printf("Start 7.1 \n"); free(grayImage.data()); printf("Start 7.2 \n"); //free(histogram); printf("Start 7 .3 \n"); //free(smoothImage.data()); printf("Start 7.4 \n");*/ return 0; }