From 18868dfbd2ccbfc4018c54057555875c2b3d5e9a Mon Sep 17 00:00:00 2001 From: gyulaid Date: Wed, 30 Mar 2022 20:37:54 +0200 Subject: [PATCH] Random generators --- MonteCarlo/CTG.cpp | 67 ++++++++++++++++++ MonteCarlo/Halton.cpp | 60 ++++++++++++++++ MonteCarlo/Hybrid.cpp | 69 +++++++++++++++++++ MonteCarlo/LCG.cpp | 69 +++++++++++++++++++ MonteCarlo/LFG.cpp | 74 ++++++++++++++++++++ MonteCarlo/MonteCarlo.cpp | 98 +++------------------------ MonteCarlo/MonteCarlo.vcxproj | 8 +++ MonteCarlo/MonteCarlo.vcxproj.filters | 20 ++++++ MonteCarlo/MonteCarloTests.h | 84 +++++++++++++++++++++++ 9 files changed, 462 insertions(+), 87 deletions(-) create mode 100644 MonteCarlo/CTG.cpp create mode 100644 MonteCarlo/Halton.cpp create mode 100644 MonteCarlo/Hybrid.cpp create mode 100644 MonteCarlo/LCG.cpp create mode 100644 MonteCarlo/LFG.cpp create mode 100644 MonteCarlo/MonteCarloTests.h diff --git a/MonteCarlo/CTG.cpp b/MonteCarlo/CTG.cpp new file mode 100644 index 0000000..7e329d7 --- /dev/null +++ b/MonteCarlo/CTG.cpp @@ -0,0 +1,67 @@ +#include + +#include "MonteCarloTests.h" +#include "Common.h" + +CTG::CTG(size_t _randomNumbers, size_t _threadCount) +{ + randomNumbers = _randomNumbers; + threadCount = _threadCount; + randomBuffer = std::vector(threadCount * randomNumbers); + seedBuffer = std::vector(threadCount); + kernel_name = "randomCTG"; + + for (int i = 0; i < threadCount; ++i) + { + seedBuffer[i] = rand(); + } +} + +void CTG::collect_results(cl::CommandQueue* queue) { + +} + +void CTG::gpu_compute(cl::Context* context, cl::CommandQueue* queue, cl::Program* program, cl::Event* Event) { + cl_int err = CL_SUCCESS; + cl::Kernel kernel = cl::Kernel(*program, kernel_name.c_str(), &err); + CheckCLError(err); + + clInputBuffer = cl::Buffer(*context, CL_MEM_READ_ONLY, sizeof(float) * threadCount, NULL, &err); + queue->enqueueWriteBuffer(clInputBuffer, true, 0, sizeof(float) * threadCount, seedBuffer.data()); + + clResultBuffer = cl::Buffer(*context, CL_MEM_WRITE_ONLY, sizeof(float) * threadCount * randomNumbers, NULL, &err); + + // Set the kernel parameters + kernel.setArg(0, randomNumbers); + kernel.setArg(1, clInputBuffer); + kernel.setArg(2, clResultBuffer); + + // Enqueue the kernel: threadCount threads in total, each generating random numbers in [0,1] randomNumbers times + queue->enqueueNDRangeKernel(kernel, + cl::NullRange, + cl::NDRange(threadCount, 1), + cl::NullRange, + NULL, + Event); +} + +void CTG::cpu_compute() +{ +} + +bool CTG::validate_results() +{ + std::ofstream ofs("randoms_" + description() + ".txt", std::ofstream::out); + for (int i = 0; i < threadCount; ++i) { + for (int j = 0; j < randomNumbers; ++j) { + ofs << j << " " << randomBuffer[i + j * threadCount] << std::endl; + } + } + ofs.close(); + return true; +} + +std::string CTG::description() +{ + return kernel_name + "(numbers=" + std::to_string(randomNumbers) + ",threads=" + std::to_string(threadCount) + ")"; +} \ No newline at end of file diff --git a/MonteCarlo/Halton.cpp b/MonteCarlo/Halton.cpp new file mode 100644 index 0000000..c7b53bb --- /dev/null +++ b/MonteCarlo/Halton.cpp @@ -0,0 +1,60 @@ +#include + +#include "MonteCarloTests.h" +#include "Common.h" + +Halton::Halton(size_t _randomNumbers, size_t _threadCount, size_t _base) { + randomNumbers = _randomNumbers; + threadCount = _threadCount; + base = _base; + randomBuffer = std::vector(threadCount * randomNumbers); + kernel_name = "haltonSequence"; +} + +void Halton::collect_results(cl::CommandQueue* queue) +{ + queue->enqueueReadBuffer(clResultBuffer, true, 0, sizeof(float) * threadCount * randomNumbers, randomBuffer.data()); +} + +void Halton::gpu_compute(cl::Context* context, cl::CommandQueue* queue, cl::Program* program, cl::Event* Event) +{ + cl_int err = CL_SUCCESS; + cl::Kernel kernel = cl::Kernel(*program, kernel_name.c_str(), &err); + CheckCLError(err); + + clResultBuffer = cl::Buffer(*context, CL_MEM_WRITE_ONLY, sizeof(float) * threadCount * randomNumbers, NULL, &err); + + // Set the kernel parameters + kernel.setArg(0, randomNumbers); + kernel.setArg(1, base); + kernel.setArg(2, clResultBuffer); + + // Enqueue the kernel: threadCount threads in total, each generating random numbers in [0,1] randomNumbers times + queue->enqueueNDRangeKernel(kernel, + cl::NullRange, + cl::NDRange(threadCount, 1), + cl::NullRange, + NULL, + Event); +} + +void Halton::cpu_compute() +{ +} + +bool Halton::validate_results() +{ + std::ofstream ofs("randoms_" + description() + ".txt", std::ofstream::out); + for (int i = 0; i < threadCount; ++i) { + for (int j = 0; j < randomNumbers; ++j) { + ofs << j << " " << randomBuffer[i + j * threadCount] << std::endl; + } + } + ofs.close(); + return true; +} + +std::string Halton::description() +{ + return kernel_name + "(numbers=" + std::to_string(randomNumbers) + ",threads=" + std::to_string(threadCount) + ",base=" + std::to_string(base) + ")"; +} \ No newline at end of file diff --git a/MonteCarlo/Hybrid.cpp b/MonteCarlo/Hybrid.cpp new file mode 100644 index 0000000..6729718 --- /dev/null +++ b/MonteCarlo/Hybrid.cpp @@ -0,0 +1,69 @@ +#include + +#include "MonteCarloTests.h" +#include "Common.h" + +Hybrid::Hybrid(size_t _randomNumbers, size_t _threadCount) +{ + randomNumbers = _randomNumbers; + threadCount = _threadCount; + randomBuffer = std::vector(threadCount * randomNumbers); + seedBuffer = std::vector(threadCount); + kernel_name = "hybridRNG"; + + for (int i = 0; i < threadCount; ++i) + { + seedBuffer[i] = rand(); + } +} + +void Hybrid::collect_results(cl::CommandQueue* queue) +{ + queue->enqueueReadBuffer(clResultBuffer, true, 0, sizeof(float) * threadCount * randomNumbers, randomBuffer.data()); +} + +void Hybrid::gpu_compute(cl::Context* context, cl::CommandQueue* queue, cl::Program* program, cl::Event* Event) +{ + cl_int err = CL_SUCCESS; + cl::Kernel kernel = cl::Kernel(*program, kernel_name.c_str(), &err); + CheckCLError(err); + + clInputBuffer = cl::Buffer(*context, CL_MEM_READ_ONLY, sizeof(float) * threadCount, NULL, &err); + queue->enqueueWriteBuffer(clInputBuffer, true, 0, sizeof(float) * threadCount, seedBuffer.data()); + + clResultBuffer = cl::Buffer(*context, CL_MEM_WRITE_ONLY, sizeof(float) * threadCount * randomNumbers, NULL, &err); + + // Set the kernel parameters + kernel.setArg(0, randomNumbers); + kernel.setArg(1, clInputBuffer); + kernel.setArg(2, clResultBuffer); + + // Enqueue the kernel: threadCount threads in total, each generating random numbers in [0,1] randomNumbers times + queue->enqueueNDRangeKernel(kernel, + cl::NullRange, + cl::NDRange(threadCount, 1), + cl::NullRange, + NULL, + Event); +} + +void Hybrid::cpu_compute() +{ +} + +bool Hybrid::validate_results() +{ + std::ofstream ofs("randoms_" + description() + ".txt", std::ofstream::out); + for (int i = 0; i < threadCount; ++i) { + for (int j = 0; j < randomNumbers; ++j) { + ofs << j << " " << randomBuffer[i + j * threadCount] << std::endl; + } + } + ofs.close(); + return true; +} + +std::string Hybrid::description() +{ + return kernel_name + "(numbers=" + std::to_string(randomNumbers) + ",threads=" + std::to_string(threadCount) + ")"; +} diff --git a/MonteCarlo/LCG.cpp b/MonteCarlo/LCG.cpp new file mode 100644 index 0000000..dcec479 --- /dev/null +++ b/MonteCarlo/LCG.cpp @@ -0,0 +1,69 @@ +#include + +#include "MonteCarloTests.h" +#include "Common.h" + +LCG::LCG(size_t _randomNumbers, size_t _threadCount) +{ + randomNumbers = _randomNumbers; + threadCount = _threadCount; + randomBuffer = std::vector(threadCount * randomNumbers); + seedBuffer = std::vector(threadCount); + kernel_name = "randomLCG"; + + for (int i = 0; i < threadCount; ++i) + { + seedBuffer[i] = rand(); + } +} + +void LCG::collect_results(cl::CommandQueue* queue) +{ + queue->enqueueReadBuffer(clResultBuffer, true, 0, sizeof(float) * threadCount * randomNumbers, randomBuffer.data()); +} + +void LCG::gpu_compute(cl::Context* context, cl::CommandQueue* queue, cl::Program* program, cl::Event* Event) +{ + cl_int err = CL_SUCCESS; + cl::Kernel kernel = cl::Kernel(*program, kernel_name.c_str(), &err); + CheckCLError(err); + + clInputBuffer = cl::Buffer(*context, CL_MEM_READ_ONLY, sizeof(float) * threadCount, NULL, &err); + queue->enqueueWriteBuffer(clInputBuffer, true, 0, sizeof(float) * threadCount, seedBuffer.data()); + + clResultBuffer = cl::Buffer(*context, CL_MEM_WRITE_ONLY, sizeof(float) * threadCount * randomNumbers, NULL, &err); + + // Set the kernel parameters + kernel.setArg(0, randomNumbers); + kernel.setArg(1, clInputBuffer); + kernel.setArg(2, clResultBuffer); + + // Enqueue the kernel: threadCount threads in total, each generating random numbers in [0,1] randomNumbers times + queue->enqueueNDRangeKernel(kernel, + cl::NullRange, + cl::NDRange(threadCount, 1), + cl::NullRange, + NULL, + Event); +} + +void LCG::cpu_compute() +{ +} + +bool LCG::validate_results() +{ + std::ofstream ofs("randoms_" + description() + ".txt", std::ofstream::out); + for (int i = 0; i < threadCount; ++i) { + for (int j = 0; j < randomNumbers; ++j) { + ofs << j << " " << randomBuffer[i + j * threadCount] << std::endl; + } + } + ofs.close(); + return true; +} + +std::string LCG::description() +{ + return kernel_name + "(numbers=" + std::to_string(randomNumbers) +",threads="+std::to_string(threadCount) + ")"; +} diff --git a/MonteCarlo/LFG.cpp b/MonteCarlo/LFG.cpp new file mode 100644 index 0000000..7e28f13 --- /dev/null +++ b/MonteCarlo/LFG.cpp @@ -0,0 +1,74 @@ +#include + +#include "MonteCarloTests.h" +#include "Common.h" + +LFG::LFG(size_t _randomNumbers, size_t _threadCount, size_t _randomStateSize) +{ + randomNumbers = _randomNumbers; + threadCount = _threadCount; + randomBuffer = std::vector(threadCount * randomNumbers); + seedBuffer = std::vector(threadCount); + kernel_name = "randomLFG"; + randomStateSize = _randomStateSize; + + for (int i = 0; i < threadCount; ++i) + { + seedBuffer[i] = rand(); + } +} + +void LFG::collect_results(cl::CommandQueue* queue) +{ + queue->enqueueReadBuffer(clResultBuffer, true, 0, sizeof(float) * threadCount * randomNumbers, randomBuffer.data()); +} + +void LFG::gpu_compute(cl::Context* context, cl::CommandQueue* queue, cl::Program* program, cl::Event* Event) +{ + cl_int err = CL_SUCCESS; + cl::Kernel kernel = cl::Kernel(*program, kernel_name.c_str(), &err); + CheckCLError(err); + + clInputBuffer = cl::Buffer(*context, CL_MEM_READ_ONLY, sizeof(float) * threadCount, NULL, &err); + queue->enqueueWriteBuffer(clInputBuffer, true, 0, sizeof(float) * threadCount, seedBuffer.data()); + + clResultBuffer = cl::Buffer(*context, CL_MEM_WRITE_ONLY, sizeof(float) * threadCount * randomNumbers, NULL, &err); + cl::Buffer randomStates(*context, CL_MEM_WRITE_ONLY, sizeof(float) * threadCount * randomStateSize, NULL, &err); + + // Set the kernel parameters + kernel.setArg(0, randomNumbers); + kernel.setArg(1, clInputBuffer); + kernel.setArg(2, randomStateSize); + kernel.setArg(3, randomStates); + kernel.setArg(4, clResultBuffer); + + // Enqueue the kernel: threadCount threads in total, each generating random numbers in [0,1] randomNumbers times + queue->enqueueNDRangeKernel(kernel, + cl::NullRange, + cl::NDRange(threadCount, 1), + cl::NullRange, + NULL, + Event); +} + +void LFG::cpu_compute() +{ +} + +bool LFG::validate_results() +{ + std::ofstream ofs("randoms_" + description() + ".txt", std::ofstream::out); + for (int i = 0; i < threadCount; ++i) { + for (int j = 0; j < randomNumbers; ++j) { + ofs << j << " " << randomBuffer[i + j * threadCount] << std::endl; + } + } + ofs.close(); + return true; +} + +std::string LFG::description() +{ + return kernel_name + "(numbers=" + std::to_string(randomNumbers) + ",threads=" + std::to_string(threadCount) + ",state="+std::to_string(randomStateSize) + ")"; +} + diff --git a/MonteCarlo/MonteCarlo.cpp b/MonteCarlo/MonteCarlo.cpp index 3f0501d..efd6b45 100644 --- a/MonteCarlo/MonteCarlo.cpp +++ b/MonteCarlo/MonteCarlo.cpp @@ -11,95 +11,13 @@ #include "cl.hpp" #include #include +#include +#include "MonteCarloTests.h" const bool writeOutRandoms = true; const size_t randomNumbers = 1024; const size_t threadCount = 512; -void capi() -{ - // Get a platform ID - cl_platform_id platformID; - clGetPlatformIDs(1, &platformID, NULL); - - // Get a device ID - cl_device_id deviceID; - clGetDeviceIDs(platformID, CL_DEVICE_TYPE_GPU, 1, &deviceID, NULL); - - // Create a context - cl_context context; - cl_context_properties contextProperties[] = - { CL_CONTEXT_PLATFORM, (cl_context_properties)platformID, 0 }; - context = clCreateContext(contextProperties, 1, &deviceID, NULL, NULL, NULL); - - // Create a command queue - cl_command_queue queue; - queue = clCreateCommandQueue(context, deviceID, CL_QUEUE_PROFILING_ENABLE, NULL); - - // Create an OpenCL program - std::string source = FileToString("../kernels/montecarlo.cl"); - const char* csource = source.c_str(); - cl_program program = clCreateProgramWithSource(context, 1, &csource, NULL, NULL); - cl_int err = clBuildProgram(program, 1, &deviceID, NULL, NULL, NULL); - if (err != CL_SUCCESS) - { - cl_uint logLength; - clGetProgramBuildInfo(program, deviceID, CL_PROGRAM_BUILD_LOG, 0, NULL, &logLength); - char* log = new char[logLength]; - clGetProgramBuildInfo(program, deviceID, CL_PROGRAM_BUILD_LOG, logLength, log, 0); - std::cout << log << std::endl; - delete[] log; - exit(-1); - } - - // Get the kernel handle - cl_kernel kernel = clCreateKernel(program, "randomLCG", &err); - if(!CheckCLError(err)) exit(-1); - - // Allocate memory for random numbers and random seeds - // Every thread receives a private seed, stored in seedBuffer/clSeedBuffer - // Every thread is supposed to generate randomNumbers random numbers - std::vector randomBuffer(threadCount * randomNumbers); - std::vector seedBuffer(threadCount); - for (int i = 0; i < threadCount; ++i) - { - seedBuffer[i] = rand(); - } - - cl_mem randomBufferDev; - randomBufferDev = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float)* threadCount * randomNumbers, NULL, &err); - if (!CheckCLError(err)) exit(-1); - cl_mem seedBufferDev; - seedBufferDev = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * threadCount, NULL, &err); - if (!CheckCLError(err)) exit(-1); - - clEnqueueWriteBuffer(queue, seedBufferDev, CL_TRUE, 0, sizeof(float) * threadCount, seedBuffer.data(), 0, NULL, NULL); - - // Set the kernel paramateres - clSetKernelArg(kernel, 0, sizeof(int), &randomNumbers); - clSetKernelArg(kernel, 1, sizeof(cl_mem), &seedBufferDev); - clSetKernelArg(kernel, 2, sizeof(cl_mem), &randomBufferDev); - // Enqueue the kernel: threadCount threads in total, each generating randomNumbers random numbers in [0,1] - clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &threadCount, NULL, 0, NULL, NULL); - // Copy the result back to the host - clEnqueueReadBuffer(queue, randomBufferDev, CL_TRUE, 0, sizeof(float) * threadCount * randomNumbers, randomBuffer.data(), 0, NULL, NULL); - - // Write out the output - if(writeOutRandoms == true) - { - std::ofstream ofs("randoms.txt", std::ofstream::out); - for (int i = 0; i < threadCount; ++i) { - for (int j = 0; j < randomNumbers; ++j) { - ofs << j << " " << randomBuffer[i + j * threadCount] << std::endl; - } - ofs << std::endl; - } - ofs.close(); - } - - std::cout << "Finished" << std::endl; -} - void cppapi() { cl_int err = CL_SUCCESS; @@ -126,7 +44,7 @@ void cppapi() cl::CommandQueue queue(context, devices[0], 0, &err); // Create the OpenCL program - std::string programSource = FileToString("../kernels/programs.cl"); + std::string programSource = FileToString("../kernels/montecarlo.cl"); cl::Program program = cl::Program(context, programSource); err = program.build(devices); if (!CheckCLError(err)) @@ -195,8 +113,14 @@ void cppapi() int main() { - capi(); - cppapi(); + //cppapi(); + OpenCLHandler handler("../kernels/montecarlo.cl"); + + //handler.run_test(new LCG(randomNumbers, threadCount)); + //handler.run_test(new LFG(randomNumbers, threadCount, 256)); + //handler.run_test(new CTG(randomNumbers, threadCount)); + //handler.run_test(new Hybrid(randomNumbers, threadCount)); + handler.run_test(new Halton(randomNumbers, threadCount, 2)); return 0; } diff --git a/MonteCarlo/MonteCarlo.vcxproj b/MonteCarlo/MonteCarlo.vcxproj index 3429527..554cb7a 100644 --- a/MonteCarlo/MonteCarlo.vcxproj +++ b/MonteCarlo/MonteCarlo.vcxproj @@ -72,6 +72,11 @@ + + + + + @@ -83,6 +88,9 @@ {f66311cb-c60d-43de-890c-7e6d8179ca44} + + + diff --git a/MonteCarlo/MonteCarlo.vcxproj.filters b/MonteCarlo/MonteCarlo.vcxproj.filters index 8c74616..4876bc4 100644 --- a/MonteCarlo/MonteCarlo.vcxproj.filters +++ b/MonteCarlo/MonteCarlo.vcxproj.filters @@ -21,6 +21,21 @@ Source Files + + Source Files + + + Source Files + + + Source Files + + + Source Files + + + Source Files + @@ -30,4 +45,9 @@ Kernels + + + Header Files + + \ No newline at end of file diff --git a/MonteCarlo/MonteCarloTests.h b/MonteCarlo/MonteCarloTests.h new file mode 100644 index 0000000..3f8383b --- /dev/null +++ b/MonteCarlo/MonteCarloTests.h @@ -0,0 +1,84 @@ +#pragma once +#include "Tests.h" + + +class LCG : public TestCase { +private: + size_t randomNumbers; + size_t threadCount; + std::vector randomBuffer; + std::vector seedBuffer; + std::string kernel_name; +public: + LCG(size_t _randomNumbers, size_t _threadCount); + void collect_results(cl::CommandQueue* queue); + void gpu_compute(cl::Context* context, cl::CommandQueue* queue, cl::Program* program, cl::Event* Event); + void cpu_compute(); + bool validate_results(); + std::string description(); +}; + +class LFG : public TestCase { +private: + size_t randomNumbers; + size_t threadCount; + std::vector randomBuffer; + std::vector seedBuffer; + std::string kernel_name; + size_t randomStateSize; +public: + LFG(size_t _randomNumbers, size_t _threadCount, size_t _randomStateSize); + void collect_results(cl::CommandQueue* queue); + void gpu_compute(cl::Context* context, cl::CommandQueue* queue, cl::Program* program, cl::Event* Event); + void cpu_compute(); + bool validate_results(); + std::string description(); +}; + +class CTG : public TestCase { +private: + size_t randomNumbers; + size_t threadCount; + std::vector randomBuffer; + std::vector seedBuffer; + std::string kernel_name; +public: + CTG(size_t _randomNumbers, size_t _threadCount); + void collect_results(cl::CommandQueue* queue); + void gpu_compute(cl::Context* context, cl::CommandQueue* queue, cl::Program* program, cl::Event* Event); + void cpu_compute(); + bool validate_results(); + std::string description(); +}; + +class Hybrid : public TestCase { +private: + size_t randomNumbers; + size_t threadCount; + std::vector randomBuffer; + std::vector seedBuffer; + std::string kernel_name; +public: + Hybrid(size_t _randomNumbers, size_t _threadCount); + void collect_results(cl::CommandQueue* queue); + void gpu_compute(cl::Context* context, cl::CommandQueue* queue, cl::Program* program, cl::Event* Event); + void cpu_compute(); + bool validate_results(); + std::string description(); +}; + +class Halton : public TestCase { +private: + size_t randomNumbers; + size_t threadCount; + size_t base; + std::vector randomBuffer; + std::string kernel_name; +public: + Halton(size_t _randomNumbers, size_t _threadCount, size_t _base); + void collect_results(cl::CommandQueue* queue); + void gpu_compute(cl::Context* context, cl::CommandQueue* queue, cl::Program* program, cl::Event* Event); + void cpu_compute(); + bool validate_results(); + std::string description(); +}; \ No newline at end of file