Now have flux calculation kernels working, but still no flux from the…

… OpenGL texturebuffer.

src/CLibOI.cpp

@@ -68,6 +68,7 @@ void CLibOI::InitRoutines()
 	mImage_flux->Init(mImageWidth * mImageHeight, true);
 
 	mrCopyImage = new CRoutine_ImageToBuffer(mOCL->GetDevice(), mOCL->GetContext(), mOCL->GetQueue());
+	mrCopyImage->SetSourcePath(mKernelSourcePath);
 	mrCopyImage->Init();
 }
 
@@ -80,7 +81,7 @@ float CLibOI::TotalFlux(bool return_value)
 	{
 		// Wait for the OpenGL queue to finish.
 		glFinish();
-		err |= clEnqueueAcquireGLObjects (mOCL->GetQueue(), 1, &mGLImage, 0, NULL, NULL);
+		err |= clEnqueueAcquireGLObjects(mOCL->GetQueue(), 1, &mGLImage, 0, NULL, NULL);
 		COpenCL::CheckOCLError("Could not acquire OpenGL Memory object.", err);
 
 		// TODO: Implement depth channel for 3D images
@@ -121,6 +122,7 @@ void   CLibOI::RegisterImage_CLMEM(cl_mem image)
 /// Registers it as the currentt image object against which liboi operations will be undertaken.
 void CLibOI::RegisterImage_GLFB(GLuint framebuffer)
 {
+	this->mImageType = OpenGLBuffer;
 	int err = CL_SUCCESS;
 	mGLImage = clCreateFromGLBuffer(mOCL->GetContext(), CL_MEM_READ_ONLY, framebuffer, &err);
 	COpenCL::CheckOCLError("Could not create OpenCL image object from framebuffer", err);
@@ -134,7 +136,7 @@ void CLibOI::RegisterImage_GLFB(GLuint framebuffer)
 void CLibOI::RegisterImage_GLTB(GLuint texturebuffer)
 {
 	// TODO: Permit loading of 3D textures for spectral imaging.
-
+	this->mImageType = OpenGLBuffer;
 	int err = CL_SUCCESS;
 
 	// TODO: note that the clCreateFromGLTexture2D was depreciated in the OpenCL 1.2 specifications.

src/COpenCLRoutine.cpp

@@ -39,7 +39,7 @@ COpenCLRoutine::~COpenCLRoutine()
 
 /// Builds the kernel from the specified string.
 /// Appends the compiled kernel to mKernels and mPrograms, returns the index at which this kernel is located.
-int COpenCLRoutine::BuildKernel(string source)
+int COpenCLRoutine::BuildKernel(string source, string kernel_name)
 {
     const char * tmp = source.c_str();
     cl_program program;
@@ -86,8 +86,8 @@ int COpenCLRoutine::BuildKernel(string source)
     mPrograms.push_back(program);
 
     // Create the compute kernel from within the program
-    kernel = clCreateKernel(program, "reduce", &err);
-	COpenCL::CheckOCLError("Failed to create parallel sum kernel.", err);
+    kernel = clCreateKernel(program, kernel_name.c_str(), &err);
+	COpenCL::CheckOCLError("Failed to create kernel named " + kernel_name, err);
 
 	// All is well, push the kernel onto the vector
 	mKernels.push_back(kernel);

src/COpenCLRoutine.h

@@ -31,7 +31,7 @@ class COpenCLRoutine
 	COpenCLRoutine(cl_device_id mDevice, cl_context mContext, cl_command_queue mQueue);
 	~COpenCLRoutine();
 
-	int BuildKernel(string source);
+	int BuildKernel(string source, string kernel_name);
 
 	string ReadSource(string filename);
 

src/CRoutine_ImageToBuffer.cpp

@@ -6,6 +6,7 @@
  */
 
 #include "CRoutine_ImageToBuffer.h"
+#include <cstdio>
 
 CRoutine_ImageToBuffer::CRoutine_ImageToBuffer(cl_device_id device, cl_context context, cl_command_queue queue)
 	:COpenCLRoutine(device, context, queue)
@@ -22,17 +23,17 @@ CRoutine_ImageToBuffer::~CRoutine_ImageToBuffer()
 // Read in the kernel source and build program object.
 void CRoutine_ImageToBuffer::Init()
 {
+#ifdef DEBUG
+    string message = "Loading and Compiling program " + mSource[0] + "\n";
+	printf("%s\n", message.c_str());
+#endif //DEBUG
+
 	string source = ReadSource(mSource[0]);
-	BuildKernel(source);
+	BuildKernel(source, "image2buf_GL_R");
 }
 
 void CRoutine_ImageToBuffer::CopyImage(cl_mem gl_image, cl_mem cl_buffer, int width, int height, int depth)
 {
-	// Create the image size memory object
-	cl_int2 tmp;
-	tmp.x = width;
-	tmp.y = height;
-
 	// TODO: We need to redo this for 3D data sets and for non-square images.
 	// TODO: Figure out how to determine these sizes dynamically.
 	size_t * global = new size_t[2];
@@ -44,14 +45,33 @@ void CRoutine_ImageToBuffer::CopyImage(cl_mem gl_image, cl_mem cl_buffer, int wi
 	int err = CL_SUCCESS;
     err |= clSetKernelArg(mKernels[0],  0, sizeof(cl_mem), &gl_image);
     err |= clSetKernelArg(mKernels[0],  1, sizeof(cl_mem), &cl_buffer);
-    err |= clSetKernelArg(mKernels[0],  2, sizeof(cl_int2), &tmp);
+    err |= clSetKernelArg(mKernels[0],  2, sizeof(int), &width);
 	COpenCL::CheckOCLError("Failed to set gl_image to cl_buffer kernel arguments.", err);
 
 
     err = CL_SUCCESS;
     err |= clEnqueueNDRangeKernel(mQueue, mKernels[0], 2, NULL, global, local, 0, NULL, NULL);
     COpenCL::CheckOCLError("Failed to enqueue image copying kernel.", err);
 
+#ifdef DEBUG
+        	// Copy back the input/output buffers.
+        	float tmp_sum = 0;
+        	int num_elements = width * height;
+        	cl_float * tmp = new cl_float[num_elements];
+        	err = clEnqueueReadBuffer(mQueue, cl_buffer, CL_TRUE, 0, num_elements * sizeof(cl_float), tmp, 0, NULL, NULL);
+
+        	for(int i = 0; i < num_elements; i++)
+        	{
+        		if(i % 10 == 0)
+        			printf("%f ", tmp[i]);
+
+        		tmp_sum += tmp[i];
+        	}
+
+        	printf("\n");
+        	printf("Sum of Image copied to CPU: %f \n", tmp_sum);
+#endif //DEBUG
+
     // Free memory
     delete global;
     delete local;

src/CRoutine_Normalize.cpp

@@ -23,7 +23,7 @@ CRoutine_Normalize::~CRoutine_Normalize()
 void CRoutine_Normalize::Init()
 {
 	string source = ReadSource(mSource[0]);
-	BuildKernel(source);
+	BuildKernel(source, "normalize");
 }
 
 void CRoutine_Normalize::Normalize(cl_mem image, cl_mem divisor, int width, int height, int depth)

src/CRoutine_Reduce.cpp

@@ -128,7 +128,7 @@ void CRoutine_Reduce::BuildKernels()
         tmp << "#define GROUP_SIZE " << group_counts[i] << "\n" << "#define OPERATIONS " << operation_counts[i] << "\n\n";
         tmp << source;
 
-        BuildKernel(tmp.str());
+        BuildKernel(tmp.str(), "reduce");
     }
 
 }
@@ -215,24 +215,6 @@ float CRoutine_Reduce::ComputeSum(bool copy_back, cl_mem final_buffer, cl_mem in
         // Check for a NaN:
         if(sum != sum)
         {
-#ifdef DEBUG
-        	// Copy back the input/output buffers.
-        	float tmp_sum = 0;
-        	cl_float * tmp = new cl_float[num_elements];
-        	err = clEnqueueReadBuffer(mQueue, input_buffer, CL_TRUE, 0, num_elements * sizeof(cl_float), tmp, 0, NULL, NULL);
-
-        	for(int i = 0; i < num_elements; i++)
-        	{
-        		if(i % 10 == 0)
-        			printf("%f ", tmp[i]);
-
-        		tmp_sum += tmp[i];
-        	}
-
-        	printf("\n");
-        	printf("Sum of Image copied to CPU: %f \n", tmp_sum);
-#endif //DEBUG
-
         	COpenCL::CheckOCLError("Error: Calculation yielded NAN.", 1);
         }
 

src/kernels/image2buf_GL_R.cl

@@ -4,13 +4,17 @@
 
 const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
 
-__kernel void normalize(__read_only  image2d_t image,
-    __global float * buffer
-    __global int image_width)
+__kernel void image2buf_GL_R(__read_only image2d_t image,
+    __global float * buffer,
+    __private int image_width)
 {
     int x = get_global_id(0);
     int y = get_global_id(1);
 
-    float4 colors = readimagef(image, sampler, float2 {x, y});
-    buffer[y * image_width + x] = colors[0];
+    float2 coords;
+    coords.x = x;
+    coords.y = y;
+
+    float4 colors = read_imagef(image, sampler, coords);
+    buffer[y * image_width + x] = colors.s0;
 }