From 619360f46e8d6af51a9c9710deeadaec7ea8adcb Mon Sep 17 00:00:00 2001 From: Peter Heywood Date: Thu, 20 Apr 2023 17:52:18 +0100 Subject: [PATCH] CMake: Demonstrate use of flamegpu_add_library on the host_fucntions example. Maybe don't merge? --- examples/cpp/host_functions/CMakeLists.txt | 17 ++- examples/cpp/host_functions/src/main.cu | 131 +------------------- examples/cpp/host_functions/src/model.cu | 132 +++++++++++++++++++++ 3 files changed, 146 insertions(+), 134 deletions(-) create mode 100644 examples/cpp/host_functions/src/model.cu diff --git a/examples/cpp/host_functions/CMakeLists.txt b/examples/cpp/host_functions/CMakeLists.txt index 356a372e1..5176b8993 100644 --- a/examples/cpp/host_functions/CMakeLists.txt +++ b/examples/cpp/host_functions/CMakeLists.txt @@ -1,4 +1,4 @@ -# Minimum CMake version 3.18 for CUDA --std=c++17 +# Minimum CMake version 3.18 for CUDA --std=c++17 cmake_minimum_required(VERSION 3.18...3.25 FATAL_ERROR) # Set the location of the ROOT flame gpu project relative to this CMakeList.txt @@ -19,12 +19,19 @@ SET(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/bin/${CMAKE_BUILD_TYPE}/) # Prepare list of source files # Can't do this automatically, as CMake wouldn't know when to regen (as CMakeLists.txt would be unchanged) -SET(ALL_SRC - ${CMAKE_CURRENT_SOURCE_DIR}/src/main.cu +SET(STATIC_SRC + ${CMAKE_CURRENT_SOURCE_DIR}/src/model.cu ) -# Add the executable and set required flags for the target -flamegpu_add_executable("${PROJECT_NAME}" "${ALL_SRC}" "${FLAMEGPU_ROOT}" "${PROJECT_BINARY_DIR}" TRUE) +# define the model at a static library target, demonstrating flamegpu_add_library +flamegpu_add_library("${PROJECT_NAME}_static" "${STATIC_SRC}" "${FLAMEGPU_ROOT}" "${PROJECT_BINARY_DIR}" TRUE) + +# Create the actual executable, and link against the static lib +set(MAIN_SRC + ${CMAKE_CURRENT_SOURCE_DIR}/src/main.cu +) +flamegpu_add_executable("${PROJECT_NAME}" "${MAIN_SRC}" "${FLAMEGPU_ROOT}" "${PROJECT_BINARY_DIR}" TRUE) +target_link_libraries(${PROJECT_NAME} PRIVATE ${PROJECT_NAME}_static) # Also set as startup project (if top level project) set_property(DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}" PROPERTY VS_STARTUP_PROJECT "${PROJECT_NAME}") \ No newline at end of file diff --git a/examples/cpp/host_functions/src/main.cu b/examples/cpp/host_functions/src/main.cu index 5bd2a0c2f..91c2371fb 100644 --- a/examples/cpp/host_functions/src/main.cu +++ b/examples/cpp/host_functions/src/main.cu @@ -1,132 +1,5 @@ -#include "flamegpu/flamegpu.h" - -const unsigned int AGENT_COUNT = 1024; - -FLAMEGPU_AGENT_FUNCTION(device_function, flamegpu::MessageNone, flamegpu::MessageNone) { - const float prop_float = FLAMEGPU->environment.getProperty("float"); - const int16_t prop_int16 = FLAMEGPU->environment.getProperty("int16_t"); - const uint64_t prop_uint64_0 = FLAMEGPU->environment.getProperty("uint64_t", 0); - const uint64_t prop_uint64_1 = FLAMEGPU->environment.getProperty("uint64_t", 1); - const uint64_t prop_uint64_2 = FLAMEGPU->environment.getProperty("uint64_t", 2); - if (blockIdx.x * blockDim.x + threadIdx.x == 0) { - printf("Agent Function[Thread 0]! Properties(Float: %g, int16: %hd, uint64[3]: {%llu, %llu, %llu})\n", prop_float, prop_int16, prop_uint64_0, prop_uint64_1, prop_uint64_2); - } - return flamegpu::ALIVE; -} -FLAMEGPU_INIT_FUNCTION(init_function) { - flamegpu::HostAgentAPI agent = FLAMEGPU->agent("agent"); - float min_x = agent.min("x"); - float max_x = agent.max("x"); - printf("Init Function! (AgentCount: %u, Min: %g, Max: %g)\n", FLAMEGPU->agent("agent").count(), min_x, max_x); - for (unsigned int i = AGENT_COUNT / 2; i < AGENT_COUNT; i++) { - flamegpu::HostNewAgentAPI instance = agent.newAgent(); - instance.setVariable("x", static_cast(i)); - instance.setVariable("a", i % 2 == 0 ? 1 : 0); - } - printf("Init Function! Created %u more agents\n", AGENT_COUNT / 2); -} -FLAMEGPU_CUSTOM_REDUCTION(customSum, a, b) { - return a + b; -} -FLAMEGPU_CUSTOM_TRANSFORM(customTransform, a) { - return (a == 0 || a == 1) ? 1 : 0; -} -FLAMEGPU_STEP_FUNCTION(step_function) { - auto agent = FLAMEGPU->agent("agent"); - int sum_a = agent.sum("a"); - int custom_sum_a = agent.reduce("a", customSum, 0); - unsigned int count_a = agent.count("a", 1); - unsigned int countif_a = agent.transformReduce("a", customTransform, customSum, 0u); - printf("Step Function! (AgentCount: %u, Sum: %d, CustomSum: %d, Count: %u, CustomCountIf: %u)\n", agent.count(), sum_a, custom_sum_a, count_a, countif_a); -} -FLAMEGPU_EXIT_FUNCTION(exit_function) { - float uniform_real = FLAMEGPU->random.uniform(); - int uniform_int = FLAMEGPU->random.uniform(1, 10); - float normal = FLAMEGPU->random.normal(); - float logNormal = FLAMEGPU->random.logNormal(1, 1); - printf("Exit Function! (%g, %i, %g, %g)\n", - uniform_real, uniform_int, normal, logNormal); -} -FLAMEGPU_HOST_FUNCTION(host_function) { - std::vector hist_x = FLAMEGPU->agent("agent").histogramEven("x", 8, -0.5, 1023.5); - printf("Host Function! (Hist: [%u, %u, %u, %u, %u, %u, %u, %u]\n", - hist_x[0], hist_x[1], hist_x[2], hist_x[3], hist_x[4], hist_x[5], hist_x[6], hist_x[7]); - FLAMEGPU->environment.setProperty("int16_t", FLAMEGPU->environment.getProperty("int16_t") + 1); -} -FLAMEGPU_EXIT_CONDITION(exit_condition) { - const float CHANCE = 0.15f; - float uniform_real = FLAMEGPU->random.uniform(); - printf("Exit Condition! (Rolled: %g)\n", uniform_real); - if (uniform_real < CHANCE) { - printf("Rolled number is less than %g, exiting!\n", CHANCE); - return flamegpu::EXIT; - } - return flamegpu::CONTINUE; -} - +int wrapped_main(int argc, const char ** argv); int main(int argc, const char ** argv) { - flamegpu::ModelDescription model("host_functions_example"); - - { // agent - flamegpu::AgentDescription agent = model.newAgent("agent"); - agent.newVariable("x"); - agent.newVariable("a"); - agent.newFunction("device_function", device_function); - } - - /** - * GLOBALS - */ - { - flamegpu::EnvironmentDescription envProperties = model.Environment(); - envProperties.newProperty("float", 12.0f); - envProperties.newProperty("int16_t", 0); - envProperties.newProperty("uint64_t", {11llu, 12llu, 13llu}); - } - /** - * Control flow - */ - { // Attach init/step/exit functions and exit condition - model.addInitFunction(init_function); - model.addStepFunction(step_function); - model.addExitFunction(exit_function); - model.addExitCondition(exit_condition); - } - - { - flamegpu::LayerDescription devicefn_layer = model.newLayer("devicefn_layer"); - devicefn_layer.addAgentFunction(device_function); - } - - { - flamegpu::LayerDescription hostfn_layer = model.newLayer("hostfn_layer"); - hostfn_layer.addHostFunction(host_function); - } - - /** - * Initialisation - */ - flamegpu::AgentVector population(model.Agent("agent"), AGENT_COUNT/2); - for (unsigned int i = 0; i < AGENT_COUNT/2; i++) { - flamegpu::AgentVector::Agent instance = population[i]; - instance.setVariable("x", static_cast(i)); - instance.setVariable("a", i % 2 == 0 ? 1 : 0); - } - - /** - * Execution - */ - flamegpu::CUDASimulation cudaSimulation(model); - cudaSimulation.SimulationConfig().steps = 0; - cudaSimulation.setPopulationData(population); - cudaSimulation.initialise(argc, argv); - cudaSimulation.simulate(); - - cudaSimulation.getPopulationData(population); - - // Ensure profiling / memcheck work correctly - flamegpu::util::cleanup(); - - return 0; + wrapped_main(argc, argv); } diff --git a/examples/cpp/host_functions/src/model.cu b/examples/cpp/host_functions/src/model.cu new file mode 100644 index 000000000..173d6d352 --- /dev/null +++ b/examples/cpp/host_functions/src/model.cu @@ -0,0 +1,132 @@ +#include "flamegpu/flamegpu.h" + +const unsigned int AGENT_COUNT = 1024; + +FLAMEGPU_AGENT_FUNCTION(device_function, flamegpu::MessageNone, flamegpu::MessageNone) { + const float prop_float = FLAMEGPU->environment.getProperty("float"); + const int16_t prop_int16 = FLAMEGPU->environment.getProperty("int16_t"); + const uint64_t prop_uint64_0 = FLAMEGPU->environment.getProperty("uint64_t", 0); + const uint64_t prop_uint64_1 = FLAMEGPU->environment.getProperty("uint64_t", 1); + const uint64_t prop_uint64_2 = FLAMEGPU->environment.getProperty("uint64_t", 2); + if (blockIdx.x * blockDim.x + threadIdx.x == 0) { + printf("Agent Function[Thread 0]! Properties(Float: %g, int16: %hd, uint64[3]: {%llu, %llu, %llu})\n", prop_float, prop_int16, prop_uint64_0, prop_uint64_1, prop_uint64_2); + } + return flamegpu::ALIVE; +} +FLAMEGPU_INIT_FUNCTION(init_function) { + flamegpu::HostAgentAPI agent = FLAMEGPU->agent("agent"); + float min_x = agent.min("x"); + float max_x = agent.max("x"); + printf("Init Function! (AgentCount: %u, Min: %g, Max: %g)\n", FLAMEGPU->agent("agent").count(), min_x, max_x); + for (unsigned int i = AGENT_COUNT / 2; i < AGENT_COUNT; i++) { + flamegpu::HostNewAgentAPI instance = agent.newAgent(); + instance.setVariable("x", static_cast(i)); + instance.setVariable("a", i % 2 == 0 ? 1 : 0); + } + printf("Init Function! Created %u more agents\n", AGENT_COUNT / 2); +} +FLAMEGPU_CUSTOM_REDUCTION(customSum, a, b) { + return a + b; +} +FLAMEGPU_CUSTOM_TRANSFORM(customTransform, a) { + return (a == 0 || a == 1) ? 1 : 0; +} +FLAMEGPU_STEP_FUNCTION(step_function) { + auto agent = FLAMEGPU->agent("agent"); + int sum_a = agent.sum("a"); + int custom_sum_a = agent.reduce("a", customSum, 0); + unsigned int count_a = agent.count("a", 1); + unsigned int countif_a = agent.transformReduce("a", customTransform, customSum, 0u); + printf("Step Function! (AgentCount: %u, Sum: %d, CustomSum: %d, Count: %u, CustomCountIf: %u)\n", agent.count(), sum_a, custom_sum_a, count_a, countif_a); +} +FLAMEGPU_EXIT_FUNCTION(exit_function) { + float uniform_real = FLAMEGPU->random.uniform(); + int uniform_int = FLAMEGPU->random.uniform(1, 10); + float normal = FLAMEGPU->random.normal(); + float logNormal = FLAMEGPU->random.logNormal(1, 1); + printf("Exit Function! (%g, %i, %g, %g)\n", + uniform_real, uniform_int, normal, logNormal); +} +FLAMEGPU_HOST_FUNCTION(host_function) { + std::vector hist_x = FLAMEGPU->agent("agent").histogramEven("x", 8, -0.5, 1023.5); + printf("Host Function! (Hist: [%u, %u, %u, %u, %u, %u, %u, %u]\n", + hist_x[0], hist_x[1], hist_x[2], hist_x[3], hist_x[4], hist_x[5], hist_x[6], hist_x[7]); + FLAMEGPU->environment.setProperty("int16_t", FLAMEGPU->environment.getProperty("int16_t") + 1); +} +FLAMEGPU_EXIT_CONDITION(exit_condition) { + const float CHANCE = 0.15f; + float uniform_real = FLAMEGPU->random.uniform(); + printf("Exit Condition! (Rolled: %g)\n", uniform_real); + if (uniform_real < CHANCE) { + printf("Rolled number is less than %g, exiting!\n", CHANCE); + return flamegpu::EXIT; + } + return flamegpu::CONTINUE; +} + + +int wrapped_main(int argc, const char ** argv) { + flamegpu::ModelDescription model("host_functions_example"); + + { // agent + flamegpu::AgentDescription agent = model.newAgent("agent"); + agent.newVariable("x"); + agent.newVariable("a"); + agent.newFunction("device_function", device_function); + } + + /** + * GLOBALS + */ + { + flamegpu::EnvironmentDescription envProperties = model.Environment(); + envProperties.newProperty("float", 12.0f); + envProperties.newProperty("int16_t", 0); + envProperties.newProperty("uint64_t", {11llu, 12llu, 13llu}); + } + /** + * Control flow + */ + { // Attach init/step/exit functions and exit condition + model.addInitFunction(init_function); + model.addStepFunction(step_function); + model.addExitFunction(exit_function); + model.addExitCondition(exit_condition); + } + + { + flamegpu::LayerDescription devicefn_layer = model.newLayer("devicefn_layer"); + devicefn_layer.addAgentFunction(device_function); + } + + { + flamegpu::LayerDescription hostfn_layer = model.newLayer("hostfn_layer"); + hostfn_layer.addHostFunction(host_function); + } + + /** + * Initialisation + */ + flamegpu::AgentVector population(model.Agent("agent"), AGENT_COUNT/2); + for (unsigned int i = 0; i < AGENT_COUNT/2; i++) { + flamegpu::AgentVector::Agent instance = population[i]; + instance.setVariable("x", static_cast(i)); + instance.setVariable("a", i % 2 == 0 ? 1 : 0); + } + + /** + * Execution + */ + flamegpu::CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 0; + cudaSimulation.setPopulationData(population); + cudaSimulation.initialise(argc, argv); + cudaSimulation.simulate(); + + cudaSimulation.getPopulationData(population); + + // Ensure profiling / memcheck work correctly + flamegpu::util::cleanup(); + + return 0; +}