diff --git a/Serie09/AUTHORS b/Serie09/AUTHORS new file mode 100644 index 0000000..9253635 --- /dev/null +++ b/Serie09/AUTHORS @@ -0,0 +1,4 @@ +Authors: + - Nicolas Richart + - Vincent Keller + . Vittoria Rezzonico diff --git a/Serie09/COPYING b/Serie09/COPYING new file mode 100644 index 0000000..e551330 --- /dev/null +++ b/Serie09/COPYING @@ -0,0 +1,24 @@ +Copyright (c) 2017, Ecole Polytechnique Federal de Lausanne - Scientific IT and +Application Support (SCITAS) All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions are met: + * Redistributions of source code must retain the above copyright + notice, this list of conditions and the following disclaimer. + * Redistributions in binary form must reproduce the above copyright + notice, this list of conditions and the following disclaimer in the + documentation and/or other materials provided with the distribution. + * Neither the name of the nor the + names of its contributors may be used to endorse or promote products + derived from this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND +ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED +WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +DISCLAIMED. IN NO EVENT SHALL BE LIABLE FOR ANY +DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES +(INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; +LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND +ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS +SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. diff --git a/Serie09/Makefile b/Serie09/Makefile new file mode 100644 index 0000000..c1fbe1f --- /dev/null +++ b/Serie09/Makefile @@ -0,0 +1,17 @@ +CXX=nvcc +LD=${CXX} +CXXFLAGS+=-Xcompiler="-Wall -Wextra -Werror" -std c++11 -O3 +LDFLAGS+=-lm $(CXXFLAGS) + +OBJS=poisson.o simulation.o simulation_gpu.o double_buffer.o dumpers.o + +all: poisson + +poisson: $(OBJS) + $(LD) -o $@ $(OBJS) $(LDFLAGS) + +%.o:%.cu + $(CXX) $(CXXFLAGS) -c $< -o $@ + +clean: + rm -f hello poisson *.o *~ diff --git a/Serie09/double_buffer.cc b/Serie09/double_buffer.cc new file mode 100644 index 0000000..e6ecf0f --- /dev/null +++ b/Serie09/double_buffer.cc @@ -0,0 +1,19 @@ +/* -------------------------------------------------------------------------- */ +#include "double_buffer.hh" +#include "grid.hh" +/* -------------------------------------------------------------------------- */ + +/* -------------------------------------------------------------------------- */ +DoubleBuffer::DoubleBuffer(int m, int n) + : m_current(new Grid(m, n)), m_old(new Grid(m, n)) {} + +/* -------------------------------------------------------------------------- */ +Grid & DoubleBuffer::current() { return *m_current; } + +/* -------------------------------------------------------------------------- */ +Grid & DoubleBuffer::old() { return *m_old; } + +/* -------------------------------------------------------------------------- */ +void DoubleBuffer::swap() { + m_current.swap(m_old); +} diff --git a/Serie09/double_buffer.hh b/Serie09/double_buffer.hh new file mode 100644 index 0000000..67a1405 --- /dev/null +++ b/Serie09/double_buffer.hh @@ -0,0 +1,24 @@ +#ifndef DOUBLE_BUFFER +#define DOUBLE_BUFFER + +/* -------------------------------------------------------------------------- */ +#include +/* -------------------------------------------------------------------------- */ +#include "grid.hh" +/* -------------------------------------------------------------------------- */ + +class DoubleBuffer { +public: + DoubleBuffer(int m, int n); + + Grid & current(); + Grid & old(); + + void swap(); +private: + + std::unique_ptr m_current; + std::unique_ptr m_old; +}; + +#endif /* DOUBLE_BUFFER */ diff --git a/Serie09/dumpers.cc b/Serie09/dumpers.cc new file mode 100644 index 0000000..de7eb79 --- /dev/null +++ b/Serie09/dumpers.cc @@ -0,0 +1,120 @@ +/* -------------------------------------------------------------------------- */ +#include "dumpers.hh" +#include "grid.hh" +/* -------------------------------------------------------------------------- */ +#include +#include +#include +#include +/* -------------------------------------------------------------------------- */ + +/* -------------------------------------------------------------------------- */ +void Dumper::set_min(float min) { + m_min = min; +} + +void Dumper::set_max(float max) { + m_max = max; +} + +float Dumper::min() const { + return m_min; +} + +float Dumper::max() const { + return m_max; +} + +/* -------------------------------------------------------------------------- */ +void DumperASCII::dump(int step) { + std::ofstream fout; + std::stringstream sfilename; + + sfilename << "output/out_" << std::setfill('0') << std::setw(5) << step << ".pgm"; + fout.open(sfilename.str()); + + int m = m_grid.m(); + int n = m_grid.n(); + + fout << "P2" << std::endl << "# CREATOR: Poisson program" << std::endl; + fout << m << " " << n << std::endl; + fout << 255 << std::endl; + + for (int i = 0; i < m; i++) { + for (int j = 0; j < n; j++) { + int v = 255. * (m_grid(i, j) - m_min) / (m_max - m_min); + v = std::min(v, 255); + fout << v << std::endl; + } + } +} + +/* -------------------------------------------------------------------------- */ +void DumperBinary::dump(int step) { + std::ofstream fout; + std::stringstream sfilename; + + sfilename << "out_" << std::setfill('0') << std::setw(5) << step << ".bmp"; + fout.open(sfilename.str(), std::ios_base::binary); + + int h = m_grid.m(); + int w = m_grid.n(); + + int row_size = 3 * w; + // if the file width (3*w) is not a multiple of 4 adds enough bytes to make it + // a multiple of 4 + int padding = (4 - (row_size) % 4) % 4; + row_size += padding; + + int filesize = 54 + (row_size)*h; + + std::vector img(row_size*h); + std::fill(img.begin(), img.end(), 0); + + + for (int i = 0; i < h; i++) { + for (int j = 0; j < w; j++) { + float v = ((m_grid(h - 1 - i, j) - m_min) / (m_max - m_min)); + + float r = v * 255; // Red channel + float g = v * 255; // Green channel + float b = v * 255; // Red channel + + r = std::min(r, 255.f); + g = std::min(g, 255.f); + b = std::min(b, 255.f); + + img[row_size * i + 3 * j + 2] = r; + img[row_size * i + 3 * j + 1] = g; + img[row_size * i + 3 * j + 0] = b; + } + } + + std::array bmpfileheader = {'B', 'M', 0, 0, 0, 0, 0, + 0, 0, 0, 54, 0, 0, 0}; + std::array bmpinfoheader = {40, 0, 0, 0, 0, 0, 0, 0, + 0, 0, 0, 0, 1, 0, 24, 0}; + + bmpfileheader[2] = filesize; + bmpfileheader[3] = filesize >> 8; + bmpfileheader[4] = filesize >> 16; + bmpfileheader[5] = filesize >> 24; + + bmpinfoheader[4] = w; + bmpinfoheader[5] = w >> 8; + bmpinfoheader[6] = w >> 16; + bmpinfoheader[7] = w >> 24; + bmpinfoheader[8] = h; + bmpinfoheader[9] = h >> 8; + bmpinfoheader[10] = h >> 16; + bmpinfoheader[11] = h >> 24; + bmpinfoheader[20] = (filesize - 54); + bmpinfoheader[21] = (filesize - 54) >> 8; + bmpinfoheader[22] = (filesize - 54) >> 16; + bmpinfoheader[23] = (filesize - 54) >> 24; + + fout.write(bmpfileheader.data(), 14); + fout.write(bmpinfoheader.data(), 40); + + fout.write(img.data(), h * row_size); +} diff --git a/Serie09/dumpers.hh b/Serie09/dumpers.hh new file mode 100644 index 0000000..03a78c5 --- /dev/null +++ b/Serie09/dumpers.hh @@ -0,0 +1,41 @@ +#ifndef DUMPERS_HH +#define DUMPERS_HH + +/* -------------------------------------------------------------------------- */ +class Grid; + +/* -------------------------------------------------------------------------- */ +class Dumper { +public: + explicit Dumper(const Grid & grid) : m_grid(grid), m_min(-1.), m_max(1.) {} + + virtual void dump(int step) = 0; + + void set_min(float min); + void set_max(float max); + + float min() const; + float max() const; + +protected: + const Grid & m_grid; + float m_min, m_max; +}; + +/* -------------------------------------------------------------------------- */ +class DumperASCII : public Dumper { +public: + explicit DumperASCII(const Grid & grid) : Dumper(grid) {} + + virtual void dump(int step); +}; + +/* -------------------------------------------------------------------------- */ +class DumperBinary : public Dumper { +public: + explicit DumperBinary(const Grid & grid) : Dumper(grid) {} + + virtual void dump(int step); +}; + +#endif /* DUMPERS_HH */ diff --git a/Serie09/grid.hh b/Serie09/grid.hh new file mode 100644 index 0000000..f0cc1ce --- /dev/null +++ b/Serie09/grid.hh @@ -0,0 +1,52 @@ +#ifndef GRID_HH +#define GRID_HH + +/* -------------------------------------------------------------------------- */ +#include +#include +#include +/* -------------------------------------------------------------------------- */ + +class Grid { +public: + Grid(int m, int n): m_m(m), m_n(n) { + cudaMallocManaged(&m_storage, m * n * sizeof(float)); + clear(); + } + + Grid(const Grid & other) : + m_m{other.m_m}, m_n{other.m_n}, m_storage{other.m_storage}, copy{true} { } + + ~Grid() { + if (not copy) { + cudaFree(m_storage); + } + } + + /// access the value [i][j] of the grid + __host__ __device__ inline float & operator()(int i, int j) { + return m_storage[i * m_n + j]; + } + + __host__ __device__ inline const float & operator()(int i, int j) const { + return m_storage[i * m_n + j]; + } + + /// set the grid to 0 + __host__ void clear() { + std::fill(m_storage, m_storage + m_m * m_n, 0.); + } + + __host__ __device__ int m() const { return m_m; } + __host__ __device__ int n() const { return m_n; } + + + __host__ __device__ float* data() { return m_storage; } + __host__ __device__ float* data() const { return m_storage; } +private: + int m_m, m_n; + float * m_storage; + bool copy{false}; +}; + +#endif /* GRID_HH */ diff --git a/Serie09/output/.to-keep b/Serie09/output/.to-keep new file mode 100644 index 0000000..e69de29 diff --git a/Serie09/poisson.cc b/Serie09/poisson.cc new file mode 100644 index 0000000..5222d5a --- /dev/null +++ b/Serie09/poisson.cc @@ -0,0 +1,82 @@ +/* -------------------------------------------------------------------------- */ +#include "simulation.hh" +/* -------------------------------------------------------------------------- */ +#include +#include +#include +#include +#include +/* -------------------------------------------------------------------------- */ +#include +/* -------------------------------------------------------------------------- */ + +#define EPSILON 0.005 + +typedef std::chrono::high_resolution_clock clk; +typedef std::chrono::duration second; + +static void usage(const std::string & prog_name) { + std::cerr << prog_name << " " << std::endl; + exit(0); +} + +int main(int argc, char * argv[]) { + if (argc < 2) usage(argv[0]); + + int N; + try { + N = std::stoi(argv[1]); + } catch(std::invalid_argument &) { + usage(argv[0]); + } + + int block_size = 32; + if (argc == 3) { + try { + block_size = std::stoi(argv[2]); + } catch(std::invalid_argument &) { + usage(argv[0]); + } + } + + // By default, we use device 0, + int dev_id = 0; + + cudaDeviceProp device_prop; + cudaGetDevice(&dev_id); + cudaGetDeviceProperties(&device_prop, dev_id); + if (device_prop.computeMode == cudaComputeModeProhibited) { + std::cerr << "Error: device is running in , no " + "threads can use ::cudaSetDevice()" + << std::endl; + return -1; + } + + auto error = cudaGetLastError(); + if(error != cudaSuccess) { + throw std::runtime_error("Error Launching Kernel: " + + std::string(cudaGetErrorName(error)) + " - " + + std::string(cudaGetErrorString(error))); + } else { + std::cout << "GPU Device " << dev_id << ": \"" << device_prop.name + << "\" with compute capability " << device_prop.major << "." + << device_prop.minor << std::endl; + } + + + Simulation simu(N, N); + + simu.set_initial_conditions(); + + auto start = clk::now(); + int k = simu.compute(block_size); + auto end = clk::now(); + + second time = end - start; + + std::cout << block_size << " " << N << " " + << k << " " << std::scientific << " " + << time.count() << std::endl; + + return 0; +} diff --git a/Serie09/simulation.cc b/Serie09/simulation.cc new file mode 100644 index 0000000..91b6b70 --- /dev/null +++ b/Serie09/simulation.cc @@ -0,0 +1,37 @@ +/* -------------------------------------------------------------------------- */ +#include "simulation.hh" +/* -------------------------------------------------------------------------- */ +#include +#include +#include +/* -------------------------------------------------------------------------- */ + +/* -------------------------------------------------------------------------- */ +Simulation::Simulation(int m, int n) + : m_global_m(m), m_global_n(n), m_h_m(1. / m), + m_h_n(1. / n), m_grids(m, n), m_f(m, n), + m_dumper(new DumperBinary(m_grids.old())) {} + +/* -------------------------------------------------------------------------- */ +void Simulation::set_initial_conditions() { + for (int i = 0; i < m_global_m; i++) { + for (int j = 0; j < m_global_n; j++) { + m_f(i, j) = -2. * 100. * M_PI * M_PI * std::sin(10. * M_PI * i * m_h_m) * + std::sin(10. * M_PI * j * m_h_n); + } + } +} + +/* -------------------------------------------------------------------------- */ +int Simulation::compute(int block_size) { + int s; + + for(s =0; s < 1000; ++s) { + compute_step(block_size); + m_grids.swap(); + } + + m_dumper->dump(s); + return s; +} + diff --git a/Serie09/simulation.hh b/Serie09/simulation.hh new file mode 100644 index 0000000..556ad5a --- /dev/null +++ b/Serie09/simulation.hh @@ -0,0 +1,45 @@ +#ifndef SIMULATION_HH +#define SIMULATION_HH + +/* -------------------------------------------------------------------------- */ +#include "double_buffer.hh" +#include "dumpers.hh" +#include "grid.hh" +/* -------------------------------------------------------------------------- */ +#include +#include +/* -------------------------------------------------------------------------- */ + +/* -------------------------------------------------------------------------- */ +class Simulation { +public: + Simulation(int m, int n); + + /// set the initial conditions, Dirichlet and source term + virtual void set_initial_conditions(); + + /// perform the simulation + int compute(int block_size = 32); + +protected: + /// compute one step and return an error + void compute_step(int block_size); +private: + /// Global problem size + int m_global_m, m_global_n; + + /// grid spacing + float m_h_m; + float m_h_n; + + /// Grids storage + DoubleBuffer m_grids; + + /// source term + Grid m_f; + + /// Dumper to use for outputs + std::unique_ptr m_dumper; +}; + +#endif /* SIMULATION_HH */ diff --git a/Serie09/simulation_gpu.cu b/Serie09/simulation_gpu.cu new file mode 100644 index 0000000..129b374 --- /dev/null +++ b/Serie09/simulation_gpu.cu @@ -0,0 +1,37 @@ +/* -------------------------------------------------------------------------- */ +#include "simulation.hh" +#include "grid.hh" +/* -------------------------------------------------------------------------- */ +#include +#include +/* -------------------------------------------------------------------------- */ + +/* -------------------------------------------------------------------------- */ +__global__ void compute_step_one_thread_per_row( + Grid uo, Grid u, Grid f, float h) { + +} + +/* -------------------------------------------------------------------------- */ +__global__ void compute_step_one_thread_per_entry( + Grid uo, Grid u, Grid f, float h) { + +} + +/* -------------------------------------------------------------------------- */ +void Simulation::compute_step(int block_size) { + Grid & u = m_grids.current(); + Grid & uo = m_grids.old(); + + int m = u.m(); + int n = u.n(); + + // add the kernel call here + + auto error = cudaGetLastError(); + if(error != cudaSuccess) { + throw std::runtime_error("Error Launching Kernel: " + + std::string(cudaGetErrorName(error)) + " - " + + std::string(cudaGetErrorString(error))); + } +}