Switched names, thinking about how to do kernel args, and buffers.

Need to do a bit more reading on how to set up interop.
Also need to figure out the buffer thing for regular primative buffers, and also image buffers

include/CL_Wrapper.h

@@ -27,18 +27,20 @@ struct device {
     cl_platform_id platform;
 };
 
-class Clapper {
+class CL_Wrapper {
 public:
 
-	enum PLATFORM_TYPE { GPU, CPU, ALL };
+	CL_Wrapper();
-
+	~CL_Wrapper();
-	Clapper();
-	~Clapper();
 
     int acquire_platform_and_device();
     int create_shared_context();
     int create_command_queue();
     int compile_kernel(std::string kernel_source, bool is_path, std::string kernel_name);
+	int set_kernel_arg(cl_kernel kernel, int index, int size, void* buffer, std::string kernel_name);
+	int store_buffer(cl_mem, std::string buffer_name);
+	int run_kernel(std::string kernel_name);
+
 
     cl_device_id getDeviceID();
     cl_platform_id getPlatformID();
@@ -53,7 +55,9 @@ private:
     cl_platform_id platform_id;
 	cl_device_id device_id;
     cl_context context;
+    cl_command_queue command_queue;
 
     std::map<std::string, cl_kernel> kernel_map;
+    std::map<std::string, cl_mem> buffer_map;
 };
 

kernels/kernel.txt

@@ -1,17 +1,125 @@
-__kernel void hello(__global char* string)
+// global  :  local  :  constant  :  private
-{
+
-string[0] = 'H';
+// Function arguments of type image2d_t, image3d_t, image2d_array_t, image1d_t, image1d_buffer_t,
-string[1] = 'e';
+// and image1d_array_t refer to image memory objects allocated in the **global** address space.
-string[2] = 'l';
+
-string[3] = 'l';
+// http://downloads.ti.com/mctools/esd/docs/opencl/memory/buffers.html
-string[4] = 'o';
+
-string[5] = ',';
+// Open CL C
-string[6] = ' ';
+// https://www.fixstars.com/en/opencl/book/OpenCLProgrammingBook/opencl-c/
-string[7] = 'W';
+
-string[8] = 'o';
+__kernel void hello(
-string[9] = 'r';
+    global int2* resolution,
-string[10] = 'l';
+    global char* map,
-string[11] = 'd';
+    global float3* projection_matrix,
-string[12] = '!';
+    global float3* cam_dir,
-string[13] = '\0';
+    global float3* cam_pos,
+    global image2d_t* canvas) {
+
+    printf("%s\n", "this is a test string\n");
+
+
+
+    const int MAX_RAY_STEPS = 64;
+
+    // The pixel coord we are at
+    int2 screenPos = (int2)(get_global_id(0) % resolution->x, get_global_id(0) / resolution->x);
+
+    // The X and Y planes
+    //float3 cameraPlaneU = vec3(1.0, 0.0, 0.0)
+
+    // Y being multiplied by the aspect ratio, usually around .5-6ish;
+    //cl_float3 cameraPlaneV = vec3(0.0, 1.0, 0.0) * iResolution.y / iResolution.x;
+
+    // So this is how they do that ray aiming! hah this is so tiny
+    // (camera direction) + (pixel.x * the X plane) + (product of pixel.y * Y plane)
+    // Oh all it's doing is adding the x and y coords of the pixel to the camera direction vector, interesting
+
+    //cl_float3 rayDir = cameraDir + screenPos.x * cameraPlaneU + screenPos.y * cameraPlaneV;
+
+    // the origin of the ray
+    // So the sign thing is for the up and down motion
+
+    //cl_float3 rayPos = vec3(0.0, 2.0 * sin(iGlobalTime * 2.7), -12.0);
+
+    // Ah, and here is where it spins around the center axis
+    // So it looks like its applying a function to rotate the x and z axis
+    //rayPos.xz = rotate2d(rayPos.xz, iGlobalTime);
+    //rayDir.xz = rotate2d(rayDir.xz, iGlobalTime);
+
+    // Just an intvec of out coords
+    //ivec3 mapPos = ivec3(floor(rayPos));
+
+    // I think this is the delta t value
+    // the magnitude of the vector divided by the rays direction. Not sure what the aim of that is
+    // The ray direction might always be normalized, so that would be the dame as my delta_T
+    //vec3 deltaDist = abs(vec3(length(rayDir)) / rayDir);
+
+    // The steps are the signs of the ray direction
+    //ivec3 rayStep = ivec3(sign(rayDir));
+
+    //  ithe sign of the rays direction
+    //  *
+    //	Convert map position to a floating point vector and take away the ray position
+    //  +
+    //  the sign of the rays direction by 0.5
+    //  +
+    //  0.5
+    //  Now multyply everything by 0.5
+    //vec3 sideDist = (sign(rayDir) * (vec3(mapPos) - rayPos) + (sign(rayDir) * 0.5) + 0.5) * deltaDist;
+
+    // A byte mask
+    //bvec3 mask;
+
+    // repeat until the max steps
+    //for (int i = 0; i < MAX_RAY_STEPS; i++) {
+
+            // If there is a voxel at the map position, continue?
+            //if (getVoxel(mapPos))
+            //        break;
+
+            //
+            // find which is smaller
+            // y ? z --> x`
+            // z ? x --> y`
+            // x ? y --> z`
+            //
+            // find which os is less or equal
+            // x` ? x --> x
+            // y` ? y --> y
+            // z` ? z --> z
+
+            // Now find which ons is
+            //mask = lessThanEqual(sideDist.xyz, min(sideDist.yzx, sideDist.zxy));
+
+
+            // Originally he used a component wise
+            /*bvec3 b1 = lessThan(sideDist.xyz, sideDist.yzx);
+            bvec3 b2 = lessThanEqual(sideDist.xyz, sideDist.zxy);
+            mask.x = b1.x && b2.x;
+            mask.y = b1.y && b2.y;
+            mask.z = b1.z && b2.z;*/
+            //Would've done mask = b1 && b2 but the compiler is making me do it component wise.
+
+            //All components of mask are false except for the corresponding largest component
+            //of sideDist, which is the axis along which the ray should be incremented.
+
+            //sideDist += vec3(mask) * deltaDist;
+            //mapPos += ivec3(mask) * rayStep;
+    //}
+
+    // Ah this is for coloring obviously, seems to be odd though, no indexing
+    //vec4 color;
+    //if (mask.x) {
+    //       color = vec4(0.5);
+    //}
+    //if (mask.y) {
+    //       color = vec4(1.0);
+    //}
+    //if (mask.z) {
+    //        color = vec4(0.75);
+    //}
+    //write_imagef(image, pixel, color);
+
+
 }

src/CL_Wrapper.cpp

@@ -1,13 +1,13 @@
-#include "Clapper.h"
+#include "CL_Wrapper.h"
 
-Clapper::Clapper() {
+CL_Wrapper::CL_Wrapper() {
 }
 
 
-Clapper::~Clapper() {
+CL_Wrapper::~CL_Wrapper() {
 }
 
-int Clapper::acquire_platform_and_device(){
+int CL_Wrapper::acquire_platform_and_device(){
 
     // Get the number of platforms
     cl_uint plt_cnt = 0;
@@ -84,7 +84,7 @@ int Clapper::acquire_platform_and_device(){
     return 0;
 };
 
-int Clapper::create_shared_context() {
+int CL_Wrapper::create_shared_context() {
 
     // Hurray for standards!
     // Setup the context properties to grab the current GL context
@@ -137,19 +137,24 @@ int Clapper::create_shared_context() {
     return 0;
 }
 
-int Clapper::create_command_queue(){
+int CL_Wrapper::create_command_queue(){
 
-    // / And the cl command queue
+    if (context && device_id) {
-    auto commandQueue = clCreateCommandQueue(context, device_id, 0, &error);
+        // And the cl command queue
+        command_queue = clCreateCommandQueue(context, device_id, 0, &error);
 
-    if (assert(error, "clCreateCommandQueue"))
+        if (assert(error, "clCreateCommandQueue"))
-        return -1;
+            return -1;
-
-    return 0;
 
+        return 0;
+    }
+    else {
+        std::cout << "Failed creating the command queue, context or device_id not initialized";
+        return -1;
+    }
 }
 
-int Clapper::compile_kernel(std::string kernel_source, bool is_path, std::string kernel_name) {
+int CL_Wrapper::compile_kernel(std::string kernel_source, bool is_path, std::string kernel_name) {
 
     const char* source;
 
@@ -202,11 +207,47 @@ int Clapper::compile_kernel(std::string kernel_source, bool is_path, std::string
     kernel_map.emplace(std::make_pair(kernel_name, kernel));
 }
 
-cl_device_id Clapper::getDeviceID(){ return device_id; };
+int CL_Wrapper::set_kernel_arg(
-cl_platform_id Clapper::getPlatformID(){ return platform_id; };
+        cl_kernel kernel,
-cl_context Clapper::getContext(){ return context; };
+        int index,
+        int size,
+        void* buffer,
+        std::string kernel_name){
+
+
+
+
+
+};
+int CL_Wrapper::store_buffer(cl_mem, std::string buffer_name){
+    buffer_map.emplace(std::make_pair(buffer_name, cl_mem));
+};
+
+int CL_Wrapper::run_kernel(std::string kernel_name){
+
+    WORKER_SIZE = 10;
+    size_t global_work_size[1] = { WORKER_SIZE };
+
+    cl_mem kernel = kernel_map.at(kernel_name);
+
+    error = clEnqueueNDRangeKernel(
+            command_queue, kernel,
+            1, NULL, global_work_size,
+            NULL, 0, NULL, NULL);
+
+    if (assert(error, "clCreateCommandQueue"))
+        return -1;
+
+
+};
+
+
+
+cl_device_id CL_Wrapper::getDeviceID(){ return device_id; };
+cl_platform_id CL_Wrapper::getPlatformID(){ return platform_id; };
+cl_context CL_Wrapper::getContext(){ return context; };
 
-bool Clapper::assert(int error_code, std::string function_name){
+bool CL_Wrapper::assert(int error_code, std::string function_name){
 
     // Just gonna do a little jump table here, just error codes so who cares
     std::string err_msg = "Error : ";

src/main.cpp

@@ -27,7 +27,7 @@
 #include "Curses.h"
 #include "util.hpp"
 #include "RayCaster.h"
-#include "Clapper.h"
+#include "CL_Wrapper.h"
 
 const int WINDOW_X = 150;
 const int WINDOW_Y = 150;
@@ -36,7 +36,7 @@ const int WINDOW_Y = 150;
 
 int main(){
 
-    Clapper c;
+    CL_Wrapper c;
     c.acquire_platform_and_device();
     c.create_shared_context();