Document the OpenCL class

main
Brett 2022-12-03 00:39:12 -05:00
parent 5b65167bd7
commit 373134a255
12 changed files with 515 additions and 297 deletions

View File

@ -0,0 +1,73 @@
<component name="ProjectCodeStyleConfiguration">
<code_scheme name="Project" version="173">
<option name="RIGHT_MARGIN" value="150" />
<Objective-C>
<option name="INDENT_CLASS_MEMBERS" value="8" />
<option name="INDENT_VISIBILITY_KEYWORDS" value="4" />
<option name="INDENT_PREPROCESSOR_DIRECTIVE" value="4" />
<option name="FUNCTION_NON_TOP_AFTER_RETURN_TYPE_WRAP" value="0" />
<option name="FUNCTION_TOP_AFTER_RETURN_TYPE_WRAP" value="0" />
<option name="FUNCTION_PARAMETERS_NEW_LINE_AFTER_LPAR" value="true" />
<option name="FUNCTION_PARAMETERS_NEW_LINE_BEFORE_RPAR" value="true" />
<option name="LAMBDA_CAPTURE_LIST_ALIGN_MULTILINE" value="true" />
<option name="STRUCTURED_BINDING_LIST_ALIGN_MULTILINE" value="true" />
<option name="FUNCTION_CALL_ARGUMENTS_NEW_LINE_AFTER_LPAR" value="true" />
<option name="FUNCTION_CALL_ARGUMENTS_NEW_LINE_BEFORE_RPAR" value="true" />
<option name="CLASS_CONSTRUCTOR_INIT_LIST_NEW_LINE_BEFORE_COLON" value="0" />
<option name="CLASS_CONSTRUCTOR_INIT_LIST_NEW_LINE_AFTER_COLON" value="1" />
<option name="SUPERCLASS_LIST_BEFORE_COLON" value="0" />
<option name="SPACE_BEFORE_INIT_LIST_COLON" value="false" />
<option name="SPACE_BEFORE_POINTER_IN_DECLARATION" value="false" />
<option name="SPACE_AFTER_POINTER_IN_DECLARATION" value="true" />
<option name="SPACE_BEFORE_REFERENCE_IN_DECLARATION" value="false" />
<option name="SPACE_AFTER_REFERENCE_IN_DECLARATION" value="true" />
</Objective-C>
<Objective-C-extensions>
<rules>
<rule entity="NAMESPACE" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="MACRO" visibility="ANY" specifier="ANY" prefix="" style="SCREAMING_SNAKE_CASE" suffix="" />
<rule entity="CLASS" visibility="ANY" specifier="ANY" prefix="" style="PASCAL_CASE" suffix="" />
<rule entity="STRUCT" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="ENUM" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="ENUMERATOR" visibility="ANY" specifier="ANY" prefix="" style="PASCAL_CASE" suffix="" />
<rule entity="TYPEDEF" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="UNION" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="CLASS_MEMBER_FUNCTION" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="STRUCT_MEMBER_FUNCTION" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="CLASS_MEMBER_FIELD" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="STRUCT_MEMBER_FIELD" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="GLOBAL_FUNCTION" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="GLOBAL_VARIABLE" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="PARAMETER" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="LOCAL_VARIABLE" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
</rules>
</Objective-C-extensions>
<Objective-C-extensions>
<rules>
<rule entity="NAMESPACE" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="MACRO" visibility="ANY" specifier="ANY" prefix="" style="SCREAMING_SNAKE_CASE" suffix="" />
<rule entity="CLASS" visibility="ANY" specifier="ANY" prefix="" style="PASCAL_CASE" suffix="" />
<rule entity="STRUCT" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="ENUM" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="ENUMERATOR" visibility="ANY" specifier="ANY" prefix="" style="PASCAL_CASE" suffix="" />
<rule entity="TYPEDEF" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="UNION" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="CLASS_MEMBER_FUNCTION" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="STRUCT_MEMBER_FUNCTION" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="CLASS_MEMBER_FIELD" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="STRUCT_MEMBER_FIELD" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="GLOBAL_FUNCTION" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="GLOBAL_VARIABLE" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="PARAMETER" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
<rule entity="LOCAL_VARIABLE" visibility="ANY" specifier="ANY" prefix="" style="CAMEL_CASE" suffix="" />
</rules>
</Objective-C-extensions>
<codeStyleSettings language="ObjectiveC">
<option name="RIGHT_MARGIN" value="150" />
<indentOptions>
<option name="LABEL_INDENT_SIZE" value="-4" />
<option name="KEEP_INDENTS_ON_EMPTY_LINES" value="true" />
</indentOptions>
</codeStyleSettings>
</code_scheme>
</component>

View File

@ -1,5 +1,5 @@
<component name="ProjectCodeStyleConfiguration">
<state>
<option name="PREFERRED_PROJECT_CODE_STYLE" value="Default" />
<option name="USE_PER_PROJECT_SETTINGS" value="true" />
</state>
</component>

View File

@ -0,0 +1,5 @@
<component name="CopyrightManager">
<copyright>
<option name="myName" value="GPL3" />
</copyright>
</component>

View File

@ -0,0 +1,3 @@
<component name="DependencyValidationManager">
<scope name="GPL_3" pattern="file[Step 2]:*/&amp;&amp;!file[Step 2]:.idea//*&amp;&amp;!file[Step 2]:Old//*&amp;&amp;!file[Step 2]:Old 2//*&amp;&amp;!file[Step 2]:resources//*&amp;&amp;!file:cmake_make_command.sh&amp;&amp;!file:CMakeLists.txt&amp;&amp;!file:Dev Log Week of 2022-10-23.odt&amp;&amp;!file:st3-rd-debug.cap&amp;&amp;!file:st3-rd-release.cap&amp;&amp;!file:Submission.zip" />
</component>

View File

@ -0,0 +1,3 @@
<component name="DependencyValidationManager">
<scope name="GPL_3_LOCAL" pattern="file[Step 2]:*/&amp;&amp;!file[Step 2]:.idea//*&amp;&amp;!file[Step 2]:Old//*&amp;&amp;!file[Step 2]:Old 2//*&amp;&amp;!file[Step 2]:resources//*&amp;&amp;!file:cmake_make_command.sh&amp;&amp;!file:CMakeLists.txt&amp;&amp;!file:Dev Log Week of 2022-10-23.odt&amp;&amp;!file:st3-rd-debug.cap&amp;&amp;!file:st3-rd-release.cap&amp;&amp;!file:Submission.zip" />
</component>

View File

@ -141,7 +141,7 @@ namespace Raytracing {
int width, height, x, y;
};
class Raycaster {
class RayCaster {
private:
int maxBounceDepth = 50;
int raysPerPixel = 50;
@ -189,7 +189,7 @@ namespace Raytracing {
// the second creates better results but is 18% slower (better defined shadows)
// likely due to not over generating unit vectors biased towards the corners
}
Raycaster(Camera& c, Image& i, World& world, const Parser& p): camera(c), image(i), world(world) {
RayCaster(Camera& c, Image& i, World& world, const Parser& p): camera(c), image(i), world(world) {
world.generateBVH();
}
void runSTDThread(int threads = -1);
@ -211,7 +211,7 @@ namespace Raytracing {
// since executors contains the only reference to the unique_ptr it will be deleted automatically
executors.clear();
}
~Raycaster() {
~RayCaster() {
deleteThreads();
delete (unprocessedQuads);
}

View File

@ -138,7 +138,7 @@ namespace Raytracing {
World& m_world;
Shader& m_imageShader;
Shader& m_worldShader;
Raycaster& m_raycaster;
RayCaster& m_raycaster;
Parser& m_parser;
Camera& m_camera;
public:
@ -147,7 +147,7 @@ namespace Raytracing {
World& world,
Shader& mImageShader,
Shader& mWorldShader,
Raycaster& mRaycaster,
RayCaster& mRaycaster,
Parser& mParser,
Camera& mCamera)
: m_window(mWindow), m_mainImage(mMainImage), m_imageShader(mImageShader), m_worldShader(mWorldShader), m_raycaster(mRaycaster),

View File

@ -9,6 +9,7 @@
// OpenCL includes
#include <CL/cl.h>
#include <CL/cl_gl.h>
#include <engine/image/image.h>
#include <config.h>
#ifdef COMPILE_GUI
@ -34,26 +35,146 @@ namespace Raytracing {
std::unordered_map<std::string, cl_mem> buffers;
std::unordered_map<std::string, cl_kernel> kernels;
/**
* Checks for some basic errors after calling OpenCL commands. Stuff like GPU out of memory... etc.
*/
void checkBasicErrors() const;
public:
/**
* Loads the shader from a file on class creation
* @param file file to load OpenCL "shader" (code) file
*/
explicit CLProgram(const std::string& file);
/**
* Used by the OpenCL class to create a basic OpenCL program
* @param context provided by the OpenCL class.
* @param deviceID provided by the OpenCL class.
*/
void loadCLShader(cl_context context, cl_device_id deviceID);
/**
* Kernels are the entry points in OpenCL. You can have multiple of them in a single program.
* @param kernelName both the name of the kernel function in the source and the reference to the kernel object used in other functions in this class.
*/
void createKernel(const std::string& kernelName);
/**
* Buffers are the quintessential datastructures in OpenCL. They are basically regions of memory allocated to a program.
* @param bufferName the name of the buffer used to store internally
* @param flags read write flags for the buffer. One of CL_MEM_READ_ONLY | CL_MEM_WRITE_ONLY | CL_MEM_READ_WRITE
* @param bytes the number of bytes to be allocated.
*/
void createBuffer(const std::string& bufferName, cl_mem_flags flags, size_t bytes);
/**
* Creates a buffer on the GPU using the data pointed to by the supplied pointer. This copy happens as soon as this is called.
* @param bufferName the name of the buffer used to store internally
* @param flags One of CL_MEM_READ_ONLY | CL_MEM_WRITE_ONLY | CL_MEM_READ_WRITE
* @param bytes the number of bytes to be allocated. Must be less than equal to the number of bytes at ptr
* @param ptr the pointer to copy to the GPU.
*/
void createBuffer(const std::string& bufferName, cl_mem_flags flags, size_t bytes, void* ptr);
/**
* Creates a buffer on the GPU using the data pointed to by the supplied pointer. This copy happens as soon as this is called.
* @param bufferName the name of the buffer used to store internally
* @param flags One of CL_MEM_READ_ONLY | CL_MEM_WRITE_ONLY | CL_MEM_READ_WRITE
* @param bytes the number of bytes to be allocated. Must be less than equal to the number of bytes at ptr
* @param ptr the pointer to copy to the GPU.
*/
void createImage(const std::string& imageName, int width, int height);
/**
* Allows you to bind certain buffers to a specific index in the kernel's argument list.
* @param kernel kernel to bind to
* @param buffer buffer to bind to argIndex
* @param argIndex the index of the argument for this buffer.
*/
void setKernelArgument(const std::string& kernel, const std::string& buffer, int argIndex);
/**
* Runs the kernel code on the GPU. Is blocking.
* @param kernel kernel function name to call
* @param globalWorkSize the total number of times to execute the kernel function code. Corresponds to the result of get_global_id
* @param localWorkSize how many work items make up a work group to be executed by a kernel. 64 is recommended, must not exceed the printed value "device max workgroup size"
* @param globalWorkOffset not used. can be used to set an offset to the result of get_global_id
*/
void runKernel(const std::string& kernel, size_t globalWorkSize, size_t localWorkSize, const size_t* globalWorkOffset = NULL);
void runKernel(const std::string& kernel, size_t* globalWorkSize, size_t* localWorkSize, cl_uint workDim = 1, const size_t* globalWorkOffset = NULL);
/**
* Runs the kernel code on the GPU. Is blocking.
* This version allows you to specify the number of work dimensions.
* globalWorkSize and localWorkSize must be an array of workDim size which specify the work size for each kernel
* For example a work dim of 2 allows for two separate work sizes to be set per dimension.
* An image is two dimensional and so global work size would be {width of image, height of image}
* and local work size would be {8, 8} for a total of 64 (again recommended). Alternatively specify CL_D2_64_LOCAL_SIZE
* The resulting execution causes get_global_id(0) to run [0, width) times and get_global_id(1) to run [0, height) times
* @param kernel kernel function name to call
* @param globalWorkSize the total number of times to execute the kernel function code. Corresponds to the result of get_global_id(dim)
* @param localWorkSize how many work items make up a work group to be executed by a kernel. total 64 is recommended, total must not exceed the printed value "device max workgroup size"
* @param workDim number of dimensions to the work group being executed.
* @param globalWorkOffset not used. can be used to set an offset to the result of get_global_id
*/
void runKernel(
const std::string& kernel, size_t* globalWorkSize, size_t* localWorkSize, cl_uint workDim = 1,
const size_t* globalWorkOffset = NULL
);
/**
* Enqueues a write command to the buffer specified by the buffer name,
* @param buffer the buffer to write to
* @param bytes the number of bytes to be copied
* @param ptr the pointer to copy from. Must have at least bytes available
* @param blocking should this function wait for the bytes to be uploaded to the GPU?
* @param offset offset in the buffer object to write to
*/
void writeBuffer(const std::string& buffer, size_t bytes, void* ptr, cl_bool blocking = CL_TRUE, size_t offset = 0);
void readBuffer(const std::string& buffer, size_t bytes, void* ptr, cl_bool blocking = CL_TRUE, size_t offset = 0);
void readImage(const std::string& imageName, size_t width, size_t height, void* ptr, cl_bool blocking = CL_TRUE, size_t x = 0, size_t y = 0);
/**
* Enqueues a read command from the buffered specified by the buffer name.
* Defaults to blocking but can be set to be non-blocking.
* @param buffer buffer to read from
* @param bytes the number of bytes to read. Make sure ptr has at least those bytes available.
* @param ptr the ptr to write the read bytes to.
* @param blocking should we wait for the read or do it async?
* @param offset offset in the buffer to read from.
*/
void readBuffer(const std::string& buffer, size_t bytes, void* ptr, cl_bool blocking = CL_TRUE, size_t offset = 0);
/**
* Reads an image from the GPU into the memory region specified. Allocated memory region must be large enough to hold the image.
* @param imageName name of the buffer to read from
* @param width width of the image. Must be less than or equal to the width of the image on the GPU
* @param height height of the image. Also must be less than or equal to the height of the image on the GPU
* @param ptr pointer to the memory region to read into
* @param blocking should we wait for the read operation to complete? Defaults to yes.
* @param x x coordinate to start the read from. Defaults to zero since it's unlikely to be needed here. Included for possible future use.
* @param y y coordinate to start the read from.
*/
void readImage(
const std::string& imageName, size_t width, size_t height, void* ptr, cl_bool blocking = CL_TRUE, size_t x = 0, size_t y = 0
);
/**
* Reads an image buffer into a RayCasting Image class.
* Image supplied must have a with and height that matches the width and height of the image buffer specified by the name.
* @param imageName name of the buffer you wish to read from
* @param image reference to an image that you want the GPU data read into.
*/
void readImage(const std::string& imageName, Image& image);
/**
* Issues all previously queued OpenCL commands in a command-queue to the device associated with the command-queue.
*/
void flushCommands();
/**
* Blocks until all previously queued OpenCL commands in a command-queue are issued to the associated device and have completed.
*/
void finishCommands();
~CLProgram();
@ -76,12 +197,38 @@ namespace Raytracing {
cl_context m_context;
/**
* prints out the important info about the specified device.
* @param device device to data dump
*/
void printDeviceInfo(cl_device_id device);
public:
/**
* creates an opencl instance on the specified platform and device. Defaults to the first GPU device
*/
explicit OpenCL(int platformID = 0, int deviceID = 0);
/**
* Creates the global OpenCL instance for the engine
*/
static void init();
/**
* Creates an OpenCL program object using the global OpenCL connection
* @param program
*/
static void createCLProgram(CLProgram& program);
/**
* @return the number of compute units the device has
*/
static cl_uint activeDeviceComputeUnits();
/**
* the frequency in megahertz of the device
* @return
*/
static cl_uint activeDeviceFrequency();
~OpenCL();

View File

@ -176,7 +176,7 @@ int main(int argc, char** args) {
if (parser.hasOption("--gui") || parser.hasOption("-g")) {
#ifdef COMPILE_GUI
Raytracing::Raycaster raycaster{camera, image, world, parser};
Raytracing::RayCaster rayCaster{camera, image, world, parser};
Texture mainImage(&image);
CLProgram program(parser.getOptionValue("--resources") + "opencl/image.cl");
@ -189,37 +189,27 @@ int main(int argc, char** args) {
size_t works[2]{(size_t) image.getWidth(), (size_t) image.getHeight()};
size_t localWorks[2]{8, 8};
unsigned char bytes[image.getWidth() * image.getHeight() * 4];
for (int i = 0; i < image.getWidth() * image.getHeight() * 4; i++)
bytes[i] = 0;
Shader shader("../resources/shaders/basic.vs", "../resources/shaders/basic.fs");
Raytracing::DisplayRenderer renderer{*window, mainImage, world, shader, worldShader, raycaster, parser, camera};
Raytracing::DisplayRenderer renderer{*window, mainImage, world, shader, worldShader, rayCaster, parser, camera};
while (!window->shouldWindowClose()) {
window->beginUpdate();
renderer.draw();
program.runKernel("drawImage", works, localWorks, 2);
program.readImage("mainImage", image.getWidth(), image.getHeight(), bytes);
const PRECISION_TYPE colorFactor = 1.0 / 255.0;
for (int i = 0; i < image.getWidth(); i++){
for (int j = 0; j < image.getHeight(); j++){
const auto pixelData = bytes + (j * 4 * image.getWidth() + i * 4);
//tlog << (int)pixelData[0] << " " << (int)pixelData[1] << " " << (int)pixelData[2] << "\n";
image.setPixelColor(i, j, {pixelData[0] * colorFactor, pixelData[1] * colorFactor, pixelData[2] * colorFactor});
}
}
program.readImage("mainImage", image);
glPolygonMode(GL_FRONT_AND_BACK, GL_LINE);
world.drawBVH(worldShader);
glPolygonMode(GL_FRONT_AND_BACK, GL_FILL);
window->endUpdate();
}
RTSignal->haltExecution = true;
raycaster.join();
rayCaster.join();
#else
flog << "Program not compiled with GUI support! Unable to open GUI\n";
#endif
} else {
Raytracing::Raycaster rayCaster{camera, image, world, parser};
Raytracing::RayCaster rayCaster{camera, image, world, parser};
ilog << "Running RayCaster (NO_GUI)!\n";
// we don't actually have to check for --single since it's implied to be default true.
int threads = std::stoi(parser.getOptionValue("--threads"));

View File

@ -97,11 +97,11 @@ namespace Raytracing {
Vec4 color;
};
Vec4 Raycaster::raycasti(const Ray& ray, int depth){
Vec4 RayCaster::raycasti(const Ray& ray, int depth) {
return {};
}
Vec4 Raycaster::raycast(const Ray& ray) {
Vec4 RayCaster::raycast(const Ray& ray) {
Ray localRay = ray;
Vec4 color{1.0, 1.0, 1.0};
for (int CURRENT_BOUNCE = 0; CURRENT_BOUNCE < maxBounceDepth; CURRENT_BOUNCE++) {
@ -139,7 +139,7 @@ namespace Raytracing {
return color;
}
void Raycaster::runRaycastingAlgorithm(RaycasterImageBounds imageBounds, int loopX, int loopY) {
void RayCaster::runRaycastingAlgorithm(RaycasterImageBounds imageBounds, int loopX, int loopY) {
try {
int x = imageBounds.x + loopX;
int y = imageBounds.y + loopY;
@ -162,7 +162,7 @@ namespace Raytracing {
}
void Raycaster::runSTDThread(int threads){
void RayCaster::runSTDThread(int threads) {
setupQueue(partitionScreen(threads));
ilog << "Running std::thread\n";
for (int i = 0; i < threads; i++) {
@ -196,7 +196,7 @@ namespace Raytracing {
}
}
void Raycaster::runOpenMP(int threads){
void RayCaster::runOpenMP(int threads) {
setupQueue(partitionScreen(threads));
#ifdef USE_OPENMP
ilog << "Running OpenMP\n";
@ -242,7 +242,7 @@ namespace Raytracing {
system_threads;
#endif
}
void Raycaster::runMPI(std::queue<RaycasterImageBounds> bounds){
void RayCaster::runMPI(std::queue<RaycasterImageBounds> bounds) {
ilog << "Running MPI\n";
dlog << "We have " << bounds.size() << " bounds currently pending!\n";
while (!bounds.empty()) {
@ -259,7 +259,7 @@ namespace Raytracing {
#endif
}
std::vector<RaycasterImageBounds> Raycaster::partitionScreen(int threads) {
std::vector<RaycasterImageBounds> RayCaster::partitionScreen(int threads) {
// if we are running single threaded, disable everything special
// the reason we run single threaded in a seperate thread is because the GUI requires its own set of updating commands
// which cannot be blocked by the raytracer, otherwise it would become unresponsive.
@ -294,7 +294,7 @@ namespace Raytracing {
return bounds;
}
void Raycaster::setupQueue(const std::vector<RaycasterImageBounds>& bounds) {
void RayCaster::setupQueue(const std::vector<RaycasterImageBounds>& bounds) {
delete (unprocessedQuads);
unprocessedQuads = new std::queue<RaycasterImageBounds>();
for (auto& b: bounds)

View File

@ -112,7 +112,7 @@ namespace Raytracing {
}
ScatterResults DiffuseMaterial::scatter(const Ray& ray, const HitData& hitData) const {
Vec4 newRay = hitData.normal + Raytracing::Raycaster::randomUnitVector().normalize();
Vec4 newRay = hitData.normal + Raytracing::RayCaster::randomUnitVector().normalize();
// rays that are close to zero are liable to floating point precision errors
if (newRay.x() < EPSILON && newRay.y() < EPSILON && newRay.z() < EPSILON && newRay.w() < EPSILON)
@ -134,11 +134,11 @@ namespace Raytracing {
Vec4 newRay = reflect(ray.getDirection().normalize(), hitData.normal);
// make sure our reflected ray is outside the sphere and doesn't point inwards
bool shouldReflect = Vec4::dot(newRay, hitData.normal) > 0;
return {shouldReflect, Ray{hitData.hitPoint, newRay + Raycaster::randomUnitVector() * fuzzyness}, getBaseColor()};
return {shouldReflect, Ray{hitData.hitPoint, newRay + RayCaster::randomUnitVector() * fuzzyness}, getBaseColor()};
}
ScatterResults TexturedMaterial::scatter(const Ray& ray, const HitData& hitData) const {
Vec4 newRay = hitData.normal + Raytracing::Raycaster::randomUnitVector().normalize();
Vec4 newRay = hitData.normal + Raytracing::RayCaster::randomUnitVector().normalize();
// rays that are close to zero are liable to floating point precision errors
if (newRay.x() < EPSILON && newRay.y() < EPSILON && newRay.z() < EPSILON && newRay.w() < EPSILON)

View File

@ -9,12 +9,15 @@
#ifdef COMPILE_GUI
#define GLFW_EXPOSE_NATIVE_X11
#define GLFW_EXPOSE_NATIVE_GLX
#include <GLFW/glfw3.h>
#include <GLFW/glfw3native.h>
#endif
#include <cstddef>
#include <utility>
#include "engine/math/vectors.h"
namespace Raytracing {
@ -23,7 +26,9 @@ namespace Raytracing {
void OpenCL::init() {
openCl = std::make_shared<OpenCL>(0, 0);
}
OpenCL::OpenCL(int platformID, int deviceID): m_activePlatform(platformID) {
OpenCL::OpenCL(int platformID, int deviceID):
m_activePlatform(platformID) {
m_CL_ERR = CL_SUCCESS;
m_numPlatforms = 0;
m_CL_ERR = clGetPlatformIDs(0, NULL, &m_numPlatforms);
@ -88,17 +93,21 @@ namespace Raytracing {
if (!textureSupport)
elog << "Warning! The OpenCL device lacks texture support!\n";
}
void OpenCL::createCLProgram(CLProgram& program) {
program.loadCLShader(openCl->m_context, openCl->m_deviceID);
}
OpenCL::~OpenCL() {
delete[](m_platformIDs);
clReleaseDevice(m_deviceID);
clReleaseContext(m_context);
}
cl_uint OpenCL::activeDeviceComputeUnits() {
return openCl->m_computeUnits;
}
cl_uint OpenCL::activeDeviceFrequency() {
return openCl->m_deviceClockFreq;
}
@ -107,6 +116,7 @@ namespace Raytracing {
CLProgram::CLProgram(const std::string& file) {
m_source = ShaderLoader::loadShaderFile(file);
}
void CLProgram::loadCLShader(cl_context context, cl_device_id deviceID) {
this->m_context = context;
this->m_deviceID = deviceID;
@ -133,31 +143,23 @@ namespace Raytracing {
}
}
/**
* Buffers are the quintessential datastructures in OpenCL. They are basically regions of memory allocated to a program.
* @param bufferName the name of the buffer used to store internally
* @param flags read write flags for the buffer. One of CL_MEM_READ_ONLY | CL_MEM_WRITE_ONLY | CL_MEM_READ_WRITE
* @param bytes the number of bytes to be allocated.
*/
void CLProgram::createBuffer(const std::string& bufferName, cl_mem_flags flags, size_t bytes) {
// create the buffer on the GPU
cl_mem buff = clCreateBuffer(m_context, flags, bytes, NULL, &m_CL_ERR);
// then store it in our buffer map for easy access.
buffers.insert({bufferName, buff});
}
/**
* Creates a buffer on the GPU using the data pointed to by the supplied pointer. This copy happens as soon as this is called.
* @param bufferName the name of the buffer used to store internally
* @param flags One of CL_MEM_READ_ONLY | CL_MEM_WRITE_ONLY | CL_MEM_READ_WRITE
* @param bytes the number of bytes to be allocated. Must be less than equal to the number of bytes at ptr
* @param ptr the pointer to copy to the GPU.
*/
void CLProgram::createBuffer(const std::string& bufferName, cl_mem_flags flags, size_t bytes, void* ptr) {
// create the buffer on the GPU
cl_mem buff = clCreateBuffer(m_context, CL_MEM_COPY_HOST_PTR | flags, bytes, ptr, &m_CL_ERR);
// then store it in our buffer map for easy access.
buffers.insert({bufferName, buff});
}
void CLProgram::createImage(const std::string& imageName, int width, int height) {
// create the texture on the GPU
cl_image_format format{CL_RGBA, CL_UNORM_INT8};
@ -208,63 +210,37 @@ namespace Raytracing {
// then store it in our buffer map for easy access.
buffers.insert({imageName, tex});
}
/**
* Kernels are the entry points in OpenCL. You can have multiple of them in a single program.
* @param kernelName both the name of the kernel function in the source and the reference to the kernel object used in other functions in this class.
*/
void CLProgram::createKernel(const std::string& kernelName) {
auto kernel = clCreateKernel(m_program, kernelName.c_str(), &m_CL_ERR);
if (m_CL_ERR != CL_SUCCESS)
elog << "Unable to create CL kernel " << kernelName << "!\n";
kernels.insert({kernelName, kernel});
}
/**
* Allows you to bind certain buffers to a specific index in the kernel's argument list.
* @param kernel kernel to bind to
* @param buffer buffer to bind to argIndex
* @param argIndex the index of the argument for this buffer.
*/
void CLProgram::setKernelArgument(const std::string& kernel, const std::string& buffer, int argIndex) {
m_CL_ERR = clSetKernelArg(kernels[kernel], argIndex, sizeof(cl_mem), (void*) &buffers[buffer]);
if (m_CL_ERR != CL_SUCCESS)
elog << "Unable to bind argument " << buffer << " to CL kernel " << kernel << "!\n";
}
/**
* Enqueues a write command to the buffer specified by the buffer name,
* @param buffer the buffer to write to
* @param bytes the number of bytes to be copied
* @param ptr the pointer to copy from. Must have at least bytes available
* @param blocking should this function wait for the bytes to be uploaded to the GPU?
* @param offset offset in the buffer object to write to
*/
void CLProgram::writeBuffer(const std::string& buffer, size_t bytes, void* ptr, cl_bool blocking, size_t offset) {
m_CL_ERR = clEnqueueWriteBuffer(m_commandQueue, buffers[buffer], blocking, offset, bytes, ptr, 0, NULL, NULL);
if (m_CL_ERR != CL_SUCCESS)
elog << "Unable to enqueue write to " << buffer << " buffer!\n";
}
/**
* Enqueues a read command from the buffered specified by the buffer name.
* Defaults to blocking but can be set to be non-blocking.
* @param buffer buffer to read from
* @param bytes the number of bytes to read. Make sure ptr has at least those bytes available.
* @param ptr the ptr to write the read bytes to.
* @param blocking should we wait for the read or do it async?
* @param offset offset in the buffer to read from.
*/
void CLProgram::readBuffer(const std::string& buffer, size_t bytes, void* ptr, cl_bool blocking, size_t offset) {
m_CL_ERR = clEnqueueReadBuffer(m_commandQueue, buffers[buffer], blocking, offset, bytes, ptr, 0, NULL, NULL);
if (m_CL_ERR != CL_SUCCESS)
elog << "Unable to enqueue read from " << buffer << " buffer!\n";
}
/**
* Issues all previously queued OpenCL commands in a command-queue to the device associated with the command-queue.
*/
void CLProgram::flushCommands() {
clFlush(m_commandQueue);
}
/**
* Blocks until all previously queued OpenCL commands in a command-queue are issued to the associated device and have completed.
*/
void CLProgram::finishCommands() {
flushCommands();
clFinish(m_commandQueue);
@ -274,7 +250,8 @@ namespace Raytracing {
m_CL_ERR = clEnqueueNDRangeKernel(m_commandQueue, kernels[kernel], 1, globalWorkOffset, &globalWorkSize, &localWorkSize, 0, NULL, NULL);
}
void CLProgram::runKernel(const std::string& kernel, size_t* globalWorkSize, size_t* localWorkSize, cl_uint workDim, const size_t* globalWorkOffset) {
void
CLProgram::runKernel(const std::string& kernel, size_t* globalWorkSize, size_t* localWorkSize, cl_uint workDim, const size_t* globalWorkOffset) {
m_CL_ERR = clEnqueueNDRangeKernel(m_commandQueue, kernels[kernel], workDim, globalWorkOffset, globalWorkSize, localWorkSize, 0, NULL, NULL);
}
@ -305,6 +282,7 @@ namespace Raytracing {
clReleaseMemObject(buffer.second);
clReleaseCommandQueue(m_commandQueue);
}
void CLProgram::checkBasicErrors() const {
switch (m_CL_ERR) {
case CL_OUT_OF_HOST_MEMORY:
@ -331,4 +309,23 @@ namespace Raytracing {
}
}
void CLProgram::readImage(const std::string& imageName, Image& image) {
// Create an array for copying bytes from the GPU to the CPU with
unsigned char bytes[image.getWidth() * image.getHeight() * 4];
for (int i = 0; i < image.getWidth() * image.getHeight() * 4; i++)
bytes[i] = 0;
// Call to OpenCL to read the image from the GPU and send it to the location of our buffer
this->readImage(imageName, image.getWidth(), image.getHeight(), bytes);
// multiplication is much faster than division,
// so to speed up the function precalculate the color multiplier used to convert from RGBA8 format to floating point format used by the raytracer.
const PRECISION_TYPE colorFactor = 1.0 / 255.0;
// copy the data from the buffer into the image by calculating its byte offset based on its pixel position.
for (int i = 0; i < image.getWidth(); i++) {
for (int j = 0; j < image.getHeight(); j++) {
const auto pixelData = bytes + (j * 4 * image.getWidth() + i * 4);
image.setPixelColor(i, j, {pixelData[0] * colorFactor, pixelData[1] * colorFactor, pixelData[2] * colorFactor});
}
}
}
}