diff --git a/src/main/java/com/rayx/RayX.java b/src/main/java/com/rayx/RayX.java index a605905..58fd7a4 100644 --- a/src/main/java/com/rayx/RayX.java +++ b/src/main/java/com/rayx/RayX.java @@ -13,7 +13,6 @@ public class RayX { static CLContext context; public static void main(String[] args) { - WindowManager manager = WindowManager.getInstance(); TestOGLWindow window = new TestOGLWindow(800, 600, "Test"); @@ -38,7 +37,7 @@ public static void main(String[] args) { window.setCallback((texture) -> { glFinish(); - CLContext.CLKernel kernel = context.getKernelObject("test"); + CLContext.CLKernel kernel = context.getKernelObject("testKernel"); CLManager.allocateMemory(context, CL_MEM_READ_ONLY, new double[]{-1.5, 0.0, .5}, "center"); CLManager.createCLFromGLTexture(context, CL_MEM_WRITE_ONLY, GL_TEXTURE_2D, texture, "tex"); context.getMemoryObject("tex").acquireFromGL(); @@ -55,55 +54,6 @@ public static void main(String[] args) { freeAll(); } - static final String KERNEL_SOURCE = """ - float4 color(int, int); - int calc(double2, int); - - __kernel void testKernel(__write_only image2d_t image, __global double* center) { - int2 pixCo = (int2){get_global_id(0),get_global_id(1)}; - int w = get_image_width(image); - int h = get_image_height(image); - if(pixCo.x >= w | pixCo.y >= h) { - return; - } - double2 absCo = (double2) - {center[2] * ((1.0 * pixCo.x / w) - .5) + center[0], - center[2] * ((1.0 * (h - pixCo.y - 1) / h) - .5) + center[1]}; - - write_imagef(image, pixCo, color(calc(absCo, 100000), 100000)); - } - - int calc(double2 coo, int maxIter) { - double a = 0, b = 0, ta, tb; - int i = 0; - for(;(i < maxIter) & !(a * a + b * b > 4); i++) { - ta = a*a - b*b + coo.x; - tb = 2 * a * b + coo.y; - a = ta; - b = tb; - } - return i == maxIter ? -1 : i; - } - - float4 color(int value, int maxIterations) { - int color; - if(value == -1) { - color = 0x0; - } else if(value == 0) { - color = 0x0; - } else { - color = 0xffffff * (1.0*log((double)value)/log((double)maxIterations)); - } - - return (float4){ - 1.0 * ((color >> 0) & 0xFF) / 0xFF, - 1.0 * ((color >> 8) & 0xFF) / 0xFF, - 1.0 * ((color >> 16) & 0xFF) / 0xFF, - 1 - }; - } - """; - static CLContext treatDevice(OpenGLWindow window, long device) { long t = System.currentTimeMillis(); CLContext context = CLManager.createContext(device, window.getWindow()); @@ -111,7 +61,25 @@ static CLContext treatDevice(OpenGLWindow window, long device) { System.out.println("Context: " + (System.currentTimeMillis() - t)); t = System.currentTimeMillis(); - CLManager.addProgramAndKernelToContext(context, KERNEL_SOURCE, "testKernel", "test"); + CLManager.putProgramFromFile(context, null, + "clcode/headers/utility/mandelbrot.h"); + CLManager.putProgramFromFile(context, new String[]{ + "clcode/headers/utility/mandelbrot.h" + }, + "clcode/utility/mandelbrot.cl"); + CLManager.putProgramFromFile(context, new String[]{ + "clcode/headers/utility/mandelbrot.h" + }, + "clcode/main.cl"); + CLManager.putExecutableProgram(context, new String[] { + "clcode/utility/mandelbrot.cl", + //Can be added, does not change anything + //"clcode/headers/utility/mandelbrot.h", + "clcode/main.cl", + },"testProgram"); + + CLManager.putKernel(context, "testKernel", "testKernel", "testProgram"); + System.out.println("Kernel: " + (System.currentTimeMillis() - t)); return context; } @@ -119,7 +87,7 @@ static CLContext treatDevice(OpenGLWindow window, long device) { static void freeAll() { long t = System.currentTimeMillis(); - context.destroyKernel("test"); + context.destroyKernel("testKernel"); context.destroy(); System.out.println("Destroy: " + (System.currentTimeMillis() - t)); } diff --git a/src/main/java/com/rayx/examples/TestOGLWindow.java b/src/main/java/com/rayx/examples/TestOGLWindow.java index da72882..b8fa5d7 100644 --- a/src/main/java/com/rayx/examples/TestOGLWindow.java +++ b/src/main/java/com/rayx/examples/TestOGLWindow.java @@ -42,7 +42,6 @@ void main() } private int texture; - private long window; private int VAO, EBO, VBO; private long lastPrint = 0; private int frames; @@ -116,7 +115,6 @@ private int createShader() { int vs = compileShader(GL_VERTEX_SHADER, VS); int fs = compileShader(GL_FRAGMENT_SHADER, FS); - glAttachShader(program, vs); glAttachShader(program, fs); glLinkProgram(program); diff --git a/src/main/java/com/rayx/opencl/CLContext.java b/src/main/java/com/rayx/opencl/CLContext.java index 2062113..198917f 100644 --- a/src/main/java/com/rayx/opencl/CLContext.java +++ b/src/main/java/com/rayx/opencl/CLContext.java @@ -6,8 +6,11 @@ import java.util.HashMap; public class CLContext { - private long context, device, commandQueue; + private final long device; + private long context; + private long commandQueue; private HashMap kernels; + private HashMap programs; private HashMap memoryObjects; /** Only {@link CLManager} should create instances */ @@ -16,6 +19,7 @@ public class CLContext { this.device = device; this.commandQueue = commandQueue; kernels = new HashMap<>(); + programs = new HashMap<>(); memoryObjects = new HashMap<>(); } @@ -26,14 +30,15 @@ public void freeMemoryObject(String id) { } public void destroy() { - for(CLKernel k: kernels.values()) { - k.destroy(); - } + kernels.values().forEach(CLKernel::destroy); kernels.clear(); kernels = null; - memoryObjects.values().forEach(CLMemoryObject::delete); + programs.values().forEach(CLProgram::destroy); + programs.clear(); + programs = null; + memoryObjects.values().forEach(CLMemoryObject::delete); memoryObjects.clear(); memoryObjects = null; @@ -48,6 +53,11 @@ public void getMemoryObjectValue(String id, ByteBuffer destination) { getMemoryObject(id).getValue(destination); } + public CLProgram getProgramObject(String id) { + assert programs.containsKey(id); + return programs.get(id); + } + public CLMemoryObject getMemoryObject(String id) { assert memoryObjects.containsKey(id); return memoryObjects.get(id); @@ -70,9 +80,21 @@ public long getCommandQueue() { return commandQueue; } - public void addKernelObject(String id, CLKernel kernel) { - assert !kernels.containsKey(id); - kernels.put(id, kernel); + public void destroyProgram(String name) { + assert kernels.values().stream(). + noneMatch(k -> name.equals(k.getProgramId())) : "Kernels depend on program \"" + name + "\""; + assert kernels.containsKey(name); + kernels.remove(name).destroy(); + } + + public void addKernelObject(CLKernel kernel) { + assert !kernels.containsKey(kernel.kernelId); + kernels.put(kernel.kernelId, kernel); + } + + public void addProgramObject(CLProgram program) { + assert !programs.containsKey(program.programId); + programs.put(program.programId, program); } public CLKernel getKernelObject(String id) { @@ -91,7 +113,7 @@ public class CLMemoryObject { //In bytes //If size is -1, it indicates that the memory object should not be read to the main memory //This is the case when - private long size; + private final long size; public CLMemoryObject(long pointer, long size) { this.pointer = pointer; @@ -126,14 +148,46 @@ public long getSize() { } } - public class CLKernel { - private String name; + public class CLProgram { + private final String programId; private long program; + /** A kernel can only be created from a program if linked is set to true*/ + private final boolean linked; + + /** @param programId Name of the program. + * Should be name of the file it was generated from, + * so that it may used for include directives in OpenCL C*/ + public CLProgram(String programId, long program, boolean linked) { + this.programId = programId; + this.program = program; + this.linked = linked; + } + + public String getProgramId() { + return programId; + } + + public long getProgram() { + return program; + } + + private void destroy() { + CLManager.destroyCLProgramInternal(program); + program = 0; + } + + public boolean isLinked() { + return linked; + } + } + + public class CLKernel { + private final String kernelId, programId; private long kernel; - public CLKernel(String name, long program, long kernel) { - this.name = name; - this.program = program; + public CLKernel(String kernelId, String programId, long kernel) { + this.kernelId = kernelId; + this.programId = programId; this.kernel = kernel; } @@ -155,22 +209,20 @@ public void run(long[] globalWorkSize, long[] localWorkSize) { } private void destroy() { - CLManager.destroyCLKernelInternal(program, kernel); + CLManager.destroyCLKernelInternal(kernel); this.kernel = 0; - this.program = 0; - name = null; } public long getKernel() { return kernel; } - public long getProgram() { - return program; + public String getProgramId() { + return programId; } - public String getName() { - return name; + public String getKernelId() { + return kernelId; } } } diff --git a/src/main/java/com/rayx/opencl/CLManager.java b/src/main/java/com/rayx/opencl/CLManager.java index 245f0fb..7f3885b 100644 --- a/src/main/java/com/rayx/opencl/CLManager.java +++ b/src/main/java/com/rayx/opencl/CLManager.java @@ -1,18 +1,25 @@ package com.rayx.opencl; +import com.rayx.RayX; import org.lwjgl.PointerBuffer; import org.lwjgl.opencl.CL12GL; import org.lwjgl.opencl.CL22; +import org.lwjgl.opencl.CLProgramCallbackI; import org.lwjgl.opencl.KHRGLSharing; import org.lwjgl.opengl.CGL; import org.lwjgl.opengl.GLX14; import org.lwjgl.system.MemoryStack; import org.lwjgl.system.Platform; +import org.lwjgl.system.Pointer; +import java.io.BufferedReader; +import java.io.IOException; +import java.io.InputStreamReader; import java.nio.ByteBuffer; import java.nio.IntBuffer; import java.nio.LongBuffer; import java.util.ArrayList; +import java.util.Arrays; import java.util.List; import static org.lwjgl.glfw.GLFWNativeGLX.glfwGetGLXContext; @@ -33,6 +40,22 @@ private CLManager() {} internalMemoryStack = MemoryStack.create(); } + public static String readFromFile(String file) { + BufferedReader r = new BufferedReader(new InputStreamReader(CLManager.class.getResourceAsStream(file))); + StringBuilder b = new StringBuilder(); + String line; + try { + while((line = r.readLine()) != null) { + b.append(line); + b.append('\n'); + } + r.close(); + } catch (IOException e) { + throw new RuntimeException(e); + } + return b.toString(); + } + private static MemoryStack nextStackFrame() { return internalMemoryStack.push(); } @@ -55,23 +78,101 @@ public static CLContext createContext(long device, long glfwWindow) { } } - public static void addProgramAndKernelToContext(CLContext context, String programSourceCode, - String kernelFunction, String kernelId) { + /***/ + public static void putProgramFromFile(CLContext context, String[] dependenciesId, String sourceFile) { + try (MemoryStack stack = CLManager.nextStackFrame()) { + IntBuffer errorBuffer = stack.mallocInt(1); + //Files start without slash because the source file is the id of the program + //Because the id is also used for header includes, it cannot start with a slash + //The slash is attached internally + assert !sourceFile.startsWith("/") : "Files have to start without slash"; + String programSourceCode = readFromFile("/" + sourceFile); + long programPointer = CL22.clCreateProgramWithSource(context.getContext(), programSourceCode, errorBuffer); + checkForError(errorBuffer); + if(dependenciesId == null) { + dependenciesId = new String[0]; + } + CLContext.CLProgram[] dependencies = new CLContext.CLProgram[dependenciesId.length]; + for(int i = 0; i < dependencies.length; i++) { + dependencies[i] = context.getProgramObject(dependenciesId[i]); + } + PointerBuffer inputHeaderNames, inputHeaderPrograms; + if(dependencies.length != 0) { + inputHeaderNames = stack.pointers(Arrays.stream(dependencies). + map((CLContext.CLProgram proObj) -> stack.ASCII(proObj.getProgramId())). + toArray(ByteBuffer[]::new)); + inputHeaderPrograms = stack.pointers(Arrays.stream(dependencies). + map(CLContext.CLProgram::getProgram).mapToLong(Long::longValue).toArray()); + } else { + inputHeaderNames = null; + inputHeaderPrograms = null; + } + int error= CL22.clCompileProgram( + programPointer, + stack.pointers(context.getDevice()), "", + inputHeaderPrograms, + inputHeaderNames, + null, + 0 + ); + String buildInfo = createBuildInfo(context.getDevice(), programPointer); + System.out.println(buildInfo); + checkForError(error); + context.addProgramObject(context.new CLProgram(sourceFile, programPointer, false)); + } + } + + public static void putExecutableProgram(CLContext context, String[] programIds, String programName) { + try (MemoryStack stack = CLManager.nextStackFrame()) { + //stack.pointers(Arrays.stream(programIds).mapToLong(u -> context.getProgramObject(u).getProgram()).toArray()), + long program = CL22.clLinkProgram(context.getContext(), + stack.pointers(context.getDevice()), + "", + stack.pointers(Arrays.stream(programIds).mapToLong(u -> context.getProgramObject(u).getProgram()).toArray()), + null, + 0); + if(program == 0) { + throw new RuntimeException("CL: Linking failed"); + } + context.addProgramObject(context.new CLProgram(programName, program, true)); + } + } + + public static void putKernel(CLContext context, String kernelFunction, String kernelId, String programId) { + try (MemoryStack stack = CLManager.nextStackFrame()) { + IntBuffer error = stack.mallocInt(1); + CLContext.CLProgram program = context.getProgramObject(programId); + if(!program.isLinked()) { + throw new RuntimeException("Program is not linked"); + } + long kernel = CL22.clCreateKernel(program.getProgram(), kernelFunction, error); + checkForError(error); + + CLContext.CLKernel kernelObj = context.new CLKernel(kernelId, programId, kernel); + context.addKernelObject(kernelObj); + } + } + + public static void putKernelAndProgramFromSource(CLContext context, String programSourceCode, + String kernelFunction, String kernelId, String programId) { try (MemoryStack stack = CLManager.nextStackFrame()) { IntBuffer error = stack.mallocInt(1); - long program = CL22.clCreateProgramWithSource(context.getContext(), programSourceCode, error); + long programPointer = CL22.clCreateProgramWithSource(context.getContext(), programSourceCode, error); checkForError(error); - int programStatus = CL22.clBuildProgram(program, context.getDevice(), "", null, 0); - String buildInfo = createBuildInfo(context.getDevice(), program); + int programStatus = CL22.clBuildProgram(programPointer, context.getDevice(), "", null, 0); + String buildInfo = createBuildInfo(context.getDevice(), programPointer); System.out.println(buildInfo); checkForError(programStatus); - long kernel = CL22.clCreateKernel(program, kernelFunction, error); + long kernel = CL22.clCreateKernel(programPointer, kernelFunction, error); checkForError(error); - CLContext.CLKernel kernelObj = context.new CLKernel(kernelId, program, kernel); - context.addKernelObject(kernelId, kernelObj); + CLContext.CLProgram program = context.new CLProgram(programId, programPointer, true); + context.addProgramObject(program); + + CLContext.CLKernel kernelObj = context.new CLKernel(kernelId, programId, kernel); + context.addKernelObject(kernelObj); } } @@ -392,11 +493,14 @@ static void destroyContextInternal(CLContext context) { checkForError(CL22.clReleaseContext(context.getContext())); } - static void destroyCLKernelInternal(long program, long kernel) { + static void destroyCLProgramInternal(long program) { assert program != 0; + checkForError(CL22.clReleaseProgram(program)); + } + + static void destroyCLKernelInternal(long kernel) { assert kernel != 0; checkForError(CL22.clReleaseKernel(kernel)); - checkForError(CL22.clReleaseProgram(program)); } static void readMemoryInternal(long commandQueue, long pointer, ByteBuffer buffer) { diff --git a/src/main/java/com/rayx/shape/Sphere.java b/src/main/java/com/rayx/shape/Sphere.java new file mode 100644 index 0000000..9125e8f --- /dev/null +++ b/src/main/java/com/rayx/shape/Sphere.java @@ -0,0 +1,24 @@ +package com.rayx.shape; + +import com.rayx.opencl.CLObjectTransfer; + +import java.nio.ByteBuffer; + +public class Sphere implements CLObjectTransfer.Transferable { + private final Vector3d position; + private final double radius; + + public Sphere(Vector3d position, double radius) { + this.position = position; + this.radius = radius; + } + + @Override + public void transferTo(ByteBuffer destination) { + destination. + putDouble(position.getX()). + putDouble(position.getY()). + putDouble(position.getZ()). + putDouble(radius); + } +} diff --git a/src/main/java/com/rayx/shape/Vector3d.java b/src/main/java/com/rayx/shape/Vector3d.java new file mode 100644 index 0000000..ba82766 --- /dev/null +++ b/src/main/java/com/rayx/shape/Vector3d.java @@ -0,0 +1,23 @@ +package com.rayx.shape; + +public class Vector3d { + private final double x, y, z; + + public Vector3d(double x, double y, double z) { + this.x = x; + this.y = y; + this.z = z; + } + + public double getX() { + return x; + } + + public double getY() { + return y; + } + + public double getZ() { + return z; + } +} diff --git a/src/main/resources/clcode/headers/utility/mandelbrot.h b/src/main/resources/clcode/headers/utility/mandelbrot.h new file mode 100644 index 0000000..e94e099 --- /dev/null +++ b/src/main/resources/clcode/headers/utility/mandelbrot.h @@ -0,0 +1,7 @@ +#ifndef __HEADER_CALC_H +#define __HEADER_CALC_H + +int mandelbrotCalc(double2 coo, int maxIter); +float4 mandelbrotColor(int value, int maxIterations); + +#endif \ No newline at end of file diff --git a/src/main/resources/clcode/main.cl b/src/main/resources/clcode/main.cl new file mode 100644 index 0000000..d496678 --- /dev/null +++ b/src/main/resources/clcode/main.cl @@ -0,0 +1,18 @@ +#include +/* +float4 color(int, int); +int calc(double2, int); +*/ +__kernel void testKernel(__write_only image2d_t image, __global double* center) { + int2 pixCo = (int2){get_global_id(0),get_global_id(1)}; + int w = get_image_width(image); + int h = get_image_height(image); + if(pixCo.x >= w | pixCo.y >= h) { + return; + } + double2 absCo = (double2) + {center[2] * ((1.0 * pixCo.x / w) - .5) + center[0], + center[2] * ((1.0 * (h - pixCo.y - 1) / h) - .5) + center[1]}; + + write_imagef(image, pixCo, mandelbrotColor(mandelbrotCalc(absCo, 100), 100)); +} \ No newline at end of file diff --git a/src/main/resources/clcode/utility/mandelbrot.cl b/src/main/resources/clcode/utility/mandelbrot.cl new file mode 100644 index 0000000..98fdc87 --- /dev/null +++ b/src/main/resources/clcode/utility/mandelbrot.cl @@ -0,0 +1,31 @@ +#include + +int mandelbrotCalc(double2 coo, int maxIter) { + double a = 0, b = 0, ta, tb; + int i = 0; + for(;(i < maxIter) & !(a * a + b * b > 4); i++) { + ta = a*a - b*b + coo.x; + tb = 2 * a * b + coo.y; + a = ta; + b = tb; + } + return i == maxIter ? -1 : i; +} + +float4 mandelbrotColor(int value, int maxIterations) { + int color; + if(value == -1) { + color = 0x0; + } else if(value == 0) { + color = 0x0; + } else { + color = 0xffffff * (1.0*log((double)value)/log((double)maxIterations)); + } + + return (float4){ + 1.0 * ((color >> 0) & 0xFF) / 0xFF, + 1.0 * ((color >> 8) & 0xFF) / 0xFF, + 1.0 * ((color >> 16) & 0xFF) / 0xFF, + 1 + }; +} \ No newline at end of file diff --git a/src/main/resources/clcode/utility/shapes.cl b/src/main/resources/clcode/utility/shapes.cl new file mode 100644 index 0000000..fad35a6 --- /dev/null +++ b/src/main/resources/clcode/utility/shapes.cl @@ -0,0 +1,3 @@ +int adhlkjffjkdhlaakhjdflsjkhaflds(int2 para) { + return para.x * para.y; +} \ No newline at end of file