From c76ce8a768d7f881a152279640a12823783b50c0 Mon Sep 17 00:00:00 2001 From: Ivan Mahonin Date: Aug 27 2015 19:00:29 +0000 Subject: OpenCL --- diff --git a/c++/contourgl/Makefile b/c++/contourgl/Makefile index b550e82..9a6d776 100644 --- a/c++/contourgl/Makefile +++ b/c++/contourgl/Makefile @@ -1,9 +1,9 @@ -DEPLIBS = gl x11 +DEPLIBS = gl x11 OpenCL -CXXFLAGS = -O3 -Wall -fmessage-length=0 `pkg-config --libs $(DEPLIBS)` -DGL_GLEXT_PROTOTYPES -DEPS = contour.h contourbuilder.h geometry.h polyspan.h rendersw.h shaders.h test.h \ - contour.cpp contourbuilder.cpp polyspan.cpp rendersw.cpp shaders.cpp test.cpp -OBJS = contour.o contourgl.o contourbuilder.o polyspan.o rendersw.o shaders.o test.o +CXXFLAGS = -O0 -g -Wall -fmessage-length=0 `pkg-config --cflags $(DEPLIBS)` -DGL_GLEXT_PROTOTYPES +DEPS = clcontext.h contour.h contourgl.h contourbuilder.h geometry.h polyspan.h rendersw.h shaders.h test.h triangulator.h \ + clcontext.cpp contour.cpp contourgl.cpp contourbuilder.cpp polyspan.cpp rendersw.cpp shaders.cpp test.cpp triangulator.cpp +OBJS = clcontext.o contour.o contourgl.o contourbuilder.o polyspan.o rendersw.o shaders.o test.o triangulator.o LIBS = `pkg-config --libs $(DEPLIBS)` TARGET = contourgl diff --git a/c++/contourgl/cl/hello.cl b/c++/contourgl/cl/hello.cl new file mode 100644 index 0000000..7d453ad --- /dev/null +++ b/c++/contourgl/cl/hello.cl @@ -0,0 +1,6 @@ +//#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable +__constant char s[] = "Hello!"; +__kernel void hello(__global char *out) { + size_t tid = get_global_id(0); + out[tid] = s[tid]; +} \ No newline at end of file diff --git a/c++/contourgl/clcontext.cpp b/c++/contourgl/clcontext.cpp new file mode 100644 index 0000000..b7ee20d --- /dev/null +++ b/c++/contourgl/clcontext.cpp @@ -0,0 +1,123 @@ +/* + ......... 2015 Ivan Mahonin + + This program is free software: you can redistribute it and/or modify + it under the terms of the GNU General Public License as published by + the Free Software Foundation, either version 3 of the License, or + (at your option) any later version. + + This program is distributed in the hope that it will be useful, + but WITHOUT ANY WARRANTY; without even the implied warranty of + MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + GNU General Public License for more details. + + You should have received a copy of the GNU General Public License + along with this program. If not, see . +*/ + +#include + +#include +#include +#include + +#include "clcontext.h" + + +using namespace std; + + +ClContext::ClContext(): err(), context() { + + // platform + + cl_uint platform_count = 0; + clGetPlatformIDs(0, NULL, &platform_count); + assert(platform_count); + cout << platform_count << " platforms" << endl; + vector platforms(platform_count); + clGetPlatformIDs(platforms.size(), &platforms.front(), NULL); + cl_platform_id platform = platforms[0]; + + char vendor[256] = { }; + err = clGetPlatformInfo(platform, CL_PLATFORM_VENDOR, sizeof(vendor), vendor, NULL); + assert(!err); + cout << "Use CL platform 0 by " << vendor << endl; + + // devices + + cl_uint device_count = 0; + err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &device_count); + assert(!err); + cout << device_count << " devices" << endl; + + devices.resize(device_count); + err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, devices.size(), &devices.front(), NULL); + assert(!err); + + // context + + context = clCreateContext(0, 1, &devices.front(), NULL, NULL, &err); + assert(context); +} + +ClContext::~ClContext() { + clReleaseContext(context); +} + +void ClContext::hello() { + + // data + + char data[] = "......"; + + // buffer + + cl_mem buffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR, sizeof(data), data, NULL); + assert(buffer); + + // program + + ifstream f("cl/hello.cl"); + string text((istreambuf_iterator(f)), istreambuf_iterator()); + const char *text_pointer = text.c_str(); + cl_program program = clCreateProgramWithSource(context, 1, &text_pointer, NULL, NULL); + assert(program); + + err = clBuildProgram(program, devices.size(), &devices.front(), "", NULL, NULL); + assert(!err); + + // kernel + + cl_kernel kernel = clCreateKernel(program, "hello", NULL); + assert(kernel); + err = clSetKernelArg(kernel, 0, sizeof(buffer), &buffer); + assert(!err); + + // command queue + + cl_command_queue queue = clCreateCommandQueue(context, devices[0], 0, NULL); + assert(queue); + + size_t work_group_size = sizeof(data); + cl_event event = NULL; + err = clEnqueueNDRangeKernel( + queue, + kernel, + 1, + NULL, + &work_group_size, + NULL, + 0, + NULL, + &event ); + assert(!err); + + clWaitForEvents(1, &event); + + // read + + clEnqueueReadBuffer(queue, buffer, CL_TRUE, 0, sizeof(data), data, 0, NULL, &event); + clWaitForEvents(1, &event); + cout << data << endl; +} diff --git a/c++/contourgl/clcontext.h b/c++/contourgl/clcontext.h new file mode 100644 index 0000000..57fe157 --- /dev/null +++ b/c++/contourgl/clcontext.h @@ -0,0 +1,33 @@ +/* + ......... 2015 Ivan Mahonin + + This program is free software: you can redistribute it and/or modify + it under the terms of the GNU General Public License as published by + the Free Software Foundation, either version 3 of the License, or + (at your option) any later version. + + This program is distributed in the hope that it will be useful, + but WITHOUT ANY WARRANTY; without even the implied warranty of + MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + GNU General Public License for more details. + + You should have received a copy of the GNU General Public License + along with this program. If not, see . +*/ + +#include + +#include + + +class ClContext { +public: + cl_int err; + cl_context context; + std::vector devices; + + ClContext(); + ~ClContext(); + + void hello(); +}; diff --git a/c++/contourgl/contourgl.cpp b/c++/contourgl/contourgl.cpp index a6f6514..3e8655f 100644 --- a/c++/contourgl/contourgl.cpp +++ b/c++/contourgl/contourgl.cpp @@ -25,6 +25,7 @@ #include #include +#include "clcontext.h" #include "test.h" #include "shaders.h" @@ -93,12 +94,16 @@ int main() { int framebuffer_width = 512; int framebuffer_height = 512; int framebuffer_samples = 16; - bool antialising = true; + bool antialising = false; + bool hdr = false; + + GLenum internal_format = hdr ? GL_RGBA16F : GL_RGBA; + GLenum color_type = hdr ? GL_FLOAT : GL_UNSIGNED_BYTE; GLuint multisample_texture_id = 0; glGenTextures(1, &multisample_texture_id); glBindTexture(GL_TEXTURE_2D_MULTISAMPLE, multisample_texture_id); - glTexImage2DMultisample(GL_TEXTURE_2D_MULTISAMPLE, framebuffer_samples, GL_RGBA16F, framebuffer_width, framebuffer_height, GL_TRUE); + glTexImage2DMultisample(GL_TEXTURE_2D_MULTISAMPLE, framebuffer_samples, internal_format, framebuffer_width, framebuffer_height, GL_TRUE); GLuint multisample_renderbuffer_id = 0; glGenRenderbuffers(1, &multisample_renderbuffer_id); @@ -114,7 +119,7 @@ int main() { GLuint texture_id = 0; glGenTextures(1, &texture_id); glBindTexture(GL_TEXTURE_2D, texture_id); - glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA32F, framebuffer_width, framebuffer_height, 0, GL_RGBA, GL_FLOAT, NULL); + glTexImage2D(GL_TEXTURE_2D, 0, internal_format, framebuffer_width, framebuffer_height, 0, GL_RGBA, color_type, NULL); GLuint renderbuffer_id = 0; glGenRenderbuffers(1, &renderbuffer_id); @@ -146,7 +151,9 @@ int main() { //Test::test1(); //Test::test2(); //Test::test3(); - Test::test4(); + //Test::test4(); + + ClContext().hello(); Shaders::deinitialize();