#include <GL/glew.h>
#include <GLFW/glfw3.h>
#include <cuda_runtime.h>
#include <cuda_gl_interop.h>
#include <cstdio>
#include <omp.h>
#include <getopt.h>
#define _USE_MATH_DEFINES
#include <math.h>
#include <stdlib.h>
#include "naive_update.h"
#include "utils.h"
#include "main.h"
#include "naive_binning.cuh"
#include "shared_binning.cuh"
#define STARTING_WINDOW_HEIGHT 1080
#define STARTING_WINDOW_WIDTH 1920
#define PROJ_BUF_SIZ (MAX_PROJS * 2 * sizeof(Projectile))
#define SNAPSHOT_START MAX_PROJS
#define SNAPSHOT_END MAX_PROJS * 2
#define MAX_TRACKED_PROJS 50
#define NAIVE_GPU 0
#define NAIVE_BINNING 1
#define SHARED_BINNING 2
/* nvcc .\main.cu .\naive_binning.cu .\naive_update.cu .\shared_binning.cu
* .\utils.cu -o demo -I "C:\vcpkg\installed\x64-windows\include" -I
* "C:\vcpkg\installed\x64-windows\include\opencv4" -L
* "C:\vcpkg\installed\x64-windows\lib" -lglew32 -lglfw3dll -lopengl32
* -lopencv_core4 -lopencv_imgcodecs4 -lopencv_imgproc4 -lgetopt
*/
/* proj sim config values */
static int projectiles_count = 1;
static float arc_gap = 0.0f;
static float fire_rate = 10.0f;
static float frequency = 0.0f;
static float amplitude = 0.0f;
static float speed = 240.0f;
static float max_lifetime = 240.0f;
// static float acceleration = 0;
static float gravity = 0;
/* testing flags */
static int use_cpu = 0;
static int test_mode = 0;
static uint64_t num_proj_spawned = 0;
static int kernel = SHARED_BINNING;
static int show_bin_lines = 0;
static int threads_per_block = 128;
static int BIN_SIZE = 128;
/* 16:9 display by default, could make this scale via the window resize event
* handler but like, that is *way* out of scope for this project at this point,
* maybe worth doing in an application with users */
static int BIN_DIM_X = 32;
static int BIN_DIM_Y = 18;
/* host-side projectile array.
* exec(2) zeros out these arrays as they go in the .BSS section
* of the virtual memory adress space; no need to explicitly zero them */
static Projectile h_projs[MAX_PROJS];
static int tracked_projs[MAX_TRACKED_PROJS];
static int tracked_count = 0;
/* the size of the window and/or frame buffer can be changed during runtime
* by the user which would massively distort our simulation.
* this a function called by the event handler that fires whenever
* glfw detects that the window has been resized */
void
framebuffer_resize_handler(GLFWwindow* window, int width, int height)
{
glViewport(0, 0, width, height);
glMatrixMode(GL_PROJECTION);
glLoadIdentity();
glOrtho(0, width, height, 0, -1, 1);
glMatrixMode(GL_MODELVIEW);
glLoadIdentity();
/* as great as this is, it doesn't re-position pre-existing projectiles...
* i'll get around to this right after never. */
}
int
main(int argc, char** argv)
{
/* assuming a device is present...*/
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
for (int i = 1; i < argc;)
{
if (strcmp(argv[i], "-cpu") == 0)
{
use_cpu = 1;
for (int j = i; j < argc - 1; j++)
{
argv[j] = argv[j + 1];
}
argc--;
continue;
}
if (strcmp(argv[i], "-benchmark") == 0 || strcmp(argv[i], "-test") == 0)
{
test_mode = 1;
for (int j = i; j < argc - 1; j++)
{
argv[j] = argv[j + 1];
}
argc--;
continue;
}
if (strcmp(argv[i], "-showbinlines") == 0)
{
show_bin_lines = 1;
for (int j = i; j < argc - 1; j++)
{
argv[j] = argv[j + 1];
}
argc--;
continue;
}
i++;
}
/* default testing values; can be overriden by explicit cli args */
if (test_mode)
{
projectiles_count = 72;
speed = 240;
amplitude = 1.5f;
frequency = 0.05f;
fire_rate = 5.0f * 11;
arc_gap = 5.0f;
gravity = 240.0f;
max_lifetime = 2000;
}
int opt;
int kernel_flag_entered = 0;
int custom_bin_flag_entered = 0;
while ((opt = getopt(argc, argv, "G:c:g:f:a:s:r:l:h:B:X:Y:K:t:C:")) != -1)
{
switch (opt)
{
case 'c':
projectiles_count = atoi(optarg);
break;
case 'g':
arc_gap = atof(optarg);
break;
case 'f':
frequency = atof(optarg);
break;
case 'a':
amplitude = atof(optarg);
break;
case 's':
speed = atof(optarg);
break;
case 'r':
fire_rate = atof(optarg);
break;
case 'l':
max_lifetime = atof(optarg);
break;
// case 'A' : acceleration = atof(optarg); break;
case 'G':
gravity = atof(optarg);
break;
case 'B':
BIN_SIZE = atoi(optarg);
custom_bin_flag_entered = 1;
break;
case 'X':
BIN_DIM_X = atoi(optarg);
custom_bin_flag_entered = 1;
break;
case 'Y':
BIN_DIM_Y = atoi(optarg);
custom_bin_flag_entered = 1;
break;
case 'K':
kernel = atoi(optarg);
kernel_flag_entered = 1;
break;
case 't':
threads_per_block = atoi(optarg);
if (threads_per_block > prop.maxThreadsPerBlock ||
threads_per_block < 1)
{
fprintf(stderr,
"Error: threads per block cannot exceed %d for "
"the current device (nor can be less than 1!)\n",
prop.maxThreadsPerBlock);
return EXIT_FAILURE;
}
break;
case 'h':
default:
printf("Usage: %s [options]\n", argv[0]);
printf(" -c <int> projectile count (default: 1)\n");
printf(
" -g <float> arc gap in degrees (default: 5 if count > 1)\n");
printf(" -f <float> wave frequency (default: 0)\n");
printf(" -a <float> wave amplitude (default: 0)\n");
printf(" -s <float> projectile speed (default: 240)\n");
printf(" -r <float> fire rate (default: 10)\n");
printf(" -l <float> lifetime (default: 240)\n");
// printf(" -A <float> acceleration (default: 0)\n");
printf(" -G <float> gravity factor (default: 0)\n");
printf(" -B <int> bin size (default: 64)\n");
printf(" -X <int> bin dimension X (default: 32)\n");
printf(" -Y <int> bin dimension Y (default: 18)\n");
printf(" -K <int> kernel (0: naive gpu, 1: naive binning, 2: "
"shared binning) (default: 2)\n");
printf(" -t <int> threads per block (default: 128)\n");
printf(" -h show this\n");
printf("Additional arguments:\n");
printf(" -cpu run the simulation on the CPU (default: "
"not enabled)\n");
printf(" -<benchmark | test> run in benchmark mode with "
"pre-configured "
"default values (default: not enabled)\n");
printf(
" -showbinlines show the bin lines (default: not enabled)\n");
return opt == 'h' ? EXIT_SUCCESS : EXIT_FAILURE;
}
}
if (kernel_flag_entered && use_cpu)
{
fprintf(stderr, "Error: cannot specify GPU kernel when in CPU mode\n");
return EXIT_FAILURE;
}
if (show_bin_lines && (kernel == NAIVE_GPU || use_cpu))
{
fprintf(stderr, "Error: cannot show bin lines when not using a binning "
"kernel on the GPU\n");
return EXIT_FAILURE;
}
if (kernel == NAIVE_GPU && custom_bin_flag_entered)
{
fprintf(stderr, "Error: cannot specify custom binning parameters when "
"using the naive GPU kernel\n");
return EXIT_FAILURE;
}
const size_t shared_mem = 9 * BIN_SIZE * sizeof(float4);
/* sizeof(int) is statically allocated in the gpu kernel itself */
if (kernel == SHARED_BINNING &&
shared_mem + (sizeof(int) * 20) >= prop.sharedMemPerBlock)
{
fprintf(stderr,
"Error: shared memory usage per block (%zu bytes) "
"exceeds the maximum available (%zu bytes) for "
"this kernel configuration. Try reducing the "
"bin size.\n",
(shared_mem + (sizeof(int) * 20)), prop.sharedMemPerBlock);
return EXIT_FAILURE;
}
if (projectiles_count > 1 && arc_gap == 0.0f)
{
arc_gap = 5.0f;
}
if (projectiles_count < 1)
{
fprintf(stderr, "Error: projectile count must be >= 1\n");
return EXIT_FAILURE;
}
int iw;
int ih;
uchar4* h_img = load_image_opencv("projectile.png", &iw, &ih);
if (!h_img)
{
fprintf(stderr, "Failed to load image.\n");
return EXIT_FAILURE;
}
glfwInit();
GLFWwindow* window =
glfwCreateWindow(STARTING_WINDOW_WIDTH, STARTING_WINDOW_HEIGHT,
"Simple Projectile Simulator", nullptr, nullptr);
glfwMakeContextCurrent(window);
glewInit();
/* register the glfw event hook for window resize */
int init_width, init_height;
glfwGetFramebufferSize(window, &init_width, &init_height);
framebuffer_resize_handler(window, init_width, init_height);
glfwSetFramebufferSizeCallback(window, framebuffer_resize_handler);
/* FPS cap (in frames to wait for)
* 0 = disable vsync
*/
glfwSwapInterval(0);
/* Load the projectile sprite into OpenGL texture */
GLuint sprite_tex;
glGenTextures(1, &sprite_tex);
glBindTexture(GL_TEXTURE_2D, sprite_tex);
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, iw, ih, 0, GL_RGBA,
GL_UNSIGNED_BYTE, h_img);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_LINEAR);
free(h_img);
const int total_bins = BIN_DIM_X * BIN_DIM_Y;
Projectile* d_projs = NULL;
int* d_bins = NULL;
int* d_bin_counts = NULL;
cudaEvent_t k_start, k_stop;
if (!use_cpu)
{
/* the first half is the real array, the second half is the snapshot */
cudaMalloc(&d_projs, PROJ_BUF_SIZ);
cudaMalloc(&d_bins, total_bins * BIN_SIZE * sizeof(int));
cudaMalloc(&d_bin_counts, total_bins * sizeof(int));
cudaMemset(d_projs, 0, PROJ_BUF_SIZ);
cudaMemset(d_bins, 0, total_bins * BIN_SIZE * sizeof(int));
cudaMemset(d_bin_counts, 0, total_bins * sizeof(int));
set_bin_constants(BIN_SIZE, BIN_DIM_X, BIN_DIM_Y);
if (test_mode)
{
/* register here (and delete at end of program) so they aren't being
* created/destroyed every frame */
cudaEventCreate(&k_start);
cudaEventCreate(&k_stop);
}
}
int active_count = 0;
float last_fire_time = 0.0f;
float fire_interval = 1.0f / fire_rate;
double mouse_x = 0.0;
double mouse_y = 0.0;
float polarity =
0; /* used for alternating the last shot if projectile_count is odd */
float last_frame_time = (float)glfwGetTime();
float last_test_print_time = (float)glfwGetTime();
/* naive GPU dimensions */
dim3 naive_gpu_threads(threads_per_block);
dim3 naive_gpu_blocks(
((MAX_PROJS + naive_gpu_threads.x - 1) / naive_gpu_threads.x));
/* naive binning dimensions */
dim3 naive_binning_threads(threads_per_block);
dim3 naive_binning_blocks(
((MAX_PROJS + naive_binning_threads.x - 1) / naive_binning_threads.x));
/* shared binning dimensions */
dim3 shared_binning_threads(threads_per_block);
dim3 shared_binning_blocks(((MAX_PROJS + shared_binning_threads.x - 1) /
shared_binning_threads.x));
dim3 shared_binning_gravity_threads(BIN_SIZE); /* NOT threads_per_block*/
dim3 shared_binning_gravity_blocks(total_bins);
dim3 shared_binning_velocity_threads(threads_per_block);
dim3 shared_binning_velocity_blocks(
((MAX_PROJS + shared_binning_velocity_threads.x - 1) /
shared_binning_velocity_threads.x));
if (test_mode && !use_cpu)
{
switch (kernel)
{
case NAIVE_GPU:
print_kernel_occupancy("NAIVE GPU", (const void*)update, 0,
naive_gpu_blocks.x);
break;
case NAIVE_BINNING:
print_kernel_occupancy("NAIVE BINNING (bin_projectiles)",
(const void*)bin_projectiles, 0,
naive_binning_blocks.x);
print_kernel_occupancy("NAIVE BINNING", (const void*)update_binned,
0, naive_binning_blocks.x);
break;
default:
case SHARED_BINNING:
print_kernel_occupancy("SHARED BINNING (bin_projectiles)",
(const void*)bin_projectiles, 0,
shared_binning_threads.x);
print_kernel_occupancy(
"SHARED BINNING (gravity)", (const void*)update_binned_shared,
shared_mem, shared_binning_gravity_threads.x);
print_kernel_occupancy("SHARED BINNING (velocity)",
(const void*)update_velocity_shared, 0,
shared_binning_velocity_threads.x);
break;
}
}
while (!glfwWindowShouldClose(window))
{
float now = (float)glfwGetTime();
float delta_t = now - last_frame_time;
last_frame_time = now;
/* log every 1 second */
if (test_mode && (now - last_test_print_time >= 1.0f))
{
fps_tracker_update(1.0f / delta_t, active_count);
if (!use_cpu)
{
kernel_timer_print();
}
last_test_print_time = (float)glfwGetTime();
}
/* low fps guard */
if (delta_t > 1.0f / 60.0f)
{
delta_t = 1.0f / 60.0f;
}
int window_width;
int window_height;
glfwPollEvents();
/* the frame buffer size may not always match the window size
* here we want to check and ensure they are every frame (as the
* user may resize the window, etc.) */
glfwGetFramebufferSize(window, &window_width, &window_height);
float bin_w = (float)window_width / BIN_DIM_X;
float bin_h = (float)window_height / BIN_DIM_Y;
int lmb_pressed = glfwGetMouseButton(window, GLFW_MOUSE_BUTTON_LEFT);
int rmb_pressed = glfwGetMouseButton(window, GLFW_MOUSE_BUTTON_RIGHT);
/* projectile spawn check */
if ((lmb_pressed == GLFW_PRESS || rmb_pressed == GLFW_PRESS ||
test_mode) &&
(now - last_fire_time) >= fire_interval)
{
last_fire_time = now;
// if (!use_cpu)
// {
// cudaMemcpy(h_projs, d_projs, MAX_PROJS * sizeof(Projectile),
// cudaMemcpyDeviceToHost);
// }
static int test_spawn_count = 0;
if (!test_mode)
{
glfwGetCursorPos(window, &mouse_x, &mouse_y);
}
else
{
int cycle = test_spawn_count % 24;
{
if (cycle < 20)
{
mouse_x = window_width / 2.0f;
mouse_y = window_height / 2.0f;
}
else
{
int corner = (cycle - 20);
switch (corner)
{
case 0: // top-left-ish
mouse_x = window_width * 0.25;
mouse_y = window_height * 0.25;
break;
case 1: // bottom-left-ish
mouse_x = window_width * 0.25;
mouse_y = window_height * 0.75;
break;
case 2: // top-right-ish
mouse_x = window_width * 0.75;
mouse_y = window_height * 0.25;
break;
case 3: // bottom-right-ish
mouse_x = window_width * 0.75;
mouse_y = window_height * 0.75;
break;
}
}
test_spawn_count++;
}
}
float dx = (float)mouse_x - window_width / 2.0f;
float dy = (float)mouse_y - window_height / 2.0f;
float center_angle = atan2f(dy, dx) * (180.0f / (float)M_PI);
float total_arc = arc_gap * (projectiles_count - 1);
int successfully_added = 0;
float2 spawn_location;
/* if rmb is clicked and the cursor is within bounds */
if (test_mode)
{
spawn_location = make_float2((float)mouse_x, (float)mouse_y);
}
else if (rmb_pressed && !(mouse_x > window_width || mouse_x < 0 ||
mouse_y > window_height || mouse_y < 0))
{
spawn_location = make_float2(mouse_x, mouse_y);
}
else
{
spawn_location =
make_float2(window_width / 2.0f, window_height / 2.0f);
}
for (int i = 0; i < MAX_PROJS; i++)
{
Projectile* p = &h_projs[i];
if (p->active)
{
continue;
}
memset(p, 0, sizeof(Projectile));
p->type = 0;
p->pos = spawn_location;
p->old_pos = spawn_location;
p->frequency = frequency;
p->amplitude = amplitude;
p->phase = (successfully_added % 2 == 0) ? 0.0f : (float)M_PI;
p->wave_phase = 0;
p->origin = spawn_location;
p->max_lifetime = max_lifetime;
/* mark a projectile for tracking! */
if (num_proj_spawned % MAX_TRACKED_PROJS == 0)
{
p->tracked = 1;
}
float angle_deg = center_angle - (total_arc / 2.0f) +
(arc_gap * successfully_added);
float angle_rad = angle_deg * ((float)M_PI / 180.0f);
p->velocity = make_float2(cosf(angle_rad) * speed,
sinf(angle_rad) * speed);
// h_projs[i].acceleration = acceleration;
p->gravity = gravity;
p->active = 1;
active_count++;
successfully_added++;
num_proj_spawned++;
if (successfully_added == projectiles_count)
{
polarity = (polarity == 0.0f ? (float)M_PI : 0.0f);
p->phase = polarity;
break;
}
}
if (!use_cpu)
{
cudaMemcpy(d_projs, h_projs, MAX_PROJS * sizeof(Projectile),
cudaMemcpyHostToDevice);
}
}
if (active_count > 0)
{
if (use_cpu)
{
update_cpu(h_projs, delta_t);
}
else
{
/* create the pointers for conveinent access while in the
* kernel
*/
Projectile* d_real = d_projs;
Projectile* d_snapshot = d_projs + SNAPSHOT_START;
switch (kernel)
{
case NAIVE_GPU:
{
if (test_mode)
{
cudaEventRecord(k_start);
}
cudaMemcpy(d_snapshot, d_real,
MAX_PROJS * sizeof(Projectile),
cudaMemcpyDeviceToDevice);
update<<<naive_gpu_blocks, naive_gpu_threads>>>(
d_projs, d_snapshot, delta_t);
cudaDeviceSynchronize();
cudaMemcpy(h_projs, d_projs, MAX_PROJS * sizeof(Projectile),
cudaMemcpyDeviceToHost);
if (test_mode)
{
cudaEventRecord(k_stop);
kernel_timer_record(k_start, k_stop);
}
break;
}
case NAIVE_BINNING:
{
cudaMemcpy(d_snapshot, d_real,
MAX_PROJS * sizeof(Projectile),
cudaMemcpyDeviceToDevice);
if (test_mode)
{
cudaEventRecord(k_start);
}
/* clear bin counts from last frame (slide 83) */
cudaMemset(d_bin_counts, 0, total_bins * sizeof(int));
/* assign projectiles to bins */
bin_projectiles<<<naive_binning_blocks,
naive_binning_threads>>>(
d_snapshot, d_bins, d_bin_counts, BIN_SIZE, BIN_DIM_X,
BIN_DIM_Y, bin_w, bin_h);
cudaDeviceSynchronize();
/* update using neighborhood bins only */
update_binned<<<naive_binning_blocks,
naive_binning_threads>>>(
d_real, d_snapshot, d_bins, d_bin_counts, BIN_SIZE,
delta_t, BIN_DIM_X, BIN_DIM_Y, bin_w, bin_h);
cudaDeviceSynchronize();
cudaMemcpy(h_projs, d_projs, MAX_PROJS * sizeof(Projectile),
cudaMemcpyDeviceToHost);
if (test_mode)
{
cudaEventRecord(k_stop);
kernel_timer_record(k_start, k_stop);
}
break;
}
default:
case SHARED_BINNING:
{
if (test_mode)
{
cudaEventRecord(k_start);
}
cudaMemcpy(d_snapshot, d_real,
MAX_PROJS * sizeof(Projectile),
cudaMemcpyDeviceToDevice);
cudaMemset(d_bin_counts, 0, total_bins * sizeof(int));
bin_projectiles<<<shared_binning_blocks,
shared_binning_threads>>>(
d_real, d_bins, d_bin_counts, BIN_SIZE, BIN_DIM_X,
BIN_DIM_Y, bin_w, bin_h);
cudaDeviceSynchronize();
/* apply/search for nearby gravity utilizing shared
* memory (each block corresponds to a bin) */
/* This is already quite a bit of shared memory --
* couldn't use sizeofprojectile, and quite frankly I
* wanted to use bit fields in a dedicated
* mini-projectile struct, but apparently cuda doesn't
* like that / it doesn't align well, so packing it into
* a float4 is the best I could come up with; could be a
* little better, could be far worse */
/* shared memory defined above */
update_binned_shared<<<shared_binning_gravity_blocks,
shared_binning_gravity_threads,
shared_mem>>>(
d_real, d_snapshot, d_bins, d_bin_counts, delta_t);
cudaDeviceSynchronize();
/* basic non-gravity movement update for every
* projectile. If there are too many projectiles in a
* bin; at least these will still move based on their
* last velocity. */
update_velocity_shared<<<shared_binning_velocity_blocks,
shared_binning_velocity_threads>>>(
d_real, delta_t);
cudaDeviceSynchronize();
cudaMemcpy(h_projs, d_projs, MAX_PROJS * sizeof(Projectile),
cudaMemcpyDeviceToHost);
if (test_mode)
{
cudaEventRecord(k_stop);
kernel_timer_record(k_start, k_stop);
}
break;
}
}
}
active_count = 0;
for (int i = 0; i < MAX_PROJS; i++)
{
if (h_projs[i].active)
{
active_count++;
}
}
glClearColor(0.0f, 0.0f, 0.0f, 1.0f);
glClear(GL_COLOR_BUFFER_BIT);
glEnable(GL_TEXTURE_2D);
glBindTexture(GL_TEXTURE_2D, sprite_tex);
glEnable(GL_BLEND);
glBlendFunc(GL_SRC_ALPHA, GL_ONE_MINUS_SRC_ALPHA);
float hw = iw / 2.0f;
float hh = ih / 2.0f;
tracked_count = 0;
for (int i = 0; i < MAX_PROJS; i++)
{
if (!h_projs[i].active)
{
continue;
}
if (h_projs[i].tracked)
{
if (tracked_count < MAX_TRACKED_PROJS)
{
tracked_projs[tracked_count++] = i;
}
continue;
}
glPushMatrix();
glTranslatef(h_projs[i].pos.x, h_projs[i].pos.y, 0.0f);
glRotatef(h_projs[i].rotation, 0.0f, 0.0f, 1.0f);
glColor3f(1.0f, 1.0f, 1.0f);
glBegin(GL_QUADS);
glTexCoord2f(0, 0);
glVertex2f(-hw, -hh);
glTexCoord2f(1, 0);
glVertex2f(hw, -hh);
glTexCoord2f(1, 1);
glVertex2f(hw, hh);
glTexCoord2f(0, 1);
glVertex2f(-hw, hh);
glEnd();
glPopMatrix();
}
/* draw tracked projectiles; they must be drawn last as to be
* drawn over the others */
for (int i = 0; i < tracked_count; i++)
{
Projectile* p = &h_projs[tracked_projs[i]];
glPushMatrix();
glTranslatef(p->pos.x, p->pos.y, 0.0f);
glRotatef(p->rotation, 0.0f, 0.0f, 1.0f);
glColor3f(0.0f, 1.0f, 0.66f);
glBegin(GL_QUADS);
glTexCoord2f(0, 0);
glVertex2f(-hw, -hh);
glTexCoord2f(1, 0);
glVertex2f(hw, -hh);
glTexCoord2f(1, 1);
glVertex2f(hw, hh);
glTexCoord2f(0, 1);
glVertex2f(-hw, hh);
glEnd();
glPopMatrix();
}
/* check above */
if (show_bin_lines /* && kernel != NAIVE_GPU */)
{
draw_bin_grid(window_width, window_height, BIN_DIM_X,
BIN_DIM_Y);
}
glfwSwapBuffers(window);
}
else
{
glClearColor(0.0f, 0.0f, 0.0f, 1.0f);
glClear(GL_COLOR_BUFFER_BIT);
glfwSwapBuffers(window);
/* idle at ~60 fps */
glfwWaitEventsTimeout(1 / 60.0f);
}
}
if (!use_cpu)
{
cudaFree(d_projs);
cudaFree(d_bins);
cudaFree(d_bin_counts);
if (test_mode)
{
cudaEventDestroy(k_start);
cudaEventDestroy(k_stop);
}
}
glDeleteTextures(1, &sprite_tex);
glfwTerminate();
return EXIT_SUCCESS;
}