From f7fd0b67fdc00bc44a5bb72acce8255bbfe2cabf Mon Sep 17 00:00:00 2001 From: Rafal Bielski <76829822+rafbiels@users.noreply.github.com> Date: Tue, 9 Jan 2024 18:38:13 +0000 Subject: [PATCH] Apply clang-format (#6) Apply clang-format 17.0.6 with --style=Google --- .clang-format | 18 ++ .../scatter_reduce_gather.cpp | 12 +- src/MPI_for_CUDA_backend/send_recv_buff.cpp | 43 ++-- src/MPI_for_CUDA_backend/send_recv_usm.cpp | 17 +- src/fluid/fluid.h | 26 ++- src/fluid/main.cpp | 70 ++++--- src/game_of_life/main.cpp | 97 ++++----- src/game_of_life/sim.hpp | 21 +- src/mandelbrot/main.cpp | 60 +++--- .../matrix-multiply.cpp | 1 - src/nbody/assets/NBodyShader.vert | 15 +- src/nbody/main.cpp | 184 ++++++++++-------- src/nbody/sim.hpp | 23 ++- src/nbody/sycl_bufs.hpp | 30 ++- src/nbody/tuple_utils.hpp | 2 +- src/scan_parallel_inclusive/scan.cpp | 3 +- 16 files changed, 332 insertions(+), 290 deletions(-) create mode 100644 .clang-format diff --git a/.clang-format b/.clang-format new file mode 100644 index 0000000..aa68020 --- /dev/null +++ b/.clang-format @@ -0,0 +1,18 @@ +BasedOnStyle: Google +IncludeCategories: + - Regex: '^$' + Priority: 3 + SortPriority: 0 + CaseSensitive: false + - Regex: '^<.*\.h>' + Priority: 1 + SortPriority: 0 + CaseSensitive: false + - Regex: '^<.*' + Priority: 2 + SortPriority: 0 + CaseSensitive: false + - Regex: '.*' + Priority: 4 + SortPriority: 0 + CaseSensitive: false diff --git a/src/MPI_for_CUDA_backend/scatter_reduce_gather.cpp b/src/MPI_for_CUDA_backend/scatter_reduce_gather.cpp index cfa67cd..8628393 100644 --- a/src/MPI_for_CUDA_backend/scatter_reduce_gather.cpp +++ b/src/MPI_for_CUDA_backend/scatter_reduce_gather.cpp @@ -1,4 +1,5 @@ -// Compile with `mpicxx -fsycl -fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend --cuda-gpu-arch=sm_xx scatter_reduce_gather.cpp -o res` +// Compile with `mpicxx -fsycl -fsycl-targets=nvptx64-nvidia-cuda +// -Xsycl-target-backend --cuda-gpu-arch=sm_xx scatter_reduce_gather.cpp -o res` // Where sm_xx is the Compute Capability (CC). If the `-Xsycl-target-backend // --cuda-gpu-arch=` flags are not explicitly provided the lowest supported CC // will be used: sm_50. @@ -14,10 +15,10 @@ #include #include + #include int main(int argc, char *argv[]) { - /* ------------------------------------------------------------------------------------------- MPI Initialization. --------------------------------------------------------------------------------------------*/ @@ -32,9 +33,10 @@ int main(int argc, char *argv[]) { if (size != 2) { if (rank == 0) { - printf("This program requires exactly 2 MPI ranks, but you are " - "attempting to use %d! Exiting...\n", - size); + printf( + "This program requires exactly 2 MPI ranks, " + "but you are attempting to use %d! Exiting...\n", + size); } MPI_Finalize(); exit(0); diff --git a/src/MPI_for_CUDA_backend/send_recv_buff.cpp b/src/MPI_for_CUDA_backend/send_recv_buff.cpp index 88c03c9..d466ebe 100644 --- a/src/MPI_for_CUDA_backend/send_recv_buff.cpp +++ b/src/MPI_for_CUDA_backend/send_recv_buff.cpp @@ -1,5 +1,6 @@ -// Compile with `mpicxx -fsycl -fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend --cuda-gpu-arch=sm_xx send_recv_buff.cpp -o res` -// Where sm_xx is the Compute Capability (CC). If the `-Xsycl-target-backend +// Compile with `mpicxx -fsycl -fsycl-targets=nvptx64-nvidia-cuda +// -Xsycl-target-backend --cuda-gpu-arch=sm_xx send_recv_buff.cpp -o res` +// where sm_xx is the Compute Capability (CC). If the `-Xsycl-target-backend // --cuda-gpu-arch=` flags are not explicitly provided the lowest supported CC // will be used: sm_50. @@ -8,13 +9,13 @@ #include #include + #include int main(int argc, char *argv[]) { - - /* ------------------------------------------------------------------------------------------- - MPI Initialization. - --------------------------------------------------------------------------------------------*/ + /* --------------------------------------------------------------------------- + MPI Initialization. + ----------------------------------------------------------------------------*/ MPI_Init(&argc, &argv); @@ -26,17 +27,18 @@ int main(int argc, char *argv[]) { if (size != 2) { if (rank == 0) { - printf("This program requires exactly 2 MPI ranks, but you are " - "attempting to use %d! Exiting...\n", - size); + printf( + "This program requires exactly 2 MPI ranks, " + "but you are attempting to use %d! Exiting...\n", + size); } MPI_Finalize(); exit(0); } - /* ------------------------------------------------------------------------------------------- - SYCL Initialization, which internally sets the CUDA device. - --------------------------------------------------------------------------------------------*/ + /* --------------------------------------------------------------------------- + SYCL Initialization, which internally sets the CUDA device. + ----------------------------------------------------------------------------*/ sycl::queue q{}; @@ -46,16 +48,16 @@ int main(int argc, char *argv[]) { std::vector data(nelem, -1); { - /* ------------------------------------------------------------------------------------------- - Create a SYCL buffer in each rank. The sycl::buffer created in each rank will - manage the copy of data to and from the device as required. - --------------------------------------------------------------------------------------------*/ + /* ------------------------------------------------------------------------- + Create a SYCL buffer in each rank. The sycl::buffer created in each rank + will manage the copy of data to and from the device as required. + --------------------------------------------------------------------------*/ sycl::buffer buff(&data[0], sycl::range{nelem}); - /* ------------------------------------------------------------------------------------------- - Perform the send/receive. - --------------------------------------------------------------------------------------------*/ + /* ------------------------------------------------------------------------- + Perform the send/receive. + --------------------------------------------------------------------------*/ if (rank == 0) { // Operate on the Rank 0 data. @@ -107,8 +109,7 @@ int main(int argc, char *argv[]) { // Check the values. Since this is outside the scope where the buffer was // created, the data array is automatically updated on the host. if (rank == 1) { - for (int i = 0; i < nelem; ++i) - assert(data[i] == -2); + for (int i = 0; i < nelem; ++i) assert(data[i] == -2); } MPI_Finalize(); diff --git a/src/MPI_for_CUDA_backend/send_recv_usm.cpp b/src/MPI_for_CUDA_backend/send_recv_usm.cpp index 6e74ba1..d363d1b 100644 --- a/src/MPI_for_CUDA_backend/send_recv_usm.cpp +++ b/src/MPI_for_CUDA_backend/send_recv_usm.cpp @@ -1,5 +1,6 @@ -// Compile with `mpicxx -fsycl -fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend --cuda-gpu-arch=sm_xx send_recv_usm.cpp -o res` -// Where sm_xx is the Compute Capability (CC). If the `-Xsycl-target-backend +// Compile with `mpicxx -fsycl -fsycl-targets=nvptx64-nvidia-cuda +// -Xsycl-target-backend --cuda-gpu-arch=sm_xx send_recv_usm.cpp -o res` Where +// sm_xx is the Compute Capability (CC). If the `-Xsycl-target-backend // --cuda-gpu-arch=` flags are not explicitly provided the lowest supported CC // will be used: sm_50. @@ -8,10 +9,10 @@ #include #include + #include int main(int argc, char *argv[]) { - /* ------------------------------------------------------------------------------------------- MPI Initialization. --------------------------------------------------------------------------------------------*/ @@ -26,9 +27,10 @@ int main(int argc, char *argv[]) { if (size != 2) { if (rank == 0) { - printf("This program requires exactly 2 MPI ranks, but you are " - "attempting to use %d! Exiting...\n", - size); + printf( + "This program requires exactly 2 MPI ranks, " + "but you are attempting to use %d! Exiting...\n", + size); } MPI_Finalize(); exit(0); @@ -86,8 +88,7 @@ int main(int argc, char *argv[]) { sycl::free(devp, q); // Check the values. - for (int i = 0; i < nelem; ++i) - assert(data[i] == -2); + for (int i = 0; i < nelem; ++i) assert(data[i] == -2); } MPI_Finalize(); return 0; diff --git a/src/fluid/fluid.h b/src/fluid/fluid.h index 7fe0511..d6458d4 100644 --- a/src/fluid/fluid.h +++ b/src/fluid/fluid.h @@ -49,8 +49,7 @@ class SYCLFluidContainer { img{sycl::range<1>(size * size)}, // Initialize queue with default selector and asynchronous exception // handler. - queue{sycl::default_selector_v, - [](sycl::exception_list exceptions) { + queue{sycl::default_selector_v, [](sycl::exception_list exceptions) { for (const std::exception_ptr& e : exceptions) { try { std::rethrow_exception(e); @@ -208,10 +207,8 @@ class SYCLFluidContainer { // Update the image pixel data with the appropriate color for a given // density. queue.submit([&](sycl::handler& cgh) { - auto img_acc{ - img.template get_access(cgh, sycl::write_only)}; - auto density_a{ - density_b.template get_access(cgh, sycl::read_write)}; + auto img_acc{img.template get_access(cgh, sycl::write_only)}; + auto density_a{density_b.template get_access(cgh, sycl::read_write)}; cgh.parallel_for( sycl::range<1>(size * size), [=](sycl::item<1> item) { auto index{item.get_id(0)}; @@ -229,8 +226,8 @@ class SYCLFluidContainer { using float_buffer = sycl::buffer; using read_write_accessor = sycl::accessor; + sycl::access::target::device, + sycl::access::placeholder::false_t>; // Wrapper around queue submission. template @@ -297,9 +294,9 @@ class SYCLFluidContainer { } // Solve linear differential equation of density / velocity. (SYCL VERSION). - static void LinearSolve(int /*b*/, read_write_accessor x, read_write_accessor x0, - float a, float c_reciprocal, std::size_t N, - sycl::handler& cgh) { + static void LinearSolve(int /*b*/, read_write_accessor x, + read_write_accessor x0, float a, float c_reciprocal, + std::size_t N, sycl::handler& cgh) { cgh.parallel_for( sycl::range<2>(N - 2, N - 2), [=](sycl::item<2> item) { auto i{1 + item.get_id(0)}; @@ -402,9 +399,10 @@ class SYCLFluidContainer { } // Move density / velocity within the field to the next step. (SYCL VERSION). - static void AdvectImpl(int /*b*/, read_write_accessor d, read_write_accessor d0, - read_write_accessor u, read_write_accessor v, - float dt0, std::size_t N, sycl::handler& cgh) { + static void AdvectImpl(int /*b*/, read_write_accessor d, + read_write_accessor d0, read_write_accessor u, + read_write_accessor v, float dt0, std::size_t N, + sycl::handler& cgh) { cgh.parallel_for( sycl::range<2>(N - 2, N - 2), [=](sycl::item<2> item) { auto i{1 + item.get_id(0)}; diff --git a/src/fluid/main.cpp b/src/fluid/main.cpp index 5163f3b..43fdc90 100644 --- a/src/fluid/main.cpp +++ b/src/fluid/main.cpp @@ -18,24 +18,31 @@ * **************************************************************************/ -#include "fluid.h" - -#include +#include #include #include #include #include -#include #include -#include -#include #include +#include +#include +#include +#include #include #include +#include "fluid.h" + constexpr Magnum::PixelFormat PIXELFORMAT{Magnum::PixelFormat::RGBA8Unorm}; +// Title bar text +constexpr std::string_view WINDOWTITLE{ + "Codeplay Fluid Simulation" + " - Move mouse to add fluid" + " - Press space to clear fluid"}; + // Size of the fluid container edge (always square shaped). constexpr int SIZE{300}; @@ -45,27 +52,26 @@ constexpr int SCALE{3}; class FluidSimulationApp : public Magnum::Platform::Application { public: FluidSimulationApp(const Arguments& arguments) - : Magnum::Platform::Application{ - arguments, - Configuration{}.setTitle("Codeplay Fluid Simulation" - " - Move mouse to add fluid - Press space to clear fluid"), - GLConfiguration{}.setFlags(GLConfiguration::Flag::QuietLog)}, - size_{SIZE}, - fluid_{SIZE, 0.2f, 0.0f, 0.0000001f}, - mesh_{Magnum::MeshTools::compile(Magnum::Primitives::squareSolid( - Magnum::Primitives::SquareFlag::TextureCoordinates))}, - shader_{Magnum::Shaders::FlatGL2D::Configuration{}.setFlags( - Magnum::Shaders::FlatGL2D::Flag::Textured | - Magnum::Shaders::FlatGL2D::Flag::TextureTransformation)} { - + : Magnum::Platform::Application{arguments, + Configuration{}.setTitle(WINDOWTITLE), + GLConfiguration{}.setFlags( + GLConfiguration::Flag::QuietLog)}, + size_{SIZE}, + fluid_{SIZE, 0.2f, 0.0f, 0.0000001f}, + mesh_{Magnum::MeshTools::compile(Magnum::Primitives::squareSolid( + Magnum::Primitives::SquareFlag::TextureCoordinates))}, + shader_{Magnum::Shaders::FlatGL2D::Configuration{}.setFlags( + Magnum::Shaders::FlatGL2D::Flag::Textured | + Magnum::Shaders::FlatGL2D::Flag::TextureTransformation)} { // Set window size. - setWindowSize({SIZE*SCALE, SIZE*SCALE}); - Magnum::GL::defaultFramebuffer.setViewport({{0,0},{SIZE*SCALE, SIZE*SCALE}}); + setWindowSize({SIZE * SCALE, SIZE * SCALE}); + Magnum::GL::defaultFramebuffer.setViewport( + {{0, 0}, {SIZE * SCALE, SIZE * SCALE}}); texture_.setWrapping(Magnum::GL::SamplerWrapping::ClampToEdge) - .setMagnificationFilter(Magnum::GL::SamplerFilter::Linear) - .setMinificationFilter(Magnum::GL::SamplerFilter::Linear) - .setStorage(1, Magnum::GL::textureFormat(PIXELFORMAT), {size_, size_}); + .setMagnificationFilter(Magnum::GL::SamplerFilter::Linear) + .setMinificationFilter(Magnum::GL::SamplerFilter::Linear) + .setStorage(1, Magnum::GL::textureFormat(PIXELFORMAT), {size_, size_}); shader_.bindTexture(texture_); } @@ -89,18 +95,18 @@ class FluidSimulationApp : public Magnum::Platform::Application { // Draws fluid to the screen. void drawEvent() override final { // Clear screen. - Magnum::GL::defaultFramebuffer.clear( - Magnum::GL::FramebufferClear::Color | - Magnum::GL::FramebufferClear::Depth); + Magnum::GL::defaultFramebuffer.clear(Magnum::GL::FramebufferClear::Color | + Magnum::GL::FramebufferClear::Depth); // Update texture with pixel data array. fluid_.WithData([&](sycl::uchar4 const* data) { - Magnum::ImageView2D img{ - PIXELFORMAT, {size_, size_}, - Corrade::Containers::ArrayView{ + Magnum::ImageView2D img{ + PIXELFORMAT, + {size_, size_}, + Corrade::Containers::ArrayView{ reinterpret_cast(data), - size_*size_*Magnum::pixelFormatSize(PIXELFORMAT)}}; - texture_.setSubImage(0, {0,0}, img); + size_ * size_ * Magnum::pixelFormatSize(PIXELFORMAT)}}; + texture_.setSubImage(0, {0, 0}, img); }); // Draw texture to screen. diff --git a/src/game_of_life/main.cpp b/src/game_of_life/main.cpp index dd3eeb8..1764bb8 100644 --- a/src/game_of_life/main.cpp +++ b/src/game_of_life/main.cpp @@ -18,28 +18,28 @@ * **************************************************************************/ -#include "sim.hpp" - -#include #include #include #include #include -#include #include -#include -#include #include +#include +#include +#include +#include #include -#include -#include #include +#include + +#include + +#include "sim.hpp" constexpr Magnum::PixelFormat PIXELFORMAT{Magnum::PixelFormat::RGBA8Unorm}; -class GameOfLifeApp : public Magnum::Platform::Application -{ +class GameOfLifeApp : public Magnum::Platform::Application { /// The window dimensions int m_width; int m_height; @@ -71,27 +71,30 @@ class GameOfLifeApp : public Magnum::Platform::Application public: GameOfLifeApp(const Arguments& arguments) - : Magnum::Platform::Application{ - arguments, - Configuration{}.setTitle("Codeplay Game of Life Demo") - .addWindowFlags(Configuration::WindowFlag::Resizable), - GLConfiguration{}.setFlags(GLConfiguration::Flag::QuietLog)}, - m_width(windowSize().x()), - m_height(windowSize().y()), - m_zoom(1), - m_resized_width{m_width}, - m_resized_height{m_height}, - m_sim(m_width, m_height), - m_mesh{Magnum::MeshTools::compile(Magnum::Primitives::squareSolid( - Magnum::Primitives::SquareFlag::TextureCoordinates))}, - m_shader{Magnum::Shaders::FlatGL2D::Configuration{}.setFlags( - Magnum::Shaders::FlatGL2D::Flag::Textured | - Magnum::Shaders::FlatGL2D::Flag::TextureTransformation)} { - + : Magnum::Platform:: + Application{arguments, + Configuration{} + .setTitle("Codeplay Game of Life Demo") + .addWindowFlags( + Configuration::WindowFlag::Resizable), + GLConfiguration{}.setFlags( + GLConfiguration::Flag::QuietLog)}, + m_width(windowSize().x()), + m_height(windowSize().y()), + m_zoom(1), + m_resized_width{m_width}, + m_resized_height{m_height}, + m_sim(m_width, m_height), + m_mesh{Magnum::MeshTools::compile(Magnum::Primitives::squareSolid( + Magnum::Primitives::SquareFlag::TextureCoordinates))}, + m_shader{Magnum::Shaders::FlatGL2D::Configuration{}.setFlags( + Magnum::Shaders::FlatGL2D::Flag::Textured | + Magnum::Shaders::FlatGL2D::Flag::TextureTransformation)} { m_tex.setWrapping(Magnum::GL::SamplerWrapping::ClampToEdge) - .setMagnificationFilter(Magnum::GL::SamplerFilter::Nearest) - .setMinificationFilter(Magnum::GL::SamplerFilter::Nearest) - .setStorage(1, Magnum::GL::textureFormat(PIXELFORMAT), {m_width,m_height}); + .setMagnificationFilter(Magnum::GL::SamplerFilter::Nearest) + .setMinificationFilter(Magnum::GL::SamplerFilter::Nearest) + .setStorage(1, Magnum::GL::textureFormat(PIXELFORMAT), + {m_width, m_height}); m_shader.bindTexture(m_tex); } @@ -105,9 +108,10 @@ class GameOfLifeApp : public Magnum::Platform::Application // Reinitializes image to new size m_tex = Magnum::GL::Texture2D{}; m_tex.setWrapping(Magnum::GL::SamplerWrapping::ClampToEdge) - .setMagnificationFilter(Magnum::GL::SamplerFilter::Nearest) - .setMinificationFilter(Magnum::GL::SamplerFilter::Nearest) - .setStorage(1, Magnum::GL::textureFormat(PIXELFORMAT), {m_width,m_height}); + .setMagnificationFilter(Magnum::GL::SamplerFilter::Nearest) + .setMinificationFilter(Magnum::GL::SamplerFilter::Nearest) + .setStorage(1, Magnum::GL::textureFormat(PIXELFORMAT), + {m_width, m_height}); m_shader.bindTexture(m_tex); m_resized = false; @@ -123,19 +127,19 @@ class GameOfLifeApp : public Magnum::Platform::Application } void drawEvent() override { - Magnum::GL::defaultFramebuffer.clear( - Magnum::GL::FramebufferClear::Color | - Magnum::GL::FramebufferClear::Depth); + Magnum::GL::defaultFramebuffer.clear(Magnum::GL::FramebufferClear::Color | + Magnum::GL::FramebufferClear::Depth); m_sim.with_img( // Gets called with image data [&](sycl::uchar4 const* data) { Magnum::ImageView2D img{ - PIXELFORMAT, {m_width, m_height}, - Corrade::Containers::ArrayView{ - reinterpret_cast(data), - m_width*m_height*Magnum::pixelFormatSize(PIXELFORMAT)}}; - m_tex.setSubImage(0, {0,0}, img); + PIXELFORMAT, + {m_width, m_height}, + Corrade::Containers::ArrayView{ + reinterpret_cast(data), + m_width * m_height * Magnum::pixelFormatSize(PIXELFORMAT)}}; + m_tex.setSubImage(0, {0, 0}, img); }); m_shader.draw(m_mesh); @@ -158,8 +162,8 @@ class GameOfLifeApp : public Magnum::Platform::Application m_sim.add_click(x, y + 1, CellState::LIVE); m_sim.add_click(x + 1, y, CellState::LIVE); m_sim.add_click(x, y - 1, CellState::LIVE); - m_sim.add_click(x - 1, y - 1, CellState::LIVE); - } + m_sim.add_click(x - 1, y - 1, CellState::LIVE); + } } void mousePressEvent(MouseEvent& event) override { @@ -187,7 +191,9 @@ class GameOfLifeApp : public Magnum::Platform::Application m_zoom = sycl::clamp(m_zoom, 1.0f, 64.0f); // Restart after zooming - if (m_zoom!=prevZoom) {m_resized = true;} + if (m_zoom != prevZoom) { + m_resized = true; + } } void keyPressEvent(KeyEvent& event) override { @@ -201,7 +207,8 @@ class GameOfLifeApp : public Magnum::Platform::Application m_resized_width = event.windowSize().x(); m_resized_height = event.windowSize().y(); m_resized = true; - Magnum::GL::defaultFramebuffer.setViewport({{0,0}, event.framebufferSize()}); + Magnum::GL::defaultFramebuffer.setViewport( + {{0, 0}, event.framebufferSize()}); } }; diff --git a/src/game_of_life/sim.hpp b/src/game_of_life/sim.hpp index d91cc01..7f27c2e 100644 --- a/src/game_of_life/sim.hpp +++ b/src/game_of_life/sim.hpp @@ -20,11 +20,11 @@ #pragma once -#include "../include/double_buf.hpp" +#include #include -#include +#include "../include/double_buf.hpp" enum class CellState : unsigned int { LIVE = 1, @@ -75,18 +75,15 @@ class GameOfLifeSim { } } }) { - // Initialize game grid in a way that 1/4 of the cells will be live and rest // dead - auto acells = - m_game.read().cells.get_host_access(sycl::write_only); - auto aimg = - m_game.read().img.get_host_access(sycl::write_only); + auto acells = m_game.read().cells.get_host_access(sycl::write_only); + auto aimg = m_game.read().img.get_host_access(sycl::write_only); std::random_device rd; - std::default_random_engine re{ rd() }; - std::bernoulli_distribution dist{ 0.75 }; + std::default_random_engine re{rd()}; + std::bernoulli_distribution dist{0.75}; for (size_t y = 0; y < m_height; y++) { for (size_t x = 0; x < m_width; x++) { @@ -219,9 +216,9 @@ class GameOfLifeSim { new_vel = sycl::fabs(new_vel) * 5.0f + sycl::float2(0.2f, 0.2f); // Set image pixel to new colour decided by state and "velocity" - img[sycl::id<2>(y, x)] = sycl::uchar4( - float(int(new_state)) * new_vel.x() * 255.0f, 0, - float(int(new_state)) * new_vel.y() * 255.0f, 255); + img[sycl::id<2>(y, x)] = + sycl::uchar4(float(int(new_state)) * new_vel.x() * 255.0f, 0, + float(int(new_state)) * new_vel.y() * 255.0f, 255); }); }); diff --git a/src/mandelbrot/main.cpp b/src/mandelbrot/main.cpp index a4b452c..0f3ecb8 100644 --- a/src/mandelbrot/main.cpp +++ b/src/mandelbrot/main.cpp @@ -18,30 +18,29 @@ * **************************************************************************/ - -#include "mandel.hpp" - -#include #include #include #include #include -#include #include -#include -#include #include +#include +#include +#include +#include #include #include +#include "mandel.hpp" + constexpr size_t WIDTH = 800; constexpr size_t HEIGHT = 600; constexpr Magnum::PixelFormat PIXELFORMAT{Magnum::PixelFormat::RGBA8Unorm}; -const size_t PIXELDATASIZE{WIDTH*HEIGHT*Magnum::pixelFormatSize(PIXELFORMAT)}; +const size_t PIXELDATASIZE{WIDTH * HEIGHT * + Magnum::pixelFormatSize(PIXELFORMAT)}; -class MandelbrotApp : public Magnum::Platform::Application -{ +class MandelbrotApp : public Magnum::Platform::Application { // Use doubles for more zoom MandelbrotCalculator m_calc; @@ -65,21 +64,21 @@ class MandelbrotApp : public Magnum::Platform::Application public: MandelbrotApp(const Arguments& arguments) - : Magnum::Platform::Application{ - arguments, - Configuration{}.setTitle("Codeplay Mandelbrot Demo"), - GLConfiguration{}.setFlags(GLConfiguration::Flag::QuietLog)}, - m_calc{WIDTH, HEIGHT}, - m_mesh{Magnum::MeshTools::compile(Magnum::Primitives::squareSolid( - Magnum::Primitives::SquareFlag::TextureCoordinates))}, - m_shader{Magnum::Shaders::FlatGL2D::Configuration{}.setFlags( - Magnum::Shaders::FlatGL2D::Flag::Textured | - Magnum::Shaders::FlatGL2D::Flag::TextureTransformation)} { - + : Magnum::Platform::Application{arguments, + Configuration{}.setTitle( + "Codeplay Mandelbrot Demo"), + GLConfiguration{}.setFlags( + GLConfiguration::Flag::QuietLog)}, + m_calc{WIDTH, HEIGHT}, + m_mesh{Magnum::MeshTools::compile(Magnum::Primitives::squareSolid( + Magnum::Primitives::SquareFlag::TextureCoordinates))}, + m_shader{Magnum::Shaders::FlatGL2D::Configuration{}.setFlags( + Magnum::Shaders::FlatGL2D::Flag::Textured | + Magnum::Shaders::FlatGL2D::Flag::TextureTransformation)} { m_tex.setWrapping(Magnum::GL::SamplerWrapping::ClampToEdge) - .setMagnificationFilter(Magnum::GL::SamplerFilter::Linear) - .setMinificationFilter(Magnum::GL::SamplerFilter::Linear) - .setStorage(1, Magnum::GL::textureFormat(PIXELFORMAT), {WIDTH,HEIGHT}); + .setMagnificationFilter(Magnum::GL::SamplerFilter::Linear) + .setMinificationFilter(Magnum::GL::SamplerFilter::Linear) + .setStorage(1, Magnum::GL::textureFormat(PIXELFORMAT), {WIDTH, HEIGHT}); m_shader.bindTexture(m_tex); } @@ -104,16 +103,17 @@ class MandelbrotApp : public Magnum::Platform::Application } void drawEvent() override { - Magnum::GL::defaultFramebuffer.clear( - Magnum::GL::FramebufferClear::Color | - Magnum::GL::FramebufferClear::Depth); + Magnum::GL::defaultFramebuffer.clear(Magnum::GL::FramebufferClear::Color | + Magnum::GL::FramebufferClear::Depth); // Update GL texture with new calculation data m_calc.with_data([&](sycl::uchar4 const* data) { Magnum::ImageView2D img{ - PIXELFORMAT, {WIDTH,HEIGHT}, - Corrade::Containers::ArrayView{reinterpret_cast(data), PIXELDATASIZE}}; - m_tex.setSubImage(0, {0,0}, img); + PIXELFORMAT, + {WIDTH, HEIGHT}, + Corrade::Containers::ArrayView{reinterpret_cast(data), + PIXELDATASIZE}}; + m_tex.setSubImage(0, {0, 0}, img); }); m_shader.draw(m_mesh); diff --git a/src/matrix_multiply_omp_compare/matrix-multiply.cpp b/src/matrix_multiply_omp_compare/matrix-multiply.cpp index 4234f6f..e7e73a8 100644 --- a/src/matrix_multiply_omp_compare/matrix-multiply.cpp +++ b/src/matrix_multiply_omp_compare/matrix-multiply.cpp @@ -25,7 +25,6 @@ * See block_host for the OpenMP implementation. */ #include - #include #include #include diff --git a/src/nbody/assets/NBodyShader.vert b/src/nbody/assets/NBodyShader.vert index 8f36d31..c99e0ca 100644 --- a/src/nbody/assets/NBodyShader.vert +++ b/src/nbody/assets/NBodyShader.vert @@ -25,21 +25,18 @@ in vec3 ciVelocity; out vec3 vColor; vec3 speedColors[9] = - vec3[](vec3(0.0, 0.0, 0.2), vec3(0.0, 0.0, 0.4), - vec3(0.0, 0.0, 0.8), vec3(0.0, 0.4, 0.4), - vec3(0.0, 0.8, 0.8), vec3(0.0, 0.8, 0.4), - vec3(0.4, 0.8, 0.0), vec3(0.8, 0.6, 0.0), - vec3(0.8, 0.2, 0.0)); + vec3[](vec3(0.0, 0.0, 0.2), vec3(0.0, 0.0, 0.4), vec3(0.0, 0.0, 0.8), + vec3(0.0, 0.4, 0.4), vec3(0.0, 0.8, 0.8), vec3(0.0, 0.8, 0.4), + vec3(0.4, 0.8, 0.0), vec3(0.8, 0.6, 0.0), vec3(0.8, 0.2, 0.0)); void main() { vec4 viewSpacePos = ciModelView * vec4(ciPosition, 1.0); // The further away a point is, the smaller its sprite float scale = log2(length(viewSpacePos)); gl_PointSize = 10.0 / clamp(scale, 0.1, 1); - gl_Position = - ciModelViewProjection * vec4(ciPosition / 10.0, 1.0); + gl_Position = ciModelViewProjection * vec4(ciPosition / 10.0, 1.0); float len = length(ciVelocity); int speed = int(ceil(len)); - vColor = mix(speedColors[max(0, speed)], - speedColors[min(8, speed)], speed - len); + vColor = + mix(speedColors[max(0, speed)], speedColors[min(8, speed)], speed - len); } diff --git a/src/nbody/main.cpp b/src/nbody/main.cpp index 8823b5f..75847f5 100644 --- a/src/nbody/main.cpp +++ b/src/nbody/main.cpp @@ -18,38 +18,36 @@ * **************************************************************************/ -#include "sim.hpp" - -#include -#include -#include -#include -#include -#include -#include - +#include +#include #include -#include #include #include #include +#include #include #include #include - #include +#include +#include +#include +#include +#include #include -#include +#include #include #include -#include -#include - -#include +#include +#include #include #include +#include + +#include "sim.hpp" + using num_t = float; constexpr num_t PI{3.141592653589793238462643383279502884197169399}; @@ -57,11 +55,13 @@ class NBodyShader : public Magnum::GL::AbstractShaderProgram { public: using Position = Magnum::GL::Attribute<0, Magnum::Vector3>; using Velocity = Magnum::GL::Attribute<1, Magnum::Vector3>; - enum : unsigned int {ColourOutput = 0}; + enum : unsigned int { ColourOutput = 0 }; explicit NBodyShader() { /* Load shader sources */ - Magnum::GL::Shader vert{Magnum::GL::Version::GL330, Magnum::GL::Shader::Type::Vertex}; - Magnum::GL::Shader frag{Magnum::GL::Version::GL330, Magnum::GL::Shader::Type::Fragment}; + Magnum::GL::Shader vert{Magnum::GL::Version::GL330, + Magnum::GL::Shader::Type::Vertex}; + Magnum::GL::Shader frag{Magnum::GL::Version::GL330, + Magnum::GL::Shader::Type::Fragment}; const Corrade::Utility::Resource rs{"nbody-data"}; vert.addSource(rs.getString("NBodyShader.vert")); frag.addSource(rs.getString("NBodyShader.frag")); @@ -80,14 +80,20 @@ class NBodyShader : public Magnum::GL::AbstractShaderProgram { CORRADE_INTERNAL_ASSERT_OUTPUT(link()); } - NBodyShader& setView(Corrade::Containers::ArrayView> matrix) { - setUniform(uniformLocation("ciModelView"), matrix); - return *this; + NBodyShader& setView( + Corrade::Containers::ArrayView< + const Magnum::Math::RectangularMatrix<4, 4, Magnum::Float>> + matrix) { + setUniform(uniformLocation("ciModelView"), matrix); + return *this; } - NBodyShader& setViewProjection(Corrade::Containers::ArrayView> matrix) { - setUniform(uniformLocation("ciModelViewProjection"), matrix); - return *this; + NBodyShader& setViewProjection( + Corrade::Containers::ArrayView< + const Magnum::Math::RectangularMatrix<4, 4, Magnum::Float>> + matrix) { + setUniform(uniformLocation("ciModelViewProjection"), matrix); + return *this; } NBodyShader& bindTexture(Magnum::GL::Texture2D& tex) { @@ -97,8 +103,7 @@ class NBodyShader : public Magnum::GL::AbstractShaderProgram { } }; -class NBodyApp : public Magnum::Platform::Application -{ +class NBodyApp : public Magnum::Platform::Application { // -- GUI -- // Distribution choice enum { @@ -181,18 +186,21 @@ class NBodyApp : public Magnum::Platform::Application public: NBodyApp(const Arguments& arguments) - : Magnum::Platform::Application{ - arguments, - Configuration{}.setTitle("Codeplay NBody Demo") - .addWindowFlags(Configuration::WindowFlag::Resizable), - GLConfiguration{}.setFlags(GLConfiguration::Flag::QuietLog)}, - m_imgui{Magnum::Vector2{windowSize()}/dpiScaling(), windowSize(), framebufferSize()}, - m_sim(m_n_bodies, distrib_cylinder{}) { - + : Magnum::Platform:: + Application{arguments, + Configuration{} + .setTitle("Codeplay NBody Demo") + .addWindowFlags( + Configuration::WindowFlag::Resizable), + GLConfiguration{}.setFlags( + GLConfiguration::Flag::QuietLog)}, + m_imgui{Magnum::Vector2{windowSize()} / dpiScaling(), windowSize(), + framebufferSize()}, + m_sim(m_n_bodies, distrib_cylinder{}) { const Corrade::Utility::Resource rs{"nbody-data"}; Corrade::PluginManager::Manager manager; auto importer{manager.loadAndInstantiate("AnyImageImporter")}; - if (importer==nullptr || !importer->openData(rs.getRaw("star.png"))) { + if (importer == nullptr || !importer->openData(rs.getRaw("star.png"))) { throw std::runtime_error{"Failed to load star.png"}; } auto image{importer->image2D(0)}; @@ -202,20 +210,22 @@ class NBodyApp : public Magnum::Platform::Application // Create star texture and bind it to GL slot 0 m_star_tex.setWrapping(Magnum::GL::SamplerWrapping::ClampToEdge) - .setMagnificationFilter(Magnum::GL::SamplerFilter::Linear) - .setMinificationFilter(Magnum::GL::SamplerFilter::Linear) - .setStorage(1, Magnum::GL::textureFormat(image->format()), image->size()) - .setSubImage(0, {}, *image); + .setMagnificationFilter(Magnum::GL::SamplerFilter::Linear) + .setMinificationFilter(Magnum::GL::SamplerFilter::Linear) + .setStorage(1, Magnum::GL::textureFormat(image->format()), + image->size()) + .setSubImage(0, {}, *image); m_ui_force_coulomb_file.fill(0); // Start with a perspective camera looking at the center using namespace Magnum::Math::Literals; - m_view = Magnum::Matrix4::lookAt( - Magnum::Vector3{200}, Magnum::Vector3{0}, Magnum::Vector3{0,0,1}) - .invertedRigid(); + m_view = Magnum::Matrix4::lookAt(Magnum::Vector3{200}, Magnum::Vector3{0}, + Magnum::Vector3{0, 0, 1}) + .invertedRigid(); m_viewProjection = Magnum::Matrix4::perspectiveProjection( - 60.0_degf, Magnum::Vector2{windowSize()}.aspectRatio(), 0.01f, 50000.0f); + 60.0_degf, Magnum::Vector2{windowSize()}.aspectRatio(), 0.01f, + 50000.0f); init_gl_bufs(); @@ -228,7 +238,8 @@ class NBodyApp : public Magnum::Platform::Application // bodies void init_gl_bufs() { const size_t arraySize{m_n_bodies * sizeof(sycl::vec)}; - m_vboStorage = Corrade::Containers::Array(Corrade::ValueInit, 2*arraySize); + m_vboStorage = + Corrade::Containers::Array(Corrade::ValueInit, 2 * arraySize); m_mesh = Magnum::GL::Mesh{Magnum::GL::MeshPrimitive::Points}; m_mesh.setCount(m_n_bodies); } @@ -349,42 +360,43 @@ class NBodyApp : public Magnum::Platform::Application } void drawEvent() override { - Magnum::GL::defaultFramebuffer - .clear(Magnum::GL::FramebufferClear::Depth) - .clearColor(Magnum::Math::Color4{0.0f,0.0f,0.0f,1.0f}); + Magnum::GL::defaultFramebuffer.clear(Magnum::GL::FramebufferClear::Depth) + .clearColor(Magnum::Math::Color4{0.0f, 0.0f, 0.0f, 1.0f}); // Disable depth to avoid black outlines Magnum::GL::Renderer::disable(Magnum::GL::Renderer::Feature::DepthTest); // Colors add to a bright white with additive blending Magnum::GL::Renderer::setBlendEquation( - Magnum::GL::Renderer::BlendEquation::Add, - Magnum::GL::Renderer::BlendEquation::Add); + Magnum::GL::Renderer::BlendEquation::Add, + Magnum::GL::Renderer::BlendEquation::Add); Magnum::GL::Renderer::setBlendFunction( - Magnum::GL::Renderer::BlendFunction::SourceAlpha, - Magnum::GL::Renderer::BlendFunction::OneMinusSourceAlpha); + Magnum::GL::Renderer::BlendFunction::SourceAlpha, + Magnum::GL::Renderer::BlendFunction::OneMinusSourceAlpha); Magnum::GL::Renderer::enable(Magnum::GL::Renderer::Feature::Blending); // Update star buffer data with new positions const size_t arraySize{m_n_bodies * sizeof(sycl::vec)}; - m_sim.with_mapped( - read_bufs_t<1>{}, [&](sycl::vec const* positions) { - std::copy_n(reinterpret_cast(positions), arraySize, m_vboStorage.data()); - }); - m_sim.with_mapped( - read_bufs_t<0>{}, [&](sycl::vec const* velocities) { - std::copy_n(reinterpret_cast(velocities), arraySize, m_vboStorage.data()+arraySize); - }); + m_sim.with_mapped(read_bufs_t<1>{}, + [&](sycl::vec const* positions) { + std::copy_n(reinterpret_cast(positions), + arraySize, m_vboStorage.data()); + }); + m_sim.with_mapped(read_bufs_t<0>{}, + [&](sycl::vec const* velocities) { + std::copy_n(reinterpret_cast(velocities), + arraySize, m_vboStorage.data() + arraySize); + }); Magnum::GL::Buffer vbo{m_vboStorage}; m_mesh.addVertexBuffer(vbo, 0, NBodyShader::Position{}) - .addVertexBuffer(vbo, arraySize, NBodyShader::Velocity{}); + .addVertexBuffer(vbo, arraySize, NBodyShader::Velocity{}); // Draw bodies m_shader.setView({&m_view, 1}) - .setViewProjection({&m_viewProjection, 1}) - .bindTexture(m_star_tex) - .draw(m_mesh); + .setViewProjection({&m_viewProjection, 1}) + .bindTexture(m_star_tex) + .draw(m_mesh); // TODO: Port the Cinder arrow drawing to Magnum /* @@ -401,9 +413,9 @@ class NBodyApp : public Magnum::Platform::Application m_imgui.newFrame(); /* Enable text input, if needed */ - if(ImGui::GetIO().WantTextInput && !isTextInputActive()) { + if (ImGui::GetIO().WantTextInput && !isTextInputActive()) { startTextInput(); - } else if(!ImGui::GetIO().WantTextInput && isTextInputActive()) { + } else if (!ImGui::GetIO().WantTextInput && isTextInputActive()) { stopTextInput(); } @@ -538,11 +550,11 @@ class NBodyApp : public Magnum::Platform::Application m_imgui.updateApplicationCursor(*this); Magnum::GL::Renderer::setBlendEquation( - Magnum::GL::Renderer::BlendEquation::Add, - Magnum::GL::Renderer::BlendEquation::Add); + Magnum::GL::Renderer::BlendEquation::Add, + Magnum::GL::Renderer::BlendEquation::Add); Magnum::GL::Renderer::setBlendFunction( - Magnum::GL::Renderer::BlendFunction::SourceAlpha, - Magnum::GL::Renderer::BlendFunction::OneMinusSourceAlpha); + Magnum::GL::Renderer::BlendFunction::SourceAlpha, + Magnum::GL::Renderer::BlendFunction::OneMinusSourceAlpha); Magnum::GL::Renderer::enable(Magnum::GL::Renderer::Feature::Blending); Magnum::GL::Renderer::enable(Magnum::GL::Renderer::Feature::ScissorTest); Magnum::GL::Renderer::disable(Magnum::GL::Renderer::Feature::FaceCulling); @@ -556,25 +568,37 @@ class NBodyApp : public Magnum::Platform::Application void viewportEvent(ViewportEvent& event) override { Magnum::GL::defaultFramebuffer.setViewport({{}, event.framebufferSize()}); - m_imgui.relayout(Magnum::Vector2{event.windowSize()}/event.dpiScaling(), - event.windowSize(), event.framebufferSize()); + m_imgui.relayout(Magnum::Vector2{event.windowSize()} / event.dpiScaling(), + event.windowSize(), event.framebufferSize()); } - void keyPressEvent(KeyEvent& event) override {m_imgui.handleKeyPressEvent(event);} - void keyReleaseEvent(KeyEvent& event) override {m_imgui.handleKeyReleaseEvent(event);} + void keyPressEvent(KeyEvent& event) override { + m_imgui.handleKeyPressEvent(event); + } + void keyReleaseEvent(KeyEvent& event) override { + m_imgui.handleKeyReleaseEvent(event); + } - void mousePressEvent(MouseEvent& event) override {m_imgui.handleMousePressEvent(event);} - void mouseReleaseEvent(MouseEvent& event) override {m_imgui.handleMouseReleaseEvent(event);} - void mouseMoveEvent(MouseMoveEvent& event) override {m_imgui.handleMouseMoveEvent(event);} + void mousePressEvent(MouseEvent& event) override { + m_imgui.handleMousePressEvent(event); + } + void mouseReleaseEvent(MouseEvent& event) override { + m_imgui.handleMouseReleaseEvent(event); + } + void mouseMoveEvent(MouseMoveEvent& event) override { + m_imgui.handleMouseMoveEvent(event); + } void mouseScrollEvent(MouseScrollEvent& event) override { - if(m_imgui.handleMouseScrollEvent(event)) { + if (m_imgui.handleMouseScrollEvent(event)) { /* Prevent scrolling the page */ event.setAccepted(); return; } } - void textInputEvent(TextInputEvent& event) override {m_imgui.handleTextInputEvent(event);} + void textInputEvent(TextInputEvent& event) override { + m_imgui.handleTextInputEvent(event); + } }; MAGNUM_APPLICATION_MAIN(NBodyApp) diff --git a/src/nbody/sim.hpp b/src/nbody/sim.hpp index a3f50cc..b22b0fc 100644 --- a/src/nbody/sim.hpp +++ b/src/nbody/sim.hpp @@ -20,17 +20,17 @@ #pragma once -#include "integrator.hpp" -#include "sycl_bufs.hpp" -#include "tuple_utils.hpp" -#include "../include/double_buf.hpp" - -#include - #include #include #include +#include + +#include "../include/double_buf.hpp" +#include "integrator.hpp" +#include "sycl_bufs.hpp" +#include "tuple_utils.hpp" + // Convenience types template using vec3 = sycl::vec; @@ -124,8 +124,8 @@ class GravSim { integrator_t m_integrator; // Base constructor, does not initialize simulation values - GravSim(size_t n_bodies) : - m_q(sycl::default_selector_v, except_handler), + GravSim(size_t n_bodies) + : m_q(sycl::default_selector_v, except_handler), m_bufs(n_bodies), m_n_bodies(n_bodies), m_time(0), @@ -300,7 +300,6 @@ class GravSim { } else if (integrator == integrator_t::RK4) { std::tie(wvelTmp, wposTmp, std::ignore) = integrate_step_rk4(grav, STEP_SIZE, vel[id], pos[id], t); - } wvel[id] = wvelTmp; @@ -354,8 +353,8 @@ class GravSim { integrate_step_rk4(force, STEP_SIZE, vel[id], pos[id], t); } - wvel[id] = wvelTmp; - wpos[id] = wposTmp; + wvel[id] = wvelTmp; + wpos[id] = wposTmp; }); } break; case force_t::COULOMB: { diff --git a/src/nbody/sycl_bufs.hpp b/src/nbody/sycl_bufs.hpp index 4388b56..0b2207d 100644 --- a/src/nbody/sycl_bufs.hpp +++ b/src/nbody/sycl_bufs.hpp @@ -31,42 +31,36 @@ struct BufToReadAccFunc { template AUTO_FUNC( // pair of (buffer, handler) - operator()(In&& in), - std::forward(in) - .first.template get_access(*std::forward(in).second, - sycl::read_only)) + operator()(In && in), std::forward(in).first.template get_access( + *std::forward(in).second, sycl::read_only)) }; // Template function object which transforms buffers to device write accessors struct BufToDcdWriteAccFunc { // pair of (buffer, handler) template - AUTO_FUNC(operator()(In&& in), - std::forward(in) - .first - .template get_access(*std::forward(in).second, - sycl::write_only)) + AUTO_FUNC(operator()(In && in), + std::forward(in).first.template get_access( + *std::forward(in).second, sycl::write_only)) }; // Template function object which transforms buffers to host read accessors struct BufToHostReadAccFunc { template auto operator()(In&& in) - -> decltype(std::forward(in) - .template get_host_access(sycl::read_only)) { - return std::forward(in) - .template get_host_access(sycl::read_only); + -> decltype(std::forward(in).template get_host_access( + sycl::read_only)) { + return std::forward(in).template get_host_access(sycl::read_only); } }; // Template function object which transforms buffers to host write accessors struct BufToHostDcdWriteAccFunc { template - auto operator()(In&& in) -> decltype( - std::forward(in) - .template get_host_access(sycl::write_only)) { - return std::forward(in) - .template get_host_access(sycl::write_only); + auto operator()(In&& in) + -> decltype(std::forward(in).template get_host_access( + sycl::write_only)) { + return std::forward(in).template get_host_access(sycl::write_only); } }; diff --git a/src/nbody/tuple_utils.hpp b/src/nbody/tuple_utils.hpp index 9e9007d..2128b54 100644 --- a/src/nbody/tuple_utils.hpp +++ b/src/nbody/tuple_utils.hpp @@ -20,8 +20,8 @@ #pragma once -#include #include +#include /* Given a function name and a one-line body, defines an auto-return-type * function with these parameters. */ diff --git a/src/scan_parallel_inclusive/scan.cpp b/src/scan_parallel_inclusive/scan.cpp index 0cbda3c..b099b3d 100644 --- a/src/scan_parallel_inclusive/scan.cpp +++ b/src/scan_parallel_inclusive/scan.cpp @@ -19,7 +19,6 @@ **************************************************************************/ #include - #include #include #include @@ -111,7 +110,7 @@ void par_scan(sycl::buffer& in, sycl::queue& q) { q.submit([&](sycl::handler& cgh) { auto data = in.template get_access(cgh); - sycl::local_accessor temp(wgroup_size * 2, cgh); + sycl::local_accessor temp(wgroup_size * 2, cgh); // Use dummy struct as the unique kernel name. cgh.parallel_for>(