//(__read_only image2d_t input, __read_only image2d_t input2,__write_only image2d_t output, const int x1, const int x2, const int y1, const int y2, const int op)
void overlay(Image * I, Image * other, int x1, int y1, int op)
{
setFormat(I,1);
setFormat(other,1);
cl_int result;
size_t gsize[2];
gsize[0] = I->width;
gsize[1] = I->height;
cl_mem tmp = clCreateImage2D(context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, &format, I->width, I->height, 0,NULL,&result);
size_t lsize[2];
lsize[0] = getLocalSize(I->width);
lsize[1] = getLocalSize(I->height);
check(result);
check(clSetKernelArg(overlay_kernel,0,sizeof(cl_mem),&I->GPU_Image));
check(clSetKernelArg(overlay_kernel,1,sizeof(cl_mem),&other->GPU_Image));
check(clSetKernelArg(overlay_kernel,2,sizeof(cl_mem),&tmp));
check(clSetKernelArg(overlay_kernel,3,sizeof(int),&x1));
check(clSetKernelArg(overlay_kernel,4,sizeof(int),&other->width));
check(clSetKernelArg(overlay_kernel,5,sizeof(int),&y1));
check(clSetKernelArg(overlay_kernel,6,sizeof(int),&other->height));
check(clSetKernelArg(overlay_kernel,7,sizeof(int),&op));
check(clEnqueueNDRangeKernel(queue, overlay_kernel, 2, 0, gsize, lsize, 0 ,NULL ,NULL));
clReleaseMemObject(I->GPU_Image);
I->GPU_Image = tmp;
}
void avgCase(Image * I, Image * other, convMatrix * matrix)
{
matrix_transfer(matrix,1);
transfer(I,1);
transfer(other,1);
cl_int result;
size_t gsize[2];
gsize[0] = other->width;
gsize[1] = other->height;
I->width = other->width;
I->height = other->height;
size_t lsize[2];
lsize[0] = getLocalSize(I->width);
lsize[1] = getLocalSize(I->height);
cl_mem tmp = clCreateImage2D(context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, &format, I->width, I->height, 0, NULL, &result);
check(result);
check(clSetKernelArg(avgCase_kernel,0,sizeof(cl_mem),&I->GPU_Image));
check(clSetKernelArg(avgCase_kernel,1,sizeof(cl_mem),&other->GPU_Image));
check(clSetKernelArg(avgCase_kernel,2,sizeof(cl_mem),&tmp));
check(clSetKernelArg(avgCase_kernel,3,sizeof(cl_mem),&matrix->GPU_Matrix));
check(clSetKernelArg(avgCase_kernel,4,sizeof(int),&matrix->dim));
check(clSetKernelArg(avgCase_kernel,5,sizeof(float),&matrix->div));
check(clEnqueueNDRangeKernel(queue, avgCase_kernel , 2, 0, gsize, NULL, 0, NULL, NULL));
clReleaseMemObject(I->GPU_Image);
free(I->CPU_Image);
I->CPU_Image = NULL;
I->CPU_Image = (TYPE*) malloc(sizeof(TYPE)*I->width*I->height*DEPTH);
I->GPU_Image = tmp;
}
void bestCase(Image * I, int width_,int height_,float color)
{
transfer(I,1);
cl_int result;
size_t gsize[2];
gsize[0] = width_;
gsize[1] = height_;
I->width = width_;
I->height = height_;
size_t lsize[2];
lsize[0] = getLocalSize(I->width);
lsize[1] = getLocalSize(I->height);
cl_mem tmp = clCreateImage2D(context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, &format, I->width, I->height, 0, NULL, &result);
check(result);
check(clSetKernelArg(bestCase_kernel,0,sizeof(cl_mem),&I->GPU_Image));
check(clSetKernelArg(bestCase_kernel,1,sizeof(cl_mem),&tmp));
check(clSetKernelArg(bestCase_kernel,2,sizeof(float),&color));
check(clEnqueueNDRangeKernel(queue, bestCase_kernel , 2, 0, gsize, NULL, 0, NULL, NULL));
clReleaseMemObject(I->GPU_Image);
free(I->CPU_Image);
I->CPU_Image = NULL;
I->CPU_Image = (TYPE*) malloc(sizeof(TYPE)*I->width*I->height*DEPTH);
I->GPU_Image = tmp;
}
void PETScVector::copyValues(std::vector<double>& u) const
{
assert(u.size() == (std::size_t) (getLocalSize() + getGhostSize()));
double* loc_x = getLocalVector();
std::copy_n(loc_x, getLocalSize() + getGhostSize(), u.begin());
restoreArray(loc_x);
}
void convolve(Image * I, ConvMatrix * matrix)
{
matrix_transfer(matrix,1);
transfer(I,1);
cl_int result;
size_t gsize[2];
gsize[0] = I->width;
gsize[1] = I->height;
size_t lsize[2];
lsize[0] = getLocalSize(I->width);
lsize[1] = 1;
cl_mem tmp = clCreateImage2D(context,CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, &format, I->width, I->height, 0,NULL,&result);
check(clSetKernelArg(convolution_kernel,0,sizeof(cl_mem),&I->GPU_Image));
check(clSetKernelArg(convolution_kernel,1,sizeof(cl_mem),&tmp));
check(clSetKernelArg(convolution_kernel,2,sizeof(cl_mem),&matrix->GPU_Matrix));
check(clSetKernelArg(convolution_kernel,3,sizeof(int),&matrix->dim));
check(clSetKernelArg(convolution_kernel,4,sizeof(float),&matrix->div));
check(clEnqueueNDRangeKernel(queue, convolution_kernel, 2, 0, gsize, NULL, 0, NULL, NULL));
clReleaseMemObject(I->GPU_Image);
I->GPU_Image = tmp;
}
void toRGB(Image * I)
{
I->type = 0;
transfer(I,1);
cl_int result;
size_t gsize[2];
gsize[0] = I->width;
gsize[1] = I->height;
size_t lsize[2];
lsize[0] = getLocalSize(I->width);
lsize[1] = getLocalSize(I->height);
cl_mem tmp = clCreateImage2D(context,CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, &format, I->width, I->height, 0,NULL,&result);
check(result);
check(clSetKernelArg(HSVtoRGB_kernel,0,sizeof(cl_mem),&I->GPU_Image));
check(clSetKernelArg(HSVtoRGB_kernel,1,sizeof(cl_mem),&tmp));
check(clEnqueueNDRangeKernel(queue,HSVtoRGB_kernel,2,0,gsize,NULL,0,NULL,NULL));
clReleaseMemObject(I->GPU_Image);
I->GPU_Image = tmp;
}
void worstcase(Image * I, Image * other, convMatrix * blur,convMatrix * blur2)
{
matrix_transfer(blur,1);
matrix_transfer(blur2,1);
transfer(I,1);
transfer(other,1);
cl_int result;
size_t gsize[2];
gsize[0] = other->width;
gsize[1] = other->height;
I->width = other->width;
I->height = other->height;
size_t lsize[2];
lsize[0] = getLocalSize(I->width);
lsize[1] = getLocalSize(I->height);
cl_mem tmp = clCreateImage2D(context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, &format, I->width, I->height, 0, NULL, &result);
check(result);
//worstCase(__read_only image2d_t input,__read_only image2d_t input2, __write_only image2d_t output,__global TYPE * matrix,__global TYPE * matrix2, const int dim, const float div)
check(clSetKernelArg(worstCase_kernel,0,sizeof(cl_mem),&I->GPU_Image));
check(clSetKernelArg(worstCase_kernel,1,sizeof(cl_mem),&other->GPU_Image));
check(clSetKernelArg(worstCase_kernel,2,sizeof(cl_mem),&tmp));
check(clSetKernelArg(worstCase_kernel,3,sizeof(cl_mem),&blur->GPU_Matrix));
check(clSetKernelArg(worstCase_kernel,4,sizeof(cl_mem),&blur2->GPU_Matrix));
check(clSetKernelArg(worstCase_kernel,5,sizeof(int),&blur->dim));
check(clSetKernelArg(worstCase_kernel,6,sizeof(float),&blur->div));
check(clEnqueueNDRangeKernel(queue, worstCase_kernel , 2, 0, gsize, NULL, 0, NULL, NULL));
clReleaseMemObject(I->GPU_Image);
free(I->CPU_Image);
I->CPU_Image = NULL;
I->CPU_Image = (TYPE*) malloc(sizeof(TYPE)*I->width*I->height*DEPTH);
I->GPU_Image = tmp;
}
void crop(Image * I, int x1, int x2, int y1, int y2)
{
transfer(I,1);
cl_int result;
size_t gsize[2];
gsize[0] = x2 - x1;
gsize[1] = y2 - y1;
I->width = x2 - x1;
I->height = y2 - y1;
cl_mem tmp = clCreateImage2D(context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, &format, I->width, I->height, 0,NULL,&result);
check(result);
size_t lsize[2];
lsize[0] = getLocalSize(I->width);
lsize[1] = getLocalSize(I->height);
check(clSetKernelArg(crop_kernel, 0, sizeof(cl_mem), &I->GPU_Image));
check(clSetKernelArg(crop_kernel, 1, sizeof(cl_mem), &tmp));
check(clSetKernelArg(crop_kernel, 2, sizeof(int), &x1));
check(clSetKernelArg(crop_kernel, 3, sizeof(int), &y1));
check(clEnqueueNDRangeKernel(queue,resize_kernel,2,0,gsize,lsize,0,NULL,NULL));
clReleaseMemObject(I->GPU_Image);
I->GPU_Image = tmp;
}
