Skip to content

Instantly share code, notes, and snippets.

@ssttv

ssttv/main.cu

Last active February 27, 2019 11:09
Show Gist options

  • Select an option

    Learn more about clone URLs
  • Save ssttv/0eebc689c9238f25bfc37cfbb20bda01 to your computer and use it in GitHub Desktop.

Select an option

Learn more about clone URLs
Save ssttv/0eebc689c9238f25bfc37cfbb20bda01 to your computer and use it in GitHub Desktop.
Download ZIP
Source code for a particle fountain made with OpenGL & CUDA
Raw
main.cu
#include <stdlib.h>
#include <math.h>
#include "particle.h"
#include <iostream>
#include <algorithm>
float g_cameraPhi = 0;
float g_cameraPsi = 0;
float g_cameraR = 10;
const float g_cameraRMin = 1;
const float g_cameraRMax = 30;
const float g_xSource = 3;
const float g_ySource = -2;
const float g_zSource = 0;
int g_xOld = INT_MAX;
int g_yOld = INT_MAX;
bool g_lmbPressed = false;
ULONGLONG g_time = 0;
float g_angle = 0;
const float PI = atan(1.0f) * 4;
const float g_tAver = 1.0f;
const float g_tErr = 0.3f;
const int g_N = 900000;
int g_m = 4;
Particle *g_particles = (Particle *) malloc(2 * g_N * sizeof(Particle));
float *g_times = (float *) malloc(2 * g_N * sizeof(float));
int g_numParticles = 0;
Particle *g_particlesGPU;
float *g_timesGPU;
Point *g_normalsGPU;
Point *g_centersGPU;
Point *g_tang1GPU;
Point *g_tang2GPU;
float *g_wGPU;
float *g_hGPU;
GLuint g_vertices;
struct cudaGraphicsResource* g_verticesCuda;
bool greater(float a, float b)
{
return a > b;
}
void renderScene(void)
{
float t = 0.0f;
ULONGLONG timeCur = GetTickCount64();
if (g_time == 0)
g_time = timeCur;
t = (timeCur - g_time) / 1000.0f;
if (t == 0)
{
return;
}
std::cout << (int)(1000.0 / (timeCur - g_time)) << std::endl;
glClearColor(1.0, 1.0, 1.0, 1.0);
glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
glEnable(GL_DEPTH_TEST);
glMatrixMode(GL_PROJECTION);
glLoadIdentity();
gluPerspective(45.0f, 4.0f / 3, 0.1f, 100.0f);
glMatrixMode(GL_MODELVIEW);
glLoadIdentity();
gluLookAt(g_cameraR * cos(g_cameraPsi) * sin(g_cameraPhi), g_cameraR * sin(g_cameraPsi), g_cameraR * cos(g_cameraPsi) * cos(g_cameraPhi),
-g_cameraR * cos(g_cameraPsi) * sin(g_cameraPhi), -g_cameraR * sin(g_cameraPsi), -g_cameraR * cos(g_cameraPsi) * cos(g_cameraPhi),
0, 1, 0);
g_time = timeCur;
g_angle += PI / 4 * t;
if (g_angle > 2 * PI)
{
g_angle -= 2 * PI;
}
glPushMatrix();
glRotatef(-g_angle * 180 / PI, 0, 1.0, 0);
glBegin(GL_QUADS);
glColor3f(0.0, 1.0, 0.0); // Set The Color To Green
glVertex3f(1.0, 1.0, -1.0); // Top Right Of The Quad (Top)
glVertex3f(-1.0, 1.0, -1.0); // Top Left Of The Quad (Top)
glVertex3f(-1.0, 1.0, 1.0); // Bottom Left Of The Quad (Top)
glVertex3f(1.0, 1.0, 1.0); // Bottom Right Of The Quad (Top)
glEnd();
glBegin(GL_QUADS);
glColor3f(1.0, 0.5, 0.0); // Set The Color To Orange
glVertex3f(1.0, -1.0, 1.0); // Top Right Of The Quad (Bottom)
glVertex3f(-1.0, -1.0, 1.0); // Top Left Of The Quad (Bottom)
glVertex3f(-1.0, -1.0, -1.0); // Bottom Left Of The Quad (Bottom)
glVertex3f(1.0, -1.0, -1.0); // Bottom Right Of The Quad (Bottom)
glEnd();
glBegin(GL_QUADS);
glColor3f(1.0, 0.0, 0.0); // Set The Color To Red
glVertex3f(1.0, 1.0, 1.0); // Top Right Of The Quad (Front)
glVertex3f(-1.0, 1.0, 1.0); // Top Left Of The Quad (Front)
glVertex3f(-1.0, -1.0, 1.0); // Bottom Left Of The Quad (Front)
glVertex3f(1.0, -1.0, 1.0); // Bottom Right Of The Quad (Front)
glEnd();
glBegin(GL_QUADS);
glColor3f(1.0, 1.0, 0.0); // Set The Color To Yellow
glVertex3f(1.0, -1.0, -1.0); // Bottom Left Of The Quad (Back)
glVertex3f(-1.0, -1.0, -1.0); // Bottom Right Of The Quad (Back)
glVertex3f(-1.0, 1.0, -1.0); // Top Right Of The Quad (Back)
glVertex3f(1.0, 1.0, -1.0); // Top Left Of The Quad (Back)
glEnd();
glBegin(GL_QUADS);
glColor3f(0.0, 0.0, 1.0); // Set The Color To Blue
glVertex3f(-1.0, 1.0, 1.0); // Top Right Of The Quad (Left)
glVertex3f(-1.0, 1.0, -1.0); // Top Left Of The Quad (Left)
glVertex3f(-1.0, -1.0, -1.0); // Bottom Left Of The Quad (Left)
glVertex3f(-1.0, -1.0, 1.0); // Bottom Right Of The Quad (Left)
glEnd();
glBegin(GL_QUADS);
glColor3f(1.0, 0.0, 1.0); // Set The Color To Violet
glVertex3f(1.0, 1.0, -1.0); // Top Right Of The Quad (Right)
glVertex3f(1.0, 1.0, 1.0); // Top Left Of The Quad (Right)
glVertex3f(1.0, -1.0, 1.0); // Bottom Left Of The Quad (Right)
glVertex3f(1.0, -1.0, -1.0); // Bottom Right Of The Quad (Right)
glEnd();
glPopMatrix();
Point normals[4] = { Point(cos(g_angle), 0, sin(g_angle)), Point(-cos(g_angle), 0, -sin(g_angle)), Point(-sin(g_angle), 0, cos(g_angle)), Point(sin(g_angle), 0, -cos(g_angle)) };
Point centers[4] = { Point(cos(g_angle), 0, sin(g_angle)), Point(-cos(g_angle), 0, -sin(g_angle)), Point(-sin(g_angle), 0, cos(g_angle)), Point(sin(g_angle), 0, -cos(g_angle)) };
Point tang1[4] = { Point(-sin(g_angle), 0, cos(g_angle)), Point(sin(g_angle), 0, -cos(g_angle)), Point(cos(g_angle), 0, sin(g_angle)), Point(-cos(g_angle), 0, -sin(g_angle)) };
Point tang2[4] = { Point(0, 1, 0), Point(0, 1, 0), Point(0, 1, 0), Point(0, 1, 0) };
float w[4] = { 1, 1, 1, 1 };
float h[4] = { 1, 1, 1, 1 };
int numAdd = (int)(g_N * t / g_tAver);
if ((float) rand() / RAND_MAX < g_N * t / g_tAver - numAdd)
{
++numAdd;
}
for (int i = 0; i < numAdd; ++i)
{
float tCur = g_tAver + ((float)rand() / RAND_MAX * 2 - 1) * g_tErr;
g_times[g_numParticles] = tCur;
float vx = -(3 + (float) rand() / RAND_MAX) * 1;
float vy = (3 + (float) rand() / RAND_MAX) * 1.8f;
float vz = (-0.5f + (float)rand() / RAND_MAX) * 2;
g_particles[g_numParticles].pos = Point(g_xSource, g_ySource, g_zSource);
g_particles[g_numParticles].v = Point(vx, vy, vz);
++g_numParticles;
}
glColor3f(0.2f, 0.3f, 0.5f);
if (g_numParticles > 0)
{
int blockSize = 32;
int gridSize = 1000;
int numElemsPerThread = (g_numParticles + blockSize * gridSize - 1) / (blockSize * gridSize);
cudaMemcpy(g_particlesGPU, g_particles, g_numParticles * sizeof(Particle), cudaMemcpyHostToDevice);
cudaMemcpy(g_timesGPU, g_times, g_numParticles * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(g_normalsGPU, normals, g_m * sizeof(Point), cudaMemcpyHostToDevice);
cudaMemcpy(g_centersGPU, centers, g_m * sizeof(Point), cudaMemcpyHostToDevice);
cudaMemcpy(g_tang1GPU, tang1, g_m * sizeof(Point), cudaMemcpyHostToDevice);
cudaMemcpy(g_tang2GPU, tang2, g_m * sizeof(Point), cudaMemcpyHostToDevice);
cudaMemcpy(g_wGPU, w, g_m * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(g_hGPU, h, g_m * sizeof(float), cudaMemcpyHostToDevice);
float4* positions;
cudaGraphicsMapResources(1, &g_verticesCuda, 0);
size_t num_bytes;
cudaGraphicsResourceGetMappedPointer((void**)&positions, &num_bytes, g_verticesCuda);
moveAndDrawParticles<<<gridSize, blockSize>>>(g_numParticles, g_particlesGPU, g_timesGPU, g_m, g_normalsGPU, g_centersGPU, g_tang1GPU, g_tang2GPU, g_wGPU, g_hGPU, t, positions, numElemsPerThread);
//std::cout << cudaGetLastError() << std::endl;
cudaDeviceSynchronize();
cudaGraphicsUnmapResources(1, &g_verticesCuda, 0);
glBindBuffer(GL_ARRAY_BUFFER, g_vertices);
glVertexPointer(3, GL_FLOAT, 16, 0);
glEnableClientState(GL_VERTEX_ARRAY);
glDrawArrays(GL_POINTS, 0, g_numParticles);
glDisableClientState(GL_VERTEX_ARRAY);
cudaMemcpy(g_particles, g_particlesGPU, g_numParticles * sizeof(Particle), cudaMemcpyDeviceToHost);
cudaMemcpy(g_times, g_timesGPU, g_numParticles * sizeof(float), cudaMemcpyDeviceToHost);
for (int i = 0; i < g_numParticles; ++i)
{
if (g_times[i] - t < 0)
{
g_particles[i] = g_particles[g_numParticles - 1];
g_times[i] = g_times[g_numParticles - 1];
--g_numParticles;
--i;
}
}
}
glTranslatef(g_xSource, g_ySource, g_zSource);
glColor3f(0.9f, 0.6f, 0.3f);
glutSolidSphere(0.1f, 10, 10);
glutSwapBuffers();
glutPostRedisplay();
}
void mouseButton(int button, int state, int x, int y)
{
if (button == GLUT_LEFT_BUTTON)
{
g_lmbPressed = (state == GLUT_DOWN);
}
g_xOld = INT_MAX;
g_yOld = INT_MAX;
}
void mouseMove(int x, int y)
{
if (!g_lmbPressed)
{
return;
}
if (g_xOld == INT_MAX)
{
g_xOld = x;
g_yOld = y;
}
g_cameraPhi += (float)(x - g_xOld) / 200 * PI;
g_cameraPsi -= (float)(g_yOld - y) / 200 * PI;
g_xOld = x;
g_yOld = y;
if (g_cameraPsi > PI / 2 - 0.001f)
{
g_cameraPsi = PI / 2 - 0.001f;
}
if (g_cameraPsi < -PI / 2 + 0.001f)
{
g_cameraPsi = -PI / 2 + 0.001f;
}
}
void mouseWheel(int wheel, int dir, int x, int y)
{
g_cameraR -= 0.3f * dir;
if (g_cameraR > g_cameraRMax)
{
g_cameraR = g_cameraRMax;
}
if (g_cameraR < g_cameraRMin)
{
g_cameraR = g_cameraRMin;
}
}
void clean()
{
cudaFree(g_particlesGPU);
cudaFree(g_timesGPU);
cudaFree(g_normalsGPU);
cudaFree(g_centersGPU);
cudaFree(g_tang1GPU);
cudaFree(g_tang2GPU);
cudaFree(g_wGPU);
cudaFree(g_hGPU);
free(g_particles);
free(g_times);
cudaGraphicsUnregisterResource(g_verticesCuda);
glDeleteBuffers(1, &g_vertices);
}
int main(int argc, char **argv)
{
cudaSetDevice(0);
glutInit(&argc, argv);
glutInitDisplayMode(GLUT_DEPTH | GLUT_DOUBLE | GLUT_RGBA);
glutInitWindowPosition(300, 100);
glutInitWindowSize(800, 600);
glutCreateWindow("Particles");
if (glewInit() != GLEW_OK)
{
std::cerr << "glewInit() failed!" << std::endl;
return 1;
}
glutDisplayFunc(renderScene);
glutIdleFunc(renderScene);
glutMouseFunc(mouseButton);
glutMotionFunc(mouseMove);
glutMouseWheelFunc(mouseWheel);
glutCloseFunc(clean);
glGenBuffers(1, &g_vertices);
glBindBuffer(GL_ARRAY_BUFFER, g_vertices);
unsigned int size = 2 * g_N * 4 * sizeof(float);
glBufferData(GL_ARRAY_BUFFER, size, 0, GL_DYNAMIC_DRAW);
glBindBuffer(GL_ARRAY_BUFFER, 0);
cudaGraphicsGLRegisterBuffer(&g_verticesCuda, g_vertices, cudaGraphicsMapFlagsWriteDiscard);
cudaMalloc(&g_particlesGPU, 2 * g_N * sizeof(Particle));
cudaMalloc(&g_timesGPU, 2 * g_N * sizeof(float));
cudaMalloc(&g_normalsGPU, g_m * sizeof(Point));
cudaMalloc(&g_centersGPU, g_m * sizeof(Point));
cudaMalloc(&g_tang1GPU, g_m * sizeof(Point));
cudaMalloc(&g_tang2GPU, g_m * sizeof(Point));
cudaMalloc(&g_wGPU, g_m * sizeof(float));
cudaMalloc(&g_hGPU, g_m * sizeof(float));
glEnable(GL_DEPTH_TEST);
glutMainLoop();
return 1;
}
Raw
particle.cu
#include "particle.h"
__device__ const float gy = -9.8f;
__device__ float dot(const Point &p1, const Point &p2)
{
return p1.x * p2.x + p1.y * p2.y + p1.z * p2.z;
}
__device__ float minusDot(const Point &p11, const Point &p12, const Point &p2)
{
return (p11.x - p12.x) * p2.x + (p11.y - p12.y) * p2.y + (p11.z - p12.z) * p2.z;
}
__device__ void checkCollision(Particle *p, const int m, const Point *normals, const Point *centers, const Point *tang1, const Point *tang2, const float *w, const float *h, const float dt)
{
int iMax = -1;
float distMax = 1;
float projMax = 1;
float sMax = 0;
for (int i = 0; i < m; ++i)
{
float dist = minusDot(centers[i], p->pos, normals[i]);
float proj = dot(p->v, normals[i]);
float s = -proj / dist;
if (dist > 0 && dist < 0.1f && s > sMax)
{
iMax = i;
sMax = s;
distMax = dist;
projMax = proj;
}
}
if (iMax >= 0)
{
float dist1 = minusDot(centers[iMax], p->pos, tang1[iMax]);
if (dist1 < -w[iMax] || dist1 > w[iMax])
return;
float dist2 = minusDot(centers[iMax], p->pos, tang2[iMax]);
if (dist2 < -h[iMax] || dist2 > h[iMax])
return;
p->v.x -= 2 * projMax * normals[iMax].x;
p->v.y -= 2 * projMax * normals[iMax].y;
p->v.z -= 2 * projMax * normals[iMax].z;
p->pos.x += 2 * distMax * normals[iMax].x;
p->pos.y += 2 * distMax * normals[iMax].y;
p->pos.z += 2 * distMax * normals[iMax].z;
}
}
__global__ void moveAndDrawParticles(const int n, Particle *particles, float *g_times, const int m, const Point *normals, const Point *centers, const Point *tang1, const Point *tang2,
const float *w, const float *h, const float dt, float4 *positions, const int numElemsPerThread)
{
int i = blockIdx.x * blockDim.x * numElemsPerThread + threadIdx.x;
for (int k = 0; k < numElemsPerThread; ++k, i += blockDim.x)
{
if (i > n)
{
return;
}
Particle *p = particles + i;
p->pos.x += p->v.x * dt;
p->pos.y += p->v.y * dt;
p->pos.z += p->v.z * dt;
checkCollision(p, m, normals, centers, tang1, tang2, w, h, dt);
p->v.y += gy * dt;
g_times[i] -= dt;
positions[i] = make_float4(p->pos.x, p->pos.y, p->pos.z, 0);
}
}
Raw
particle.h
#pragma once
#ifdef __CUDACC__
#define CUDA_CALLABLE __host__ __device__
#else
#define CUDA_CALLABLE
#endif
#include <GL/glew.h>
#include <GL/freeglut.h>
#include <cuda_runtime.h>
#include <cuda_gl_interop.h>
struct Point
{
float x;
float y;
float z;
CUDA_CALLABLE Point(float _x, float _y, float _z) : x(_x), y(_y), z(_z) {}
};
struct Particle
{
Point pos;
Point v;
CUDA_CALLABLE Particle(Point _pos, Point _v) : pos(_pos), v(_v) {}
};
__global__ void moveAndDrawParticles(int n, Particle *particles, float *g_times, int m, const Point *normals, const Point *centers, const Point *tang1, const Point *tang2,
const float *w, const float *h, float dt, float4 *positions, int numElemsPerThread);
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment