summaryrefslogtreecommitdiff
path: root/Canvas/gpu/CalculateCanvas.cu
diff options
context:
space:
mode:
authorAndrew <saintruler@gmail.com>2019-05-05 22:53:12 +0400
committerAndrew <saintruler@gmail.com>2019-05-05 22:53:12 +0400
commit3a94328b3cff255ac64a208ac98919cc3e530aa0 (patch)
tree377ab395a13e95b5c12a442645c2eaa4b2a644f8 /Canvas/gpu/CalculateCanvas.cu
parent15ccb946f43283bfc91d80ddf012d9ae2a0333ed (diff)
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
Diffstat (limited to 'Canvas/gpu/CalculateCanvas.cu')
-rw-r--r--Canvas/gpu/CalculateCanvas.cu52
1 files changed, 52 insertions, 0 deletions
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<<<numBlocks, threadsPerBlock>>>(gpuCanvas, width, gpuPoints, nPoints);
+
+ cudaMemcpy(canvas, gpuCanvas, canvasSize, cudaMemcpyDeviceToHost);
+
+ cudaFree(gpuCanvas);
+ cudaFree(gpuPoints);
+}