From 3836335f414f3fcba9416bd9f0480696ce05aa6f Mon Sep 17 00:00:00 2001 From: Ivan Mahonin Date: Jul 26 2018 15:51:55 +0000 Subject: contourgl: CUDA initialization --- diff --git a/c++/contourgl/.gitignore b/c++/contourgl/.gitignore index d87bd8d..38804c5 100644 --- a/c++/contourgl/.gitignore +++ b/c++/contourgl/.gitignore @@ -1,2 +1,2 @@ /contourgl - +/*.d diff --git a/c++/contourgl/Makefile b/c++/contourgl/Makefile index 98eca6f..09ef1f2 100644 --- a/c++/contourgl/Makefile +++ b/c++/contourgl/Makefile @@ -1,22 +1,29 @@ +# config + +# just comment following line to disable CUDA +CUDA = cuda-9.2 +CUDA_PATH := /opt/$(CUDA) +CUDA_BIN := $(CUDA_PATH)/bin +CUDA_PKGCONFIG := $(CUDA_PATH)/pkgconfig + DEPLIBS = gl x11 OpenCL -CXXFLAGS = -O3 -Wall -fmessage-length=0 `pkg-config --cflags $(DEPLIBS)` -DGL_GLEXT_PROTOTYPES - -HEADERS = \ - clcontext.h \ - clrender.h \ - contour.h \ - contourbuilder.h \ - environment.h \ - geometry.h \ - glcontext.h \ - measure.h \ - polyspan.h \ - shaders.h \ - swrender.h \ - test.h \ - triangulator.h \ - utils.h + +# compute build options + +CXXFLAGS := $(CXXFLAGS) -O3 -Wall -fmessage-length=0 -DGL_GLEXT_PROTOTYPES +CXXFLAGS := $(CXXFLAGS) $(shell pkg-config --cflags $(DEPLIBS)) +LIBS := $(LIBS) $(shell pkg-config --libs $(DEPLIBS)) + +ifdef CUDA + CXXFLAGS := $(CXXFLAGS) -DCUDA $(shell PKG_CONFIG_PATH=$(CUDA_PKGCONFIG) pkg-config --cflags $(CUDA)) + LIBS := $(LIBS) $(shell PKG_CONFIG_PATH=$(CUDA_PKGCONFIG) pkg-config --libs $(CUDA)) +endif + + +# files lists + +TARGET = contourgl SOURCES = \ contourgl.cpp \ @@ -35,31 +42,44 @@ SOURCES = \ triangulator.cpp \ utils.cpp -OBJS = \ - contourgl.o \ - clcontext.o \ - clrender.o \ - contour.o \ - contourbuilder.o \ - environment.o \ - geometry.o \ - glcontext.o \ - measure.o \ - polyspan.o \ - shaders.o \ - swrender.o \ - test.o \ - triangulator.o \ - utils.o - -DEPS = $(HEADERS) $(SOURCES) -LIBS = `pkg-config --libs $(DEPLIBS)` -TARGET = contourgl +ifdef CUDA + SOURCES += cudacontext.cpp + CUDA_SOURCES = \ + cuda/contour.cu \ + cuda/hello.cu +endif -$(TARGET): $(OBJS) + +# files lists postprocessing + +OBJS = $(SOURCES:.cpp=.o) +DEPS = $(SOURCES:.cpp=.d) +PTXS = $(CUDA_SOURCES:.cu=.ptx) + +# internal targets + +%.ptx: %.cu + $(CUDA_PATH)/bin/nvcc -ptx $< -o $@ + +# rule for make *.d files with include (.h) dependencies +%.d: %.cpp + @set -e; rm -f $@; \ + $(CXX) -MM $(CXXFLAGS) $< > $@.$$$$; \ + sed 's,\($*\)\.o[ :]*,\1.o $@ : ,g' < $@.$$$$ > $@; \ + rm -f $@.$$$$ + +$(TARGET): $(PTXS) $(DEPS) $(OBJS) $(CXX) -o $(TARGET) $(OBJS) $(LIBS) -all: $(TARGET) +# include rules with include (.h) dependencies +-include $(DEPS) + +# actual targets + +# declare that 'all', 'cuda' and 'clean' are cannot be a filenames +.PHONY: all clean + +all: $(TARGET) clean: - rm -f $(OBJS) $(TARGET) + rm -f $(PTXS) $(OBJS) $(DEPS) $(TARGET) diff --git a/c++/contourgl/cuda/.gitignore b/c++/contourgl/cuda/.gitignore new file mode 100644 index 0000000..a224e7f --- /dev/null +++ b/c++/contourgl/cuda/.gitignore @@ -0,0 +1 @@ +/*.ptx diff --git a/c++/contourgl/cuda/contour.cu b/c++/contourgl/cuda/contour.cu new file mode 100644 index 0000000..e69de29 --- /dev/null +++ b/c++/contourgl/cuda/contour.cu diff --git a/c++/contourgl/cuda/hello.cu b/c++/contourgl/cuda/hello.cu new file mode 100644 index 0000000..e494eeb --- /dev/null +++ b/c++/contourgl/cuda/hello.cu @@ -0,0 +1,11 @@ +extern "C" { + +__constant__ char s[] = "Hello!"; + + +__global__ void hello(char *out) { + int i = threadIdx.x; + out[i] = s[i]; +} + +} \ No newline at end of file diff --git a/c++/contourgl/cudacontext.cpp b/c++/contourgl/cudacontext.cpp new file mode 100644 index 0000000..3901098 --- /dev/null +++ b/c++/contourgl/cudacontext.cpp @@ -0,0 +1,134 @@ +/* + ......... 2018 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 "cudacontext.h" + + +using namespace std; + + +CudaParams::CudaParams(): params_buffer_size() { + params_extra.push_back(CU_LAUNCH_PARAM_BUFFER_POINTER); + params_extra.push_back(NULL); + params_extra.push_back(CU_LAUNCH_PARAM_BUFFER_SIZE); + params_extra.push_back(¶ms_buffer_size); + params_extra.push_back(CU_LAUNCH_PARAM_END); +} + +void CudaParams::reset() { + params_buffer.clear(); + params_offsets.clear(); + params_pointers.clear(); + params_extra.clear(); + params_buffer_size = 0; +} + +CudaParams& CudaParams::add(const void* data, int size, int align) { + assert(align > 0); + + int index = params_buffer.empty() ? 0 : ((params_buffer.size() - 1)/align + 1)*align; + params_buffer.resize(index + size); + memcpy(¶ms_buffer[index], data, size); + params_buffer_size = params_buffer.size(); + + params_offsets.push_back(index); + + char *root = ¶ms_buffer.front(); + params_pointers.push_back(root + index); + if (params_pointers.front() != root) { + params_pointers.clear(); + for(std::vector::iterator i = params_offsets.begin(); i != params_offsets.end(); ++i) + params_pointers.push_back(root + *i); + } + params_extra[1] = root; + + return *this; +} + + +CudaContext::CudaContext(): + device(), + context(), + err() +{ + const int device_index = 0; + + err = cuInit(0); + assert(!err); + + err = cuDeviceGet(&device, device_index); + assert(!err); + + char device_name[1024] = {}; + err = cuDeviceGetName(device_name, sizeof(device_name), device); + assert(!err); + //cout << "CUDA device " << device_index << ": " << device_name << endl; + + err = cuCtxCreate(&context, CU_CTX_SCHED_AUTO, device); + assert(!err); + + hello(); +} + +CudaContext::~CudaContext() { + cuCtxDestroy(context); +} + +void CudaContext::hello() { + CUmodule module; + err = cuModuleLoad(&module, "cuda/hello.ptx"); + assert(!err); + + CUfunction kernel; + err = cuModuleGetFunction(&kernel, module, "hello"); + assert(!err); + + char data[] = "......"; + + CUdeviceptr buffer; + err = cuMemAlloc(&buffer, sizeof(data)); + + CudaParams params; + params.add(buffer); + + err = cuLaunchKernel( + kernel, + 1, 1, 1, + sizeof(data), 1, 1, + 0, 0, 0, + params.get_extra() ); + assert(!err); + + err = cuStreamSynchronize(0); + assert(!err); + + err = cuMemcpyDtoH(data, buffer, sizeof(data)); + assert(!err); + + err = cuMemFree(buffer); + assert(!err); + + err = cuModuleUnload(module); + assert(!err); + + cout << data << endl; +} diff --git a/c++/contourgl/cudacontext.h b/c++/contourgl/cudacontext.h new file mode 100644 index 0000000..f30f7a2 --- /dev/null +++ b/c++/contourgl/cudacontext.h @@ -0,0 +1,68 @@ +/* + ......... 2018 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 . +*/ + +#ifndef _CUDACONTEXT_H_ +#define _CUDACONTEXT_H_ + +#include + +#include + + +class CudaParams { +private: + std::vector params_buffer; + std::vector params_offsets; + std::vector params_pointers; + std::vector params_extra; + size_t params_buffer_size; + +public: + CudaParams(); + + void reset(); + + CudaParams& add(const void* data, int size, int align); + + template + CudaParams& add(const T &data, int align) + { return add(&data, sizeof(data), align); } + + template + CudaParams& add(const T &data) + { return add(data, __alignof(T)); } + + void** get_params() const + { return params_pointers.empty() ? NULL : &(const_cast(this)->params_pointers.front()); } + void** get_extra() const + { return params_extra.empty() ? NULL : &(const_cast(this)->params_extra.front()); } +}; + + +class CudaContext { +public: + CUdevice device; + CUcontext context; + CUresult err; + + CudaContext(); + ~CudaContext(); + + void hello(); +}; + +#endif diff --git a/c++/contourgl/environment.h b/c++/contourgl/environment.h index dd8b0e0..a9b3783 100644 --- a/c++/contourgl/environment.h +++ b/c++/contourgl/environment.h @@ -22,12 +22,20 @@ #include "clcontext.h" #include "shaders.h" +#ifdef CUDA +#include "cudacontext.h" +#endif + class Environment { public: GlContext gl; ClContext cl; Shaders shaders; + #ifdef CUDA + CudaContext cu; + #endif + Environment(int width, int height, bool hdr, bool multisample, int samples): gl(width, height, hdr, multisample, samples) { } void use() { gl.use(); }