inline void storeFramebuffer(int resultIndex,
const FilePath& pngDir,
const FilePath& exrDir,
const FilePath& baseName,
float gamma,
const Framebuffer& framebuffer) {
// Create output directories in case they don't exist
createDirectory(pngDir.str());
createDirectory(exrDir.str());
// Path of primary output image files
FilePath pngFile = pngDir + baseName.addExt(".png");
FilePath exrFile = exrDir + baseName.addExt(".exr");
// Make a copy of the image
Image copy = framebuffer.getChannel(0);
// Store the EXR image "as is"
storeEXRImage(exrFile.str(), copy);
// Post-process copy of image and store PNG file
copy.flipY();
copy.divideByAlpha();
copy.applyGamma(gamma);
storeImage(pngFile.str(), copy);
// Store complete framebuffer as a single EXR multi layer image
auto exrFramebufferFilePath = exrDir + baseName.addExt(".bnzframebuffer.exr");
storeEXRFramebuffer(exrFramebufferFilePath.str(), framebuffer);
// Store complete framebuffer as PNG indivual files in a dediacted subdirectory
auto pngFramebufferDirPath = pngDir + "framebuffers/";
createDirectory(pngFramebufferDirPath.str());
if(resultIndex >= 0) {
pngFramebufferDirPath = pngFramebufferDirPath + toString3(resultIndex); // Split different results of the same batch in different subdirectories
}
createDirectory(pngFramebufferDirPath.str());
// Store each channel of the framebuffer
for(auto i = 0u; i < framebuffer.getChannelCount(); ++i) {
auto name = framebuffer.getChannelName(i);
// Prefix by the index of the channel
FilePath pngFile = pngFramebufferDirPath + FilePath(toString3(i)).addExt("_" + name + ".png");
auto copy = framebuffer.getChannel(i);
copy.flipY();
copy.divideByAlpha();
copy.applyGamma(gamma);
storeImage(pngFile.str(), copy);
}
}
void renderToFile(const FileName& fileName)
{
resize(g_width,g_height);
AffineSpace3fa pixel2world = g_camera.pixel2world(g_width,g_height);
render(0.0f,pixel2world.l.vx,pixel2world.l.vy,pixel2world.l.vz,pixel2world.p);
void* ptr = map();
Ref<Image> image = new Image4c(g_width, g_height, (Col4c*)ptr);
storeImage(image, fileName);
unmap();
}
static void outputMode(const FileName& fileName)
{
if (!g_renderer) throw std::runtime_error("no renderer set");
/* render */
Ref<Device::RTCamera> camera = createCamera(AffineSpace::lookAtPoint(g_camPos,g_camLookAt,g_camUp));
Ref<Device::RTScene> scene = createScene(g_scene.cast<Scene>());
g_device->rtRenderFrame(g_renderer,camera,scene,g_frameBuffer);
/* store to disk */
void* ptr = g_device->rtMapFrameBuffer(g_frameBuffer);
Ref<Image3f> image = new Image3f(g_width,g_height);
memcpy(&image->data[0],ptr,g_width*g_height*sizeof(Col3f));
storeImage(image.cast<Image>(),fileName);
g_device->rtUnmapFrameBuffer(g_frameBuffer);
g_rendered = true;
}
void PHISHtmlGeneric::diaShow() const
{
bool useTmp( false );
int i;
QByteArray postfix;
QList <QByteArray> ids;
QString imageId;
QTransform t=transformation();
if ( _it->hasGraphicEffect() ) postfix=QByteArray::fromRawData( "_g", 2 );
if ( !t.isIdentity() ) postfix=QByteArray::fromRawData( "_t", 2 );
for ( i=0; i<_it->pictureBookIds().count(); i++ ) {
imageId=_it->pictureBookIds().at( i );
// check if we have a graphical changed picture
if ( !postfix.isNull() || imageId.startsWith( QLatin1String( "phi" ) )
|| (_it->itemProperties() & PHIItem::PNoCache) ) {
imageId=checkForImageId( postfix, i ); // image already cached?
useTmp=true;
if ( imageId.isNull() ) {
imageId=_it->pictureBookIds().at( i );
QImage img=createImageImage( imageId );
if ( _it->hasGraphicEffect() ) img=graphicsEffectImage( img );
if ( !t.isIdentity() ) img=img.transformed( t, Qt::SmoothTransformation );
imageId=storeImage( img, postfix, i );
if ( imageId.isNull() ) return;
}
}
ids.append( imageId.toUtf8() );
}
i=0;
QByteArray src, imgid;
QByteArray style=" style=\"position:absolute;"+transformationPositionStyle( t, true );
QList <QByteArray> titles=_it->toolTipData().split(':');
_out+=_indent+"<div "+id();
if ( titles.count()==1 ) _out+=title();
_out+=style+"\">\n";
style=" style=\"left:0;top:0;";
if ( t.isIdentity() ) style+="width:"+QByteArray::number( _it->width() )+"px;height:"
+QByteArray::number( _it->height() )+"px;";
foreach ( src, ids ) {
src="/phi.phis?phiimg="+src+"";
if ( useTmp ) src+="&phitmp=1";
imgid=_it->id()+"_phi_"+QByteArray::number( i );
_out+='\t';
imageSource( src, style, imgid, (titles.count()>i ? titles.at( i ): QByteArray()) );
i++;
}
static CONDITION
storageCallback(MSG_C_STORE_REQ * request, MSG_C_STORE_RESP * response,
unsigned long received, unsigned long estimate,
/* DCM_OBJECT ** object, STORAGE_PARAMS * params, */
DCM_OBJECT ** object, void *paramsPtr,
DUL_PRESENTATIONCONTEXT * pc)
{
CONDITION
cond = 0;
IE_OBJECT
* ieObject;
/* The definition and assignment help satisfy prototypes for this callback
** function (defined by dicom_services.h)
*/
STORAGE_PARAMS *params;
params = (STORAGE_PARAMS *) paramsPtr;
if (!silent)
printf("%8d bytes received of %8d estimated \n", received, estimate);
if (!silent && (object != NULL))
(void) DCM_DumpElements(object, 0);
if (object != NULL) {
if (doVerification) {
cond = IE_ExamineObject(object, &ieObject);
(void) IE_Free((void **) &ieObject);
if (cond != IE_NORMAL) { /* The image does not satisfy */
/* Part 3 requirements */
response->status = MSG_K_C_STORE_DATASETNOTMATCHSOPCLASSERROR;
return SRV_NORMAL;
}
}
cond = storeImage(params->handle, params->fileName,
params->transferSyntax, request->classUID,
params->owner, params->groupName, params->priv,
object);
}
if (response != NULL) {
if (cond == APP_NORMAL)
response->status = MSG_K_SUCCESS;
else
response->status = MSG_K_C_STORE_OUTOFRESOURCES;
}
return SRV_NORMAL;
}
void renderToFile(const FileName& fileName)
{
resize(g_width,g_height);
if (g_anim_mode) g_camera.anim = true;
do {
double msec = getSeconds();
AffineSpace3fa pixel2world = g_camera.pixel2world(g_width,g_height);
render(0.0f,pixel2world.l.vx,pixel2world.l.vy,pixel2world.l.vz,pixel2world.p);
msec = getSeconds() - msec;
std::cout << "render time " << 1.0/msec << " fps" << std::endl;
} while(g_loop_mode);
void* ptr = map();
Ref<Image> image = new Image4uc(g_width, g_height, (Col4uc*)ptr);
storeImage(image, fileName);
unmap();
cleanup();
}
//using namespace cl;
int main(int argc, char ** argv)
{
try {
cl_int err;
cl::vector<cl::Platform> platforms;
cl::Platform::get(&platforms);
std::cout << "Number of platforms:\t" << platforms.size() << std::endl;
for (cl::vector<cl::Platform>::iterator i = platforms.begin(); i != platforms.end(); ++i) {
// pick a platform and do something
std::cout << " Platform Name: " << (*i).getInfo<CL_PLATFORM_NAME>().c_str()<< std::endl;
}
float theta = 3.14159/6;
int W ;
int H ;
const char* inputFile = "input.bmp";
const char* outputFile = "output.bmp";
// Homegrown function to read a BMP from file
float* ip = readImage(inputFile, &W, &H);
float * op = new float[W*H];
//Lets choose the first platform
cl_context_properties cps[3] = {
CL_CONTEXT_PLATFORM, (cl_context_properties)(platforms[PLATFORM_TO_USE])(), 0};
// Select the default platform and create a context
// using this platform for a GPU type device
cl::Context context(DEVICE_TYPE_TO_USE, cps);
cl::vector<cl::Device> devices = context.getInfo<CL_CONTEXT_DEVICES>();
//Lets create a command queue on the first device
cl::CommandQueue queue = cl::CommandQueue(context, devices[0], 0, &err);
//[H3] Step2 – Declare Buffers and Move Data
//We assume that the input image is the array “ip”
//and the angle of rotation is theta
float cos_theta = cos(theta);
float sin_theta = sin(theta);
cl::Buffer d_ip = cl::Buffer(context, CL_MEM_READ_ONLY, W*H* sizeof(float));
cl::Buffer d_op = cl::Buffer(context, CL_MEM_READ_WRITE, W*H* sizeof(float));
queue.enqueueWriteBuffer(d_ip, CL_TRUE, 0, W*H* sizeof(float), ip);
//[H3]Step3 – Runtime kernel compilation
std::ifstream sourceFileName("rotation.cl");
std::string sourceFile(
std::istreambuf_iterator<char>( sourceFileName),
(std::istreambuf_iterator<char>())
);
//std::cout<<sourceFile;
cl::Program::Sources rotn_source(1,
std::make_pair(sourceFile.c_str(),
sourceFile.length() +1
)
);
cl::Program rotn_program(context, rotn_source);
rotn_program.build(devices);
cl::Kernel rotn_kernel(rotn_program, "img_rotate", &err);
//[H3]Step4 – Run the program
rotn_kernel.setArg(0, d_op);
rotn_kernel.setArg(1, d_ip);
rotn_kernel.setArg(2, W);
rotn_kernel.setArg(3, H);
rotn_kernel.setArg(4, sin_theta);
rotn_kernel.setArg(5, cos_theta);
// Run the kernel on specific ND range
cl::NDRange globalws(W,H);
//In this example the local work group size is not important because
//there is no communication between local work items
queue.enqueueNDRangeKernel(rotn_kernel, cl::NullRange, globalws, cl::NullRange);
//[H3]Step5 – Read result back to host
// Read buffer d_op into a local op array
queue.enqueueReadBuffer(d_op, CL_TRUE, 0, W*H*sizeof(float), op);
storeImage(op, outputFile, H, W, inputFile);
}
catch(cl::Error err)
{
std::cout << err.what() << "(" << err.err() << ")" << std::endl;
}
}
void PHISHtmlGeneric::rolloverButton() const
{
QByteArray postfix, url1, url2;
QTransform t=transformation();
if ( !_it->value().isEmpty() ) postfix=QByteArray::fromRawData( "_v", 2 );
if ( _it->hasGraphicEffect() ) postfix=QByteArray::fromRawData( "_g", 2 );
if ( !t.isIdentity() ) postfix=QByteArray::fromRawData( "_t", 2 );
QString imageId, first, second;
for ( int i=0; i<2; i++ ) {
if ( _it->pictureBookIds().count() ) { // rollover with images
imageId=checkForImageId( postfix, i );
if ( imageId.isNull() ) {
if ( i<_it->pictureBookIds().count() ) {
imageId=_it->pictureBookIds().at( i );
QImage img=createImageImage( imageId );
if ( !_it->value().isEmpty() ) img=createRolloverTextImage( i, img );
if ( _it->hasGraphicEffect() ) img=graphicsEffectImage( img );
if ( !t.isIdentity() ) img=img.transformed( t, Qt::SmoothTransformation );
imageId=storeImage( img, postfix, i );
if ( imageId.isNull() ) return;
} else {
Q_ASSERT( i==1 );
if ( _it->itemProperties() & PHIItem::PRolloverColors ) {
QImage img=createRolloverTextImage( i );
if ( _it->hasGraphicEffect() ) img=graphicsEffectImage( img );
if ( !t.isIdentity() ) img=img.transformed( t, Qt::SmoothTransformation );
imageId=storeImage( img, postfix, i );
if ( imageId.isNull() ) return;
}
}
}
} else { // rollover with text only
imageId=checkForImageId( postfix, i );
if ( imageId.isNull() ) {
if ( i==0 ) {
QImage img=createRolloverTextImage( i );
if ( _it->hasGraphicEffect() ) img=graphicsEffectImage( img );
if ( !t.isIdentity() ) img=img.transformed( t, Qt::SmoothTransformation );
imageId=storeImage( img, postfix, i );
if ( imageId.isNull() ) return;
} else {
if ( _it->itemProperties() & PHIItem::PRolloverColors ) {
QImage img=createRolloverTextImage( i );
if ( _it->hasGraphicEffect() ) img=graphicsEffectImage( img );
if ( !t.isIdentity() ) img=img.transformed( t, Qt::SmoothTransformation );
imageId=storeImage( img, postfix, i );
if ( imageId.isNull() ) return;
}
}
}
}
if ( i==0 ) first=imageId;
else second=imageId;
}
if ( first.startsWith( QLatin1String( "phi" ) ) || _it->itemProperties() & PHIItem::PNoCache ) {
url1="/phi.phis?phiimg="+first.toUtf8()+"&phitmp=1";
if ( !second.isNull() ) url2="/phi.phis?phiimg="
+second.toUtf8()+"&phitmp=1";
} else {
url1="/phi.phis?phiimg="+first.toUtf8();
if ( !second.isNull() ) url2="/phi.phis?phiimg="
+second.toUtf8();
}
QByteArray style=" style=\""+transformationPositionStyle( t, t.isIdentity() )+visibilityStyle();
return imageSource( url1, style, _it->id(), _it->toolTipData(), url2 );
}
int main(int argc, char** argv) {
// Set up the data on the host
clock_t start, start0;
start0 = clock();
start = clock();
// Rows and columns in the input image
int imageHeight;
int imageWidth;
const char* inputFile = "input.bmp";
const char* outputFile = "output.bmp";
// Homegrown function to read a BMP from file
float* inputImage = readImage(inputFile, &imageWidth,
&imageHeight);
// Size of the input and output images on the host
int dataSize = imageHeight*imageWidth*sizeof(float);
// Pad the number of columns
#ifdef NON_OPTIMIZED
int deviceWidth = imageWidth;
#else // READ_ALIGNED || READ4
int deviceWidth = roundUp(imageWidth, WGX);
#endif
int deviceHeight = imageHeight;
// Size of the input and output images on the device
int deviceDataSize = imageHeight*deviceWidth*sizeof(float);
// Output image on the host
float* outputImage = NULL;
outputImage = (float*)malloc(dataSize);
int i, j;
for(i = 0; i < imageHeight; i++) {
for(j = 0; j < imageWidth; j++) {
outputImage[i*imageWidth+j] = 0;
}
}
// 45 degree motion blur
float filter[49] =
{0, 0, 0, 0, 0, 0.0145, 0,
0, 0, 0, 0, 0.0376, 0.1283, 0.0145,
0, 0, 0, 0.0376, 0.1283, 0.0376, 0,
0, 0, 0.0376, 0.1283, 0.0376, 0, 0,
0, 0.0376, 0.1283, 0.0376, 0, 0, 0,
0.0145, 0.1283, 0.0376, 0, 0, 0, 0,
0, 0.0145, 0, 0, 0, 0, 0};
int filterWidth = 7;
int paddingPixels = (int)(filterWidth/2) * 2;
stoptime(start, "set up input, output.");
start = clock();
// Set up the OpenCL environment
// Discovery platform
cl_platform_id platform;
clGetPlatformIDs(1, &platform, NULL);
// Discover device
cl_device_id device;
clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 1, &device,
NULL);
size_t time_res;
clGetDeviceInfo(device, CL_DEVICE_PROFILING_TIMER_RESOLUTION,
sizeof(time_res), &time_res, NULL);
printf("Device profiling timer resolution: %zu ns.\n", time_res);
// Create context
cl_context_properties props[3] = {CL_CONTEXT_PLATFORM,
(cl_context_properties)(platform), 0};
cl_context context;
context = clCreateContext(props, 1, &device, NULL, NULL,
NULL);
// Create command queue
cl_ulong time_start, time_end, exec_time;
cl_event timing_event;
cl_command_queue queue;
queue = clCreateCommandQueue(context, device, CL_QUEUE_PROFILING_ENABLE, NULL);
// Create memory buffers
cl_mem d_inputImage;
cl_mem d_outputImage;
cl_mem d_filter;
d_inputImage = clCreateBuffer(context, CL_MEM_READ_ONLY,
deviceDataSize, NULL, NULL);
d_outputImage = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
deviceDataSize, NULL, NULL);
d_filter = clCreateBuffer(context, CL_MEM_READ_ONLY,
49*sizeof(float),NULL, NULL);
// Write input data to the device
#ifdef NON_OPTIMIZED
clEnqueueWriteBuffer(queue, d_inputImage, CL_TRUE, 0, deviceDataSize,
inputImage, 0, NULL, NULL);
#else // READ_ALIGNED || READ4
size_t buffer_origin[3] = {0,0,0};
size_t host_origin[3] = {0,0,0};
size_t region[3] = {deviceWidth*sizeof(float),
imageHeight, 1};
clEnqueueWriteBufferRect(queue, d_inputImage, CL_TRUE,
buffer_origin, host_origin, region,
deviceWidth*sizeof(float), 0, imageWidth*sizeof(float), 0,
inputImage, 0, NULL, NULL);
#endif
// Write the filter to the device
clEnqueueWriteBuffer(queue, d_filter, CL_TRUE, 0,
49*sizeof(float), filter, 0, NULL, NULL);
// Read in the program from file
char* source = readSource("convolution.cl");
// Create the program
cl_program program;
// Create and compile the program
program = clCreateProgramWithSource(context, 1,
(const char**)&source, NULL, NULL);
cl_int build_status;
build_status = clBuildProgram(program, 1, &device, NULL, NULL,
NULL);
// Create the kernel
cl_kernel kernel;
#if defined NON_OPTIMIZED || defined READ_ALIGNED
// Only the host-side code differs for the aligned reads
kernel = clCreateKernel(program, "convolution", NULL);
#else // READ4
kernel = clCreateKernel(program, "convolution_read4", NULL);
#endif
// Selected work group size is 16x16
int wgWidth = WGX;
int wgHeight = WGY;
// When computing the total number of work items, the
// padding work items do not need to be considered
int totalWorkItemsX = roundUp(imageWidth-paddingPixels,
wgWidth);
int totalWorkItemsY = roundUp(imageHeight-paddingPixels,
wgHeight);
// Size of a work group
size_t localSize[2] = {wgWidth, wgHeight};
// Size of the NDRange
size_t globalSize[2] = {totalWorkItemsX, totalWorkItemsY};
// The amount of local data that is cached is the size of the
// work groups plus the padding pixels
#if defined NON_OPTIMIZED || defined READ_ALIGNED
int localWidth = localSize[0] + paddingPixels;
#else // READ4
// Round the local width up to 4 for the read4 kernel
int localWidth = roundUp(localSize[0]+paddingPixels, 4);
#endif
int localHeight = localSize[1] + paddingPixels;
// Compute the size of local memory (needed for dynamic
// allocation)
size_t localMemSize = (localWidth * localHeight *
sizeof(float));
// Set the kernel arguments
clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_inputImage);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_outputImage);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &d_filter);
clSetKernelArg(kernel, 3, sizeof(int), &deviceHeight);
clSetKernelArg(kernel, 4, sizeof(int), &deviceWidth);
clSetKernelArg(kernel, 5, sizeof(int), &filterWidth);
clSetKernelArg(kernel, 6, localMemSize, NULL);
clSetKernelArg(kernel, 7, sizeof(int), &localHeight);
clSetKernelArg(kernel, 8, sizeof(int), &localWidth);
stoptime(start, "set up kernel");
start = clock();
// Execute the kernel
clEnqueueNDRangeKernel(queue, kernel, 2, NULL, globalSize,
localSize, 0, NULL, &timing_event);
// Wait for kernel to complete
clFinish(queue);
stoptime(start, "run kernel");
clGetEventProfilingInfo(timing_event, CL_PROFILING_COMMAND_START,
sizeof(time_start), &time_start, NULL);
clGetEventProfilingInfo(timing_event, CL_PROFILING_COMMAND_END,
sizeof(time_end), &time_end, NULL);
exec_time = time_end-time_start;
printf("Profile execution time = %.3lf sec.\n", (double) exec_time/1000000000);
// Read back the output image
#ifdef NON_OPTIMIZED
clEnqueueReadBuffer(queue, d_outputImage, CL_TRUE, 0,
deviceDataSize, outputImage, 0, NULL, NULL);
#else // READ_ALIGNED || READ4
// Begin reading output from (3,3) on the device
// (for 7x7 filter with radius 3)
buffer_origin[0] = 3*sizeof(float);
buffer_origin[1] = 3;
buffer_origin[2] = 0;
// Read data into (3,3) on the host
host_origin[0] = 3*sizeof(float);
host_origin[1] = 3;
host_origin[2] = 0;
// Region is image size minus padding pixels
region[0] = (imageWidth-paddingPixels)*sizeof(float);
region[1] = (imageHeight-paddingPixels);
region[2] = 1;
// Perform the read
clEnqueueReadBufferRect(queue, d_outputImage, CL_TRUE,
buffer_origin, host_origin, region,
deviceWidth*sizeof(float), 0, imageWidth*sizeof(float), 0,
outputImage, 0, NULL, NULL);
#endif
// Homegrown function to write the image to file
storeImage(outputImage, outputFile, imageHeight,
imageWidth, inputFile);
// Free OpenCL objects
clReleaseMemObject(d_inputImage);
clReleaseMemObject(d_outputImage);
clReleaseMemObject(d_filter);
clReleaseKernel(kernel);
clReleaseProgram(program);
clReleaseCommandQueue(queue);
clReleaseContext(context);
return 0;
}
int main() {
// Set the image rotation (in degrees)
float theta = 3.14159/6;
float cos_theta = cosf(theta);
float sin_theta = sinf(theta);
printf("theta = %f (cos theta = %f, sin theta = %f)\n", theta, cos_theta,
sin_theta);
// Rows and columns in the input image
int imageHeight;
int imageWidth;
const char* inputFile = "input.bmp";
const char* outputFile = "output.bmp";
// Homegrown function to read a BMP from file
float* inputImage = readImage(inputFile, &imageWidth,
&imageHeight);
// Size of the input and output images on the host
int dataSize = imageHeight*imageWidth*sizeof(float);
// Output image on the host
float* outputImage = NULL;
outputImage = (float*)malloc(dataSize);
// Set up the OpenCL environment
cl_int status;
// Discovery platform
cl_platform_id platforms[2];
cl_platform_id platform;
status = clGetPlatformIDs(2, platforms, NULL);
chk(status, "clGetPlatformIDs");
platform = platforms[PLATFORM_TO_USE];
// Discover device
cl_device_id device;
clGetDeviceIDs(platform, CL_DEVICE_TYPE_CPU, 1, &device, NULL);
chk(status, "clGetDeviceIDs");
// Create context
cl_context_properties props[3] = {CL_CONTEXT_PLATFORM,
(cl_context_properties)(platform), 0};
cl_context context;
context = clCreateContext(props, 1, &device, NULL, NULL, &status);
chk(status, "clCreateContext");
// Create command queue
cl_command_queue queue;
queue = clCreateCommandQueue(context, device, 0, &status);
chk(status, "clCreateCommandQueue");
// Create the input and output buffers
cl_mem d_input;
d_input = clCreateBuffer(context, CL_MEM_READ_ONLY, dataSize, NULL,
&status);
chk(status, "clCreateBuffer");
cl_mem d_output;
d_output = clCreateBuffer(context, CL_MEM_WRITE_ONLY, dataSize, NULL,
&status);
chk(status, "clCreateBuffer");
// Copy the input image to the device
status = clEnqueueWriteBuffer(queue, d_input, CL_TRUE, 0, dataSize,
inputImage, 0, NULL, NULL);
chk(status, "clEnqueueWriteBuffer");
const char* source = readSource("rotation.cl");
// Create a program object with source and build it
cl_program program;
program = clCreateProgramWithSource(context, 1, &source, NULL, NULL);
chk(status, "clCreateProgramWithSource");
status = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
chk(status, "clBuildProgram");
// Create the kernel object
cl_kernel kernel;
kernel = clCreateKernel(program, "img_rotate", &status);
chk(status, "clCreateKernel");
// Set the kernel arguments
status = clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_output);
status |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_input);
status |= clSetKernelArg(kernel, 2, sizeof(int), &imageWidth);
status |= clSetKernelArg(kernel, 3, sizeof(int), &imageHeight);
status |= clSetKernelArg(kernel, 4, sizeof(float), &sin_theta);
status |= clSetKernelArg(kernel, 5, sizeof(float), &cos_theta);
chk(status, "clSetKernelArg");
// Set the work item dimensions
size_t globalSize[2] = {imageWidth, imageHeight};
status = clEnqueueNDRangeKernel(queue, kernel, 2, NULL, globalSize, NULL, 0,
NULL, NULL);
chk(status, "clEnqueueNDRange");
// Read the image back to the host
status = clEnqueueReadBuffer(queue, d_output, CL_TRUE, 0, dataSize,
outputImage, 0, NULL, NULL);
chk(status, "clEnqueueReadBuffer");
// Write the output image to file
storeImage(outputImage, outputFile, imageHeight, imageWidth, inputFile);
return 0;
}
