From 3a94328b3cff255ac64a208ac98919cc3e530aa0 Mon Sep 17 00:00:00 2001 From: Andrew Date: Sun, 5 May 2019 22:53:12 +0400 Subject: Made changes to CMakeLists.txt so CUDA version of program builds when CUDA is installed. Replaced C-style approach to generating texture to OOP-style with Canvas.h --- Canvas/Canvas.cpp | 98 +++++++++++++++++++++++++++++++++++++++ Canvas/Canvas.h | 97 +++++--------------------------------- Canvas/cpu/CalculateCanvas.cpp | 60 ++++++++++++++++++++++++ Canvas/cpu/CalculateCanvas.h | 15 ++++++ Canvas/cuda/CalculateCanvasGPU.cu | 52 --------------------- Canvas/cuda/CalculateCanvasGPU.h | 20 -------- Canvas/gpu/CalculateCanvas.cu | 52 +++++++++++++++++++++ Canvas/gpu/CalculateCanvas.h | 16 +++++++ 8 files changed, 252 insertions(+), 158 deletions(-) create mode 100644 Canvas/Canvas.cpp create mode 100644 Canvas/cpu/CalculateCanvas.cpp create mode 100644 Canvas/cpu/CalculateCanvas.h delete mode 100644 Canvas/cuda/CalculateCanvasGPU.cu delete mode 100644 Canvas/cuda/CalculateCanvasGPU.h create mode 100644 Canvas/gpu/CalculateCanvas.cu create mode 100644 Canvas/gpu/CalculateCanvas.h (limited to 'Canvas') diff --git a/Canvas/Canvas.cpp b/Canvas/Canvas.cpp new file mode 100644 index 0000000..6034633 --- /dev/null +++ b/Canvas/Canvas.cpp @@ -0,0 +1,98 @@ +// +// Created by saintruler on 05.05.19. +// +#include "Canvas.h" + +Canvas::Canvas(int width, int height, Point *points, int nPoints) +{ + float vertices[] = { + // positions // colors // texture coords + 1.f, 1.f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f, // top right + 1.f, -1.f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f, // bottom right + -1.f, -1.f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, // bottom left + -1.f, 1.f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f // top left + }; + + int indices[] = { + 0, 1, 3, // first triangle + 1, 2, 3 // second triangle + }; + + glGenVertexArrays(1, &VAO); + glGenBuffers(1, &VBO); + glGenBuffers(1, &EBO); + + glBindVertexArray(VAO); + + glBindBuffer(GL_ARRAY_BUFFER, VBO); + glBufferData(GL_ARRAY_BUFFER, sizeof(vertices), vertices, GL_STATIC_DRAW); + + glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, EBO); + glBufferData(GL_ELEMENT_ARRAY_BUFFER, sizeof(indices), indices, GL_STATIC_DRAW); + + // position attribute + glVertexAttribPointer(0, 3, GL_FLOAT, GL_FALSE, 8 * sizeof(float), (void*)0); + glEnableVertexAttribArray(0); + // color attribute + glVertexAttribPointer(1, 3, GL_FLOAT, GL_FALSE, 8 * sizeof(float), (void*)(3 * sizeof(float))); + glEnableVertexAttribArray(1); + // texture coord attribute + glVertexAttribPointer(2, 2, GL_FLOAT, GL_FALSE, 8 * sizeof(float), (void*)(6 * sizeof(float))); + glEnableVertexAttribArray(2); + + this->width = width; + this->height = height; + this->nPoints = nPoints; + + this->points = points; + canvas = (unsigned char*) malloc(width * height * 3); + glGenTextures(1, &texture); +} + +void Canvas::SetPixel(int x, int y, Color color) +{ + canvas[y * width * 3 + x * 3 + 0] = color.r; + canvas[y * width * 3 + x * 3 + 1] = color.g; + canvas[y * width * 3 + x * 3 + 2] = color.b; + + UpdateTexture(); +} + +void Canvas::UpdateTexture() +{ + // all upcoming GL_TEXTURE_2D operations now have effect on this texture object + glBindTexture(GL_TEXTURE_2D, texture); + + // set texture wrapping to GL_REPEAT (default wrapping method) + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_REPEAT); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_REPEAT); + + // set texture filtering parameters + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_LINEAR); + + generate_canvas(width, height, canvas, points, nPoints); + + glTexImage2D(GL_TEXTURE_2D, 0, GL_RGB, width, width, 0, GL_RGB, GL_UNSIGNED_BYTE, canvas); + glGenerateMipmap(GL_TEXTURE_2D); +} + +void Canvas::DrawTexture() +{ + UpdateTexture(); + + glBindTexture(GL_TEXTURE_2D, texture); + shader.use(); + + glBindVertexArray(VAO); + glDrawElements(GL_TRIANGLES, 6, GL_UNSIGNED_INT, 0); +} + +void Canvas::DeleteCanvas() +{ + glDeleteVertexArrays(1, &VAO); + glDeleteBuffers(1, &VBO); + glDeleteBuffers(1, &EBO); + + delete[] canvas; +} diff --git a/Canvas/Canvas.h b/Canvas/Canvas.h index 84dbe56..ee6c844 100644 --- a/Canvas/Canvas.h +++ b/Canvas/Canvas.h @@ -7,6 +7,9 @@ #include #include +#include +#include +#include "gpu/CalculateCanvas.h" #include @@ -18,100 +21,22 @@ struct Color class Canvas { public: - Canvas(int width, int height) - { -// vertices = { -// // positions // colors // texture coords -// 1.f, 1.f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f, // top right -// 1.f, -1.f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f, // bottom right -// -1.f, -1.f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, // bottom left -// -1.f, 1.f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f // top left -// }; -// -// indices = { -// 0, 1, 3, // first triangle -// 1, 2, 3 // second triangle -// }; - - glGenVertexArrays(1, &VAO); - glGenBuffers(1, &VBO); - glGenBuffers(1, &EBO); - - glBindVertexArray(VAO); - - glBindBuffer(GL_ARRAY_BUFFER, VBO); - glBufferData(GL_ARRAY_BUFFER, sizeof(vertices), vertices, GL_STATIC_DRAW); - - glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, EBO); - glBufferData(GL_ELEMENT_ARRAY_BUFFER, sizeof(indices), indices, GL_STATIC_DRAW); - - // position attribute - glVertexAttribPointer(0, 3, GL_FLOAT, GL_FALSE, 8 * sizeof(float), (void*)0); - glEnableVertexAttribArray(0); - // color attribute - glVertexAttribPointer(1, 3, GL_FLOAT, GL_FALSE, 8 * sizeof(float), (void*)(3 * sizeof(float))); - glEnableVertexAttribArray(1); - // texture coord attribute - glVertexAttribPointer(2, 2, GL_FLOAT, GL_FALSE, 8 * sizeof(float), (void*)(6 * sizeof(float))); - glEnableVertexAttribArray(2); - - - this->width = width; - this->height = height; - - canvas = (unsigned char*) malloc(width * height * 3); - glGenTextures(1, &texture); - } - - void SetPixel(int x, int y, Color color) - { - canvas[y * width * 3 + x * 3 + 0] = color.r; - canvas[y * width * 3 + x * 3 + 1] = color.g; - canvas[y * width * 3 + x * 3 + 2] = color.b; - - UpdateTexture(); - } - - void UpdateTexture() - { - // all upcoming GL_TEXTURE_2D operations now have effect on this texture object - glBindTexture(GL_TEXTURE_2D, texture); - - // set texture wrapping to GL_REPEAT (default wrapping method) - glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_REPEAT); - glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_REPEAT); - - // set texture filtering parameters - glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR); - glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_LINEAR); - - glTexImage2D(GL_TEXTURE_2D, 0, GL_RGB, width, width, 0, GL_RGB, GL_UNSIGNED_BYTE, canvas); - glGenerateMipmap(GL_TEXTURE_2D); - } - - void DrawTexture() - { - glBindTexture(GL_TEXTURE_2D, texture); - // render container -// ourShader.use(); - glBindVertexArray(VAO); - glDrawElements(GL_TRIANGLES, 6, GL_UNSIGNED_INT, 0); - } + Canvas(int width, int height, Point* points, int nPoints); - unsigned int* GetTexture() - { - return &(this->texture); - } + void SetPixel(int x, int y, Color color); + void UpdateTexture(); + void DrawTexture(); + void DeleteCanvas(); private: - int width, height; + int width, height, nPoints; unsigned char* canvas; unsigned int texture; + Point* points; unsigned int VBO, VAO, EBO; - float* vertices; - unsigned int* indices; + Shader shader = Shader("4.1.texture.vs", "4.1.texture.fs"); }; #endif //OPENGLTEST_CANVAS_H diff --git a/Canvas/cpu/CalculateCanvas.cpp b/Canvas/cpu/CalculateCanvas.cpp new file mode 100644 index 0000000..75a9996 --- /dev/null +++ b/Canvas/cpu/CalculateCanvas.cpp @@ -0,0 +1,60 @@ +// +// Created by saintruler on 05.05.19. +// +#include "CalculateCanvas.h" + +void calculate(unsigned char* canvas, int width, Point* points, int nPoints, int canvasX, int canvasY) +{ + int pathX, pathY; + int length_squared; + double f = 0; + + for (int i = 0; i < nPoints; i++) + { + pathX = points[i].x - canvasX; + pathY = points[i].y - canvasY; + + length_squared = pathX * pathX + pathY * pathY; + f += (1e5 * 20) / (float) length_squared; + } + + int l = (int) f; + if (l > 255) l = 255; + + canvas[canvasY * width * 3 + canvasX * 3 + 0] = l; + canvas[canvasY * width * 3 + canvasX * 3 + 1] = 0; + canvas[canvasY * width * 3 + canvasX * 3 + 2] = 0; +} + +//void calculate(unsigned char* canvas, int width, Point* points, int nPoints, int canvasX, int canvasY) +// point canvas_pt(canvasX, canvasY); +// double f = 0; +// for (point pt : points) +// { +// point path = pt - canvas_pt; +// +// f += (1e5 * 20 * 1) / path.length_sqared(); +// } +// +// int l = (int) f; +// if (l > 255) l = 255; +// +// +// canvas[y * SCR_WIDTH * 3 + x * 3 + 0] = l; +// canvas[y * SCR_WIDTH * 3 + x * 3 + 1] = 0; +// canvas[y * SCR_WIDTH * 3 + x * 3 + 2] = 0; +//} + + + +void generate_canvas(int width, int height, unsigned char* canvas, Point* points, int nPoints) +{ + for (int x = 0; x < width; x++) + { + for (int y = 0; y < height; y++) + { + calculate(canvas, width, points, nPoints, x, y); + } + } +} + diff --git a/Canvas/cpu/CalculateCanvas.h b/Canvas/cpu/CalculateCanvas.h new file mode 100644 index 0000000..c1a1662 --- /dev/null +++ b/Canvas/cpu/CalculateCanvas.h @@ -0,0 +1,15 @@ +// +// Created by saintruler on 05.05.19. +// + +#ifndef OPENGLTEST_CALCULATECANVAS_H +#define OPENGLTEST_CALCULATECANVAS_H + +#include + + +void calculate(unsigned char* canvas, int width, Point* points, int nPoints, int canvasX, int canvasY); + +void generate_canvas(int width, int height, unsigned char* canvas, Point* points, int nPoints); + +#endif //OPENGLTEST_CALCULATECANVAS_H diff --git a/Canvas/cuda/CalculateCanvasGPU.cu b/Canvas/cuda/CalculateCanvasGPU.cu deleted file mode 100644 index c8c6fad..0000000 --- a/Canvas/cuda/CalculateCanvasGPU.cu +++ /dev/null @@ -1,52 +0,0 @@ -#include "CalculateCanvasGPU.h" - -__global__ -void calculate(unsigned char* canvas, int width, cuda_point* points, int nPoints) -{ - int canvasX = blockIdx.x * blockDim.x + threadIdx.x; - int canvasY = blockIdx.y * blockDim.y + threadIdx.y; - - int pathX, pathY; - int length_squared; - double f = 0; - - for (int i = 0; i < nPoints; i++) - { - pathX = points[i].x - canvasX; - pathY = points[i].y - canvasY; - - length_squared = pathX * pathX + pathY * pathY; - f += (1e5 * 20) / (float) length_squared; - } - - int l = (int) f; - if (l > 255) l = 255; - - canvas[canvasY * width * 3 + canvasX * 3 + 0] = l; - canvas[canvasY * width * 3 + canvasX * 3 + 1] = 0; - canvas[canvasY * width * 3 + canvasX * 3 + 2] = 0; -} - -void generate_canvas(int width, int height, unsigned char* canvas, cuda_point* points, int nPoints) -{ - const int canvasSize = width * height * 3 * sizeof(unsigned char); - const int pointsSize = nPoints * sizeof(cuda_point); - - unsigned char* gpuCanvas; - cuda_point* gpuPoints; - - cudaMalloc((void**) &gpuCanvas, canvasSize); - cudaMalloc((void**) &gpuPoints, pointsSize); - - cudaMemcpy(gpuCanvas, canvas, canvasSize, cudaMemcpyHostToDevice); - cudaMemcpy(gpuPoints, points, pointsSize, cudaMemcpyHostToDevice); - - dim3 threadsPerBlock(16, 16); - dim3 numBlocks(width / threadsPerBlock.x, height / threadsPerBlock.y); - calculate<<>>(gpuCanvas, width, gpuPoints, nPoints); - - cudaMemcpy(canvas, gpuCanvas, canvasSize, cudaMemcpyDeviceToHost); - - cudaFree(gpuCanvas); - cudaFree(gpuPoints); -} diff --git a/Canvas/cuda/CalculateCanvasGPU.h b/Canvas/cuda/CalculateCanvasGPU.h deleted file mode 100644 index 46f54c5..0000000 --- a/Canvas/cuda/CalculateCanvasGPU.h +++ /dev/null @@ -1,20 +0,0 @@ -// -// Created by saintruler on 05.05.19. -// - -#ifndef OPENGLTEST_CALCULATECANVASGPU_H -#define OPENGLTEST_CALCULATECANVASGPU_H - -#include -#include -#include - -typedef struct -{ - int x, y; -} cuda_point; - - -void generate_canvas(int width, int height, unsigned char* canvas, cuda_point* points, int nPoints); - -#endif //OPENGLTEST_CALCULATECANVASGPU_H diff --git a/Canvas/gpu/CalculateCanvas.cu b/Canvas/gpu/CalculateCanvas.cu new file mode 100644 index 0000000..3a73067 --- /dev/null +++ b/Canvas/gpu/CalculateCanvas.cu @@ -0,0 +1,52 @@ +#include "CalculateCanvas.h" + +__global__ +void calculate(unsigned char* canvas, int width, Point* points, int nPoints) +{ + int canvasX = blockIdx.x * blockDim.x + threadIdx.x; + int canvasY = blockIdx.y * blockDim.y + threadIdx.y; + + int pathX, pathY; + int length_squared; + double f = 0; + + for (int i = 0; i < nPoints; i++) + { + pathX = points[i].x - canvasX; + pathY = points[i].y - canvasY; + + length_squared = pathX * pathX + pathY * pathY; + f += (1e5 * 20) / (float) length_squared; + } + + int l = (int) f; + if (l > 255) l = 255; + + canvas[canvasY * width * 3 + canvasX * 3 + 0] = l; + canvas[canvasY * width * 3 + canvasX * 3 + 1] = 0; + canvas[canvasY * width * 3 + canvasX * 3 + 2] = 0; +} + +void generate_canvas(int width, int height, unsigned char* canvas, Point* points, int nPoints) +{ + const int canvasSize = width * height * 3 * sizeof(unsigned char); + const int pointsSize = nPoints * sizeof(Point); + + unsigned char* gpuCanvas; + Point* gpuPoints; + + cudaMalloc((void**) &gpuCanvas, canvasSize); + cudaMalloc((void**) &gpuPoints, pointsSize); + + cudaMemcpy(gpuCanvas, canvas, canvasSize, cudaMemcpyHostToDevice); + cudaMemcpy(gpuPoints, points, pointsSize, cudaMemcpyHostToDevice); + + dim3 threadsPerBlock(16, 16); + dim3 numBlocks(width / threadsPerBlock.x, height / threadsPerBlock.y); + calculate<<>>(gpuCanvas, width, gpuPoints, nPoints); + + cudaMemcpy(canvas, gpuCanvas, canvasSize, cudaMemcpyDeviceToHost); + + cudaFree(gpuCanvas); + cudaFree(gpuPoints); +} diff --git a/Canvas/gpu/CalculateCanvas.h b/Canvas/gpu/CalculateCanvas.h new file mode 100644 index 0000000..0de853f --- /dev/null +++ b/Canvas/gpu/CalculateCanvas.h @@ -0,0 +1,16 @@ +// +// Created by saintruler on 05.05.19. +// + +#ifndef OPENGLTEST_CALCULATECANVAS_H +#define OPENGLTEST_CALCULATECANVAS_H + +#include +#include +#include + +#include + +void generate_canvas(int width, int height, unsigned char* canvas, Point* points, int nPoints); + +#endif //OPENGLTEST_CALCULATECANVAS_H -- cgit v1.2.3