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];
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 contrast1D(unsigned char *grayImage, const int width, const int height, unsigned int *histogram, const unsigned int HISTOGRAM_SIZE, const unsigned int CONTRAST_THRESHOLD, NSTimer &timer) {
unsigned int i = 0;
NSTimer kernelTime = NSTimer("kernelTime", false, false);
while ( (i < HISTOGRAM_SIZE) && (histogram[i] < CONTRAST_THRESHOLD) ) {
i++;
}
unsigned int min = i;
i = HISTOGRAM_SIZE - 1;
while ( (i > min) && (histogram[i] < CONTRAST_THRESHOLD) ) {
i--;
}
unsigned int max = i;
float diff = max - min;
kernelTime.start();
// Kernel
for ( int y = 0; y < height; y++ ) {
for (int x = 0; x < width; x++ ) {
unsigned char pixel = grayImage[(y * width) + x];
if ( pixel < min ) {
pixel = 0;
}
else if ( pixel > max ) {
pixel = 255;
}
else {
pixel = static_cast< unsigned char >(255.0f * (pixel - min) / diff);
}
grayImage[(y * width) + x] = pixel;
}
}
// /Kernel
kernelTime.stop();
cout << fixed << setprecision(6);
cout << "contrast1D (kernel): \t\t" << kernelTime.getElapsed() << " seconds." << endl;
}
void histogram1D(unsigned char *grayImage, unsigned char *histogramImage, const int width, const int height, unsigned int *histogram, const unsigned int HISTOGRAM_SIZE, const unsigned int BAR_WIDTH, NSTimer &timer) {
unsigned int max = 0;
NSTimer kernelTime = NSTimer("kernelTime", false, false);
memset(reinterpret_cast< void * >(histogram), 0, HISTOGRAM_SIZE * sizeof(unsigned int));
kernelTime.start();
// Kernel
for ( int y = 0; y < height; y++ ) {
for ( int x = 0; x < width; x++ ) {
histogram[static_cast< unsigned int >(grayImage[(y * width) + x])] += 1;
}
}
// /Kernel
kernelTime.stop();
for ( unsigned int i = 0; i < HISTOGRAM_SIZE; i++ ) {
if ( histogram[i] > max ) {
max = histogram[i];
}
}
for ( int x = 0; x < HISTOGRAM_SIZE * BAR_WIDTH; x += BAR_WIDTH ) {
unsigned int value = HISTOGRAM_SIZE - ((histogram[x / BAR_WIDTH] * HISTOGRAM_SIZE) / max);
for ( unsigned int y = 0; y < value; y++ ) {
for ( unsigned int i = 0; i < BAR_WIDTH; i++ ) {
histogramImage[(y * HISTOGRAM_SIZE * BAR_WIDTH) + x + i] = 0;
}
}
for ( unsigned int y = value; y < HISTOGRAM_SIZE; y++ ) {
for ( unsigned int i = 0; i < BAR_WIDTH; i++ ) {
histogramImage[(y * HISTOGRAM_SIZE * BAR_WIDTH) + x + i] = 255;
}
}
}
cout << fixed << setprecision(6);
cout << "histogram1D (kernel): \t\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;
}