diff --git a/Makefile b/Makefile index 42560db..4693eff 100644 --- a/Makefile +++ b/Makefile @@ -1,12 +1,12 @@ CXX=g++ HEADERDIR=./header INCLUDE := $(shell find $(HEADERDIR) -type d | sed -e 's/^/-I/' | tr '\n' ' ' | sed 's/.$$//') -CXXFLAGS=-g $(INCLUDE) +CXXFLAGS=-g -std=c++17 $(INCLUDE) SRCDIR=./src SRC := $(shell find $(SRCDIR) -name '*.cpp' | tr '\n' ' ' | sed 's/.$$//') OBJDIR=./build/debug/obj OBJ := $(SRC:$(SRCDIR)/%.cpp=$(OBJDIR)/%.o) -LIBS=-lSDL2 +LIBS=-lSDL2 -lOpenCL EXECUTABLE=slime_mold all: debug diff --git a/gpu_kernel.clcpp b/gpu_kernel.clcpp new file mode 100644 index 0000000..75b9883 --- /dev/null +++ b/gpu_kernel.clcpp @@ -0,0 +1,135 @@ +struct Agent +{ + float x; + float y; + float direction; +}; + +static struct Agent *agents; +static float *ground; +static float *working_ground; +static int2 ground_dimensions; + +uint hash(uint input) +{ + input ^= 2747636419u; + input *= 2654435769u; + input ^= input >> 16; + input *= 2654435769u; + input ^= input >> 16; + input *= 2654435769u; + return input; +} + +void kernel initialize_globals(global struct Agent *new_agents, global float *new_ground, global float *new_working_ground, global int* new_ground_dimensions) +{ + agents = new_agents; + ground = new_ground; + working_ground = new_working_ground; + ground_dimensions[0] = new_ground_dimensions[0]; + ground_dimensions[1] = new_ground_dimensions[1]; +} + +void kernel update_ground() +{ + int current_index = get_global_id(0); + int cx = current_index % ground_dimensions[0]; + int cy = current_index / ground_dimensions[0]; + + float sum = 0.0f; + + for (int i = cx-1; i <= cx+1; i++) + { + if (i < 0 || i >= ground_dimensions[0]) + { + continue; + } + + for(int j = cy-1; j <= cy+1; j++) + { + if (j < 0 || j >= ground_dimensions[1]) + { + continue; + } + + sum += ground[i+j*ground_dimensions[0]]; + } + } + + working_ground[current_index] = sum/9.0f-0.005f; + + if (working_ground[current_index] < 0) + { + working_ground[current_index] = 0.0f; + } +} + +void kernel iterate_ground() +{ + ground[get_global_id(0)] = working_ground[get_global_id(0)]; +} + +void kernel move_agents() +{ + int current_index = get_global_id(0); + + uint psuedorandom_int = hash(agents[current_index].x+agents[current_index].y*ground_dimensions[0]); + + if (psuedorandom_int % 100 < 12) + { + psuedorandom_int = hash(psuedorandom_int); + agents[current_index].direction += (float)(psuedorandom_int%1000)*M_PI_F/3.0f/1000.0f-M_PI_F/6.0f; + } + else + { + float lsense = ground[(int)(agents[current_index].x+2.0f*cos(agents[current_index].direction+M_PI_F/6.0f)) + (int)(agents[current_index].y+2.0f*sin(agents[current_index].direction+M_PI_F/6.0f)) * ground_dimensions[0]]; + float ssense = ground[(int)(agents[current_index].x+2.0f*cos(agents[current_index].direction)) + (int)(agents[current_index].y+2.0f*sin(agents[current_index].direction)) * ground_dimensions[0]]; + float rsense = ground[(int)(agents[current_index].x+2.0f*cos(agents[current_index].direction-M_PI_F/6.0f)) + (int)(agents[current_index].y+2.0f*sin(agents[current_index].direction-M_PI_F/6.0f)) * ground_dimensions[0]]; + + if (lsense > ssense && lsense > rsense) + { + agents[current_index].direction += M_PI_F/6.0f; + } + + if (rsense > ssense && rsense > lsense) + { + agents[current_index].direction += -M_PI_F/6.0f; + } + } + + float nx = agents[current_index].x + cos(agents[current_index].direction); + float ny = agents[current_index].y + sin(agents[current_index].direction); + + if(nx < 0.0f || nx >= ground_dimensions[0] || ny < 0.0f || ny >= ground_dimensions[1]) + { + nx = (nx<0.0f)?0.0f:((nx>=ground_dimensions[0])?(ground_dimensions[0]-1.0f):nx); + ny = (ny<0.0f)?0.0f:((ny>=ground_dimensions[1])?(ground_dimensions[1]-1.0f):ny); + psuedorandom_int = hash(psuedorandom_int); + agents[current_index].direction = (float)(psuedorandom_int%1000)*2.0f*M_PI_F/1000.0f; + } + + agents[current_index].x = nx; + agents[current_index].y = ny; +} + +void kernel emit_agent_pheromones() +{ + int current_index = get_global_id(0); + int x = agents[current_index].x; + int y = agents[current_index].y; + int ground_index = x+y*ground_dimensions[0]; + + ground[ground_index] = 1.0f; +} + +void kernel fill_draw_data(global uchar* draw_data) +{ + int index = get_global_id(0); + + uchar val = (uchar)(255.0f*ground[index]); + + draw_data[4*index] = val; + draw_data[4*index+1] = val; + draw_data[4*index+2] = val; + draw_data[4*index+3] = 255; +} \ No newline at end of file diff --git a/header/Agent.hpp b/header/Agent.hpp deleted file mode 100644 index 0834bcd..0000000 --- a/header/Agent.hpp +++ /dev/null @@ -1,15 +0,0 @@ -#pragma once - -#include "Field.hpp" - -class Agent -{ -public: - Agent(float fl_x, float fl_y, float fl_w, float fl_h, float direction, float fl_pheromone_strength, float fl_random_influence, float fl_lookahead_distance); - ~Agent(); - - void tick(Field* field); - -private: - float fl_x, fl_y, fl_w, fl_h, fl_direction, fl_pheromone_strength, fl_random_influence, fl_lookahead_distance; -}; \ No newline at end of file diff --git a/header/Field.hpp b/header/Field.hpp deleted file mode 100644 index e54c5e5..0000000 --- a/header/Field.hpp +++ /dev/null @@ -1,26 +0,0 @@ -#pragma once - -#include - -class Field -{ -public: - Field(int i_width, int i_height, float fl_pheromone_max); - ~Field(); - - void tick(); - - void add_pheromone(float fl_x, float fl_y, float fl_strength); - - float get_pheromone_strength(float fl_x, float fl_y); - float get_pheromone_strength(int i_x, int i_y); - - float get_pheromone_max(); - -private: - int i_width, i_height; - - float** fl_ground; - - float fl_pheromone_max; -}; \ No newline at end of file diff --git a/large_circle_simulation.gif b/large_circle_simulation.gif new file mode 100644 index 0000000..61eccaf Binary files /dev/null and b/large_circle_simulation.gif differ diff --git a/src/Agent.cpp b/src/Agent.cpp deleted file mode 100644 index 6655496..0000000 --- a/src/Agent.cpp +++ /dev/null @@ -1,89 +0,0 @@ -#include "Agent.hpp" - -#include - -Agent::Agent(float fl_x, float fl_y, float fl_w, float fl_h, float fl_direction, float fl_pheromone_strength, float fl_random_influence, float fl_lookahead_distance) -{ - this->fl_x = fl_x; - this->fl_y = fl_y; - this->fl_w = fl_w; - this->fl_h = fl_h; - this->fl_direction = fl_direction; - this->fl_pheromone_strength = fl_pheromone_strength; - this->fl_random_influence = 100.0f*fl_random_influence; - this->fl_lookahead_distance = fl_lookahead_distance; -} - -Agent::~Agent() -{ - -} - -void Agent::tick(Field* field) -{ - ///Update direction - float fl_xd, fl_yd; - - if ((float)(rand()%100) < fl_random_influence) - { //Move randomly - switch (rand() % 3) - { - case 0: - fl_direction += M_PI_4; - break; - case 1: - break; - case 2: - fl_direction -= M_PI_4; - break; - } - } - else - { //Move based on environment - - //Gather information - float fl_l_pheromone, fl_s_pheromone, fl_r_pheromone, fl_test_dir = fl_direction; - fl_xd = fl_lookahead_distance*cos(fl_test_dir); - fl_yd = fl_lookahead_distance*sin(fl_test_dir); - - fl_s_pheromone = field->get_pheromone_strength(fl_x+fl_xd, fl_y+fl_yd); - - fl_test_dir += M_PI/6.0F; - fl_xd = fl_lookahead_distance*cos(fl_test_dir); - fl_yd = fl_lookahead_distance*sin(fl_test_dir); - fl_l_pheromone = field->get_pheromone_strength(fl_x+fl_xd, fl_y+fl_yd); - - fl_test_dir -= M_PI/3.0F; - fl_xd = fl_lookahead_distance*cos(fl_test_dir); - fl_yd = fl_lookahead_distance*sin(fl_test_dir); - fl_r_pheromone = field->get_pheromone_strength(fl_x+fl_xd, fl_y+fl_yd); - - if (fl_r_pheromone >= fl_s_pheromone && fl_r_pheromone >= fl_l_pheromone) - { - fl_direction -= M_PI/6.0; - } - else if (fl_l_pheromone >= fl_s_pheromone && fl_l_pheromone >= fl_r_pheromone) - { - fl_direction += M_PI/6.0; - } - } - - //Move - fl_xd = cos(fl_direction); - fl_yd = sin(fl_direction); - - if (fl_x + fl_xd < 0 || fl_x + fl_xd >= fl_w || fl_y + fl_yd < 0 || fl_y + fl_yd >= fl_h) - { - fl_x = fmin(fl_w,fmax(0.0f, fl_x+fl_xd)); - fl_y = fmin(fl_h,fmax(0.0f, fl_y+fl_yd)); - fl_direction = 2.0f*M_PI*((float)(rand()%100)/100.0f); - } - else - { - fl_x += fl_xd; - fl_y += fl_yd; - } - - //Add pheromone to new position - field->add_pheromone(fl_x, fl_y, fl_pheromone_strength); -} \ No newline at end of file diff --git a/src/Field.cpp b/src/Field.cpp deleted file mode 100644 index 735d173..0000000 --- a/src/Field.cpp +++ /dev/null @@ -1,118 +0,0 @@ -#include "Field.hpp" - -#include - -Field::Field(int i_width, int i_height, float fl_pheromone_max) -{ - this->i_width = i_width; - this->i_height = i_height; - - fl_ground = new float*[i_width]; - - for (int i = 0; i < i_width; i++) - { - fl_ground[i] = new float[i_height]; - - for(int j = 0; j < i_height; j++) - { - fl_ground[i][j] = 0.0f; - } - } - - this->fl_pheromone_max = fl_pheromone_max; -} - -Field::~Field() -{ - for(int i = 0; i < i_width; i++) - { - delete[] fl_ground[i]; - } - - delete[] fl_ground; -} - -void Field::tick() -{ - float** fl_temp_ground; - fl_temp_ground = new float*[i_width]; - - for (int i = 0; i < i_width; i++) - { - fl_temp_ground[i] = new float[i_height]; - } - - - for(int x = 0; x < i_width; x++) - { - for(int y = 0; y < i_height; y++) - { - - float fl_sum = 0.0f; - float fl_count = 0.0f; - - for (int i = x-1; i <= x+1; i++) - { - for (int j = y-1; j <= y+1; j++) - { - if (i < 0 || i >= i_width || j < 0 || j >= i_height) - { - continue; - } - - fl_sum += fl_ground[i][j]; - fl_count += 1.0f; - } - } - - fl_temp_ground[x][y] = fmax(0.0f, fl_sum/fl_count-0.005); - } - } - - for(int i = 0; i < i_width; i++) - { - memcpy(&fl_ground[i][0], &fl_temp_ground[i][0], sizeof(float)*i_height); - delete[] fl_temp_ground[i]; - } - - delete[] fl_temp_ground; -} - -void Field::add_pheromone(float fl_x, float fl_y, float fl_strength) -{ - int i_x = (int)fl_x; - int i_y = (int)fl_y; - - if (i_x < 0 || i_x >= i_width || i_y < 0 || i_y >= i_height) - { - return; - } - - fl_ground[i_x][i_y] = fmin(fl_ground[i_x][i_y]+fl_strength, fl_pheromone_max); -} - -float Field::get_pheromone_strength(float fl_x, float fl_y) -{ - int i_x = (int)fl_x, i_y = (int)fl_y; - if (i_x < 0 || i_x >= i_width || i_y < 0 || i_y >= i_height) - { - return 0.0f; - } - - return fl_ground[i_x][i_y]; -} - -float Field::get_pheromone_strength(int i_x, int i_y) -{ - if (i_x < 0 || i_x >= i_width || i_y < 0 || i_y >= i_height) - { - return 0.0f; - } - - return fl_ground[i_x][i_y]; -} - -float Field::get_pheromone_max() -{ - return fl_pheromone_max; -} \ No newline at end of file diff --git a/src/main.cpp b/src/main.cpp index 4b38ce9..81aad11 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -1,92 +1,221 @@ #include +#define CL_HPP_TARGET_OPENCL_VERSION 200 +#include + #include #include #include #include #include #include +#include +#include -#include "Field.hpp" -#include "Agent.hpp" - -unsigned long long SYS_TIME_US() +struct Agent { - return (unsigned long long)(std::chrono::duration_cast(std::chrono::system_clock::now().time_since_epoch()).count()); -} + float x,y,direction; +}; + +bool get_source(std::string* source); + +unsigned long long sys_time_us(); int main(int argc, char* argv[]) { srand(time(0)); - int width, height; + uint32_t u32_width, u32_height; + uint32_t u32_dimensions[2]; if (argc == 1) { - width = height = 300; + u32_width = u32_height = 300; } else { - width = atoi(argv[1]); - height = atoi(argv[2]); + u32_width = atoi(argv[1]); + u32_height = atoi(argv[2]); } + int i_size = u32_width*u32_height; + + u32_dimensions[0] = u32_width; + u32_dimensions[1] = u32_height; + if (SDL_Init(SDL_INIT_VIDEO) < 0) { std::cout << "Video Init Error." << std::endl; } - SDL_Window* window = SDL_CreateWindow("Slime Mold Simulator", SDL_WINDOWPOS_CENTERED-width/2, SDL_WINDOWPOS_CENTERED-height/2, width, height, SDL_WINDOW_SHOWN); + SDL_Window* window = SDL_CreateWindow("Slime Mold Simulator", SDL_WINDOWPOS_CENTERED-u32_width/2, SDL_WINDOWPOS_CENTERED-u32_height/2, u32_width, u32_height, SDL_WINDOW_SHOWN); SDL_Renderer* renderer = SDL_CreateRenderer(window, -1, SDL_RENDERER_ACCELERATED | SDL_RENDERER_TARGETTEXTURE); + SDL_RenderClear(renderer); + SDL_RenderPresent(renderer); - SDL_Texture* field_texture = SDL_CreateTexture(renderer, SDL_GetWindowPixelFormat(window), SDL_TEXTUREACCESS_TARGET, width, height); + SDL_Texture* field_texture = SDL_CreateTexture(renderer, SDL_GetWindowPixelFormat(window), SDL_TEXTUREACCESS_STREAMING, u32_width, u32_height); - Field field(width, height, 1.0f); + SDL_Texture* start_texture = SDL_CreateTexture(renderer, SDL_GetWindowPixelFormat(window), SDL_TEXTUREACCESS_STREAMING, u32_width, u32_height); - std::vector agents; - - bool* start_positions = new bool[width*height]; - for(int i = 0; i < width*height; i++) + float* ground = new float[i_size]; + for(int i = 0; i < i_size; i++) { - start_positions[i] = false; + ground[i] = 0.0f; } - int created = 0; - int index = 0; - int max_count = (int)sqrt(width*height)*10; - while(created < max_count) + int i_agent_count = (int)sqrt(u32_width*u32_height)*1000; + Agent* agents = new Agent[i_agent_count]; + { - int rng_cur = rand()%(width*height/max_count); - if (rng_cur == 0 && !start_positions[index]) + uint8_t* u8_pixels; + int i_pitch; + SDL_LockTexture(start_texture, NULL, (void**)&u8_pixels,&i_pitch); + + int i_created = 0; + uint32_t u32_radius = 4*((u32_width all_platforms; + cl::Platform::get(&all_platforms); - int iteration_duration = 16666/10; - int max_loops = 10000000/iteration_duration; + cl::Platform default_platform=all_platforms[0]; + + std::cout << "Using platform :" << default_platform.getInfo() << std::endl; + + std::vector all_devices; + default_platform.getDevices(CL_DEVICE_TYPE_ALL, &all_devices); + + if (all_devices.size() == 0) + { + std::cout << "Could not find OpenCL device." << std::endl; + exit(1); + } + + cl::Device default_device = all_devices[0]; + + std::cout << "Using device :" << default_device.getInfo() << std::endl; + + cl::Context context({default_device}); + + std::string kernel_code; + if (!get_source(&kernel_code)) + { + std::cout << "Error reading kernel source." << std::endl; + exit(1); + } + + cl::Program::Sources sources; + sources.push_back({kernel_code.c_str(), kernel_code.length()}); + cl::Program program(context, sources); + + if (program.build({default_device}, "-cl-std=CL2.0") != CL_SUCCESS) + { + std::cout << "Error building: " << program.getBuildInfo(default_device) << std::endl; + exit(1); + } + + cl::Buffer buffer_agents(context, CL_MEM_READ_WRITE, sizeof(Agent)*i_agent_count); + cl::Buffer buffer_ground(context, CL_MEM_READ_WRITE, sizeof(float)*i_size); + cl::Buffer buffer_ground_swap(context, CL_MEM_READ_WRITE, sizeof(float)*i_size); + cl::Buffer buffer_ground_dimensions(context, CL_MEM_READ_WRITE, sizeof(uint32_t)*2); + cl::Buffer buffer_draw_data(context, CL_MEM_READ_WRITE, sizeof(uint8_t)*4*i_size); + + cl::CommandQueue queue(context,default_device); + + queue.enqueueWriteBuffer(buffer_agents, CL_TRUE, 0, sizeof(Agent)*i_agent_count, agents); + queue.enqueueWriteBuffer(buffer_ground, CL_TRUE, 0, sizeof(float)*i_size, ground); + queue.enqueueWriteBuffer(buffer_ground_swap, CL_TRUE, 0, sizeof(float)*i_size, ground); + queue.enqueueWriteBuffer(buffer_ground_dimensions, CL_TRUE, 0, sizeof(uint32_t)*2, &u32_dimensions[0]); + + cl::Kernel kernel_initialize_globals(program, "initialize_globals"); + kernel_initialize_globals.setArg(0, buffer_agents); + kernel_initialize_globals.setArg(1, buffer_ground); + kernel_initialize_globals.setArg(2, buffer_ground_swap); + kernel_initialize_globals.setArg(3, buffer_ground_dimensions); + + queue.enqueueNDRangeKernel(kernel_initialize_globals, cl::NullRange, cl::NDRange(1), cl::NullRange); + queue.finish(); + + cl::Kernel kernel_update_ground(program, "update_ground"); + cl::Kernel kernel_iterate_ground(program, "iterate_ground"); + cl::Kernel kernel_move_agents(program, "move_agents"); + cl::Kernel kernel_emit_agent_pheromones(program, "emit_agent_pheromones"); + cl::Kernel kernel_fill_draw_data(program, "fill_draw_data"); + + queue.enqueueNDRangeKernel(kernel_emit_agent_pheromones, cl::NullRange, cl::NDRange(i_agent_count), cl::NullRange); + queue.enqueueReadBuffer(buffer_ground, CL_TRUE, 0, sizeof(float)*i_size, ground); + queue.finish();//*/ + + SDL_Rect screen_rect = {0,0,(int)u32_width,(int)u32_height}; + + unsigned long long ull_time_old, ull_time_now; + ull_time_now = ull_time_old = sys_time_us(); + + int iteration_duration = 16666; bool pause = true; while(true) @@ -104,44 +233,46 @@ int main(int argc, char* argv[]) { pause = !pause; } + else if (event.key.keysym.scancode == SDL_SCANCODE_T) + { + SDL_SetRenderDrawColor(renderer, 0,0,0,0); + SDL_RenderClear(renderer); + + SDL_RenderCopy(renderer, start_texture, &screen_rect, &screen_rect); + + SDL_RenderPresent(renderer); + } } } - time_now = SYS_TIME_US(); - if (time_now >= time_old + iteration_duration) + ull_time_now = sys_time_us(); + if (ull_time_now >= ull_time_old + iteration_duration) { - time_old = time_now; + ull_time_old = ull_time_now; if (pause) { continue; } - field.tick(); - for(unsigned int i = 0; i < agents.size(); i++) - { - agents[i]->tick(&field); - } + + queue.enqueueNDRangeKernel(kernel_update_ground, cl::NullRange, cl::NDRange(i_size), cl::NullRange); + queue.enqueueNDRangeKernel(kernel_iterate_ground, cl::NullRange, cl::NDRange(i_size), cl::NullRange); + + queue.enqueueNDRangeKernel(kernel_move_agents, cl::NullRange, cl::NDRange(i_agent_count), cl::NullRange); + queue.enqueueNDRangeKernel(kernel_emit_agent_pheromones, cl::NullRange, cl::NDRange(i_agent_count), cl::NullRange); - SDL_SetRenderTarget(renderer, field_texture); - SDL_SetRenderDrawColor(renderer, 0,0,0,0); - SDL_RenderClear(renderer); + kernel_fill_draw_data.setArg(0, buffer_draw_data); + queue.enqueueNDRangeKernel(kernel_fill_draw_data, cl::NullRange, cl::NDRange(i_size), cl::NullRange); - float max_strength = field.get_pheromone_max(); + uint8_t* u8_pixels; + int i_pitch; + SDL_LockTexture(field_texture, NULL, (void**)&u8_pixels, &i_pitch); - for (int x = 0; x < width; x++) - { - for (int y = 0; y < height; y++) - { - float cur_strength = field.get_pheromone_strength(x, y); - - int color_val = std::max(0,std::min(255,(int)(255.0f * cur_strength / max_strength))); - SDL_SetRenderDrawColor(renderer, color_val, color_val, color_val, 0xFF); + queue.enqueueReadBuffer(buffer_draw_data, CL_TRUE, 0, sizeof(uint8_t)*4*i_size, u8_pixels); + queue.finish(); - SDL_RenderDrawPoint(renderer, x, y); - } - } + SDL_UnlockTexture(field_texture); - SDL_SetRenderTarget(renderer, NULL); SDL_SetRenderDrawColor(renderer, 0,0,0,0); SDL_RenderClear(renderer); @@ -156,10 +287,30 @@ int main(int argc, char* argv[]) SDL_Quit(); - for (unsigned int i = 0; i < agents.size(); i++) - { - delete agents[i]; - } + delete[] agents; + delete[] ground; return 0; +} + +bool get_source(std::string* source) +{ + if (source == nullptr) + { + return false; + } + + std::ifstream in_file("gpu_kernel.clcpp", std::ios::in); + std::stringstream buffer; + + buffer << in_file.rdbuf(); + + *source = buffer.str(); + + return true; +} + +unsigned long long sys_time_us() +{ + return (unsigned long long)(std::chrono::duration_cast(std::chrono::system_clock::now().time_since_epoch()).count()); } \ No newline at end of file