bppt.hpp

Go to the documentation of this file.

/*************************************************************************
 * Copyright (C) 2011 by Saleh Dindar and the Swarm-NG Development Team  *
 *                                                                       *
 * This program is free software; you can redistribute it and/or modify  *
 * it under the terms of the GNU General Public License as published by  *
 * the Free Software Foundation; either version 3 of the License.        *
 *                                                                       *
 * This program is distributed in the hope that it will be useful,       *
 * but WITHOUT ANY WARRANTY; without even the implied warranty of        *
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the         *
 * GNU General Public License for more details.                          *
 *                                                                       *
 * You should have received a copy of the GNU General Public License     *
 * along with this program; if not, write to the                         *
 * Free Software Foundation, Inc.,                                       *
 * 59 Temple Place - Suite 330, Boston, MA  02111-1307, USA.             *
 ************************************************************************/

#pragma once

#include "../integrator.hpp"
#include "../log/log.hpp"
#include "../plugin.hpp"

#include "helpers.hpp"
#include "utilities.hpp"
#include "device_settings.hpp"


namespace swarm {
namespace gpu {

namespace bppt {


GPUAPI int sysid(){
        return ((blockIdx.z * gridDim.y + blockIdx.y) * gridDim.x + blockIdx.x) * blockDim.x + threadIdx.x;
}

GPUAPI int sysid_in_block(){
        return threadIdx.x;
}

GPUAPI int thread_in_system() {
        return threadIdx.y;
}


GPUAPI int system_per_block_gpu() {
    return blockDim.x;
  };

GPUAPI int thread_component_idx(int nbod) {
        return thread_in_system() / nbod;
}

GPUAPI int thread_body_idx(int nbod) {
        return thread_in_system() % nbod;
}


template< class Impl, class T> 
GPUAPI void * system_shared_data_pointer(Impl* integ, T compile_time_param) {
        extern __shared__ char shared_mem[];
        int b = sysid_in_block() / SHMEM_CHUNK_SIZE ;
        int i = sysid_in_block() % SHMEM_CHUNK_SIZE ;
        int idx = i * sizeof(double) 
                + b * SHMEM_CHUNK_SIZE 
                * Impl::shmem_per_system(compile_time_param);
        return &shared_mem[idx];
}



class integrator : public gpu::integrator  {
        typedef gpu::integrator Base;
        protected:
        int _override_system_per_block;

        public:
        integrator(const config &cfg) : Base(cfg) {
                int spb = cfg.optional("system_per_block",0);
                _override_system_per_block =(spb % SHMEM_CHUNK_SIZE == 0) ? spb : 0;
        }

         
        const int& override_system_per_block()const{
                return _override_system_per_block;
        }

        template<class T>
        static GENERIC int thread_per_system(T compile_time_param){
                const int nbod = T::n;
                const int body_comp = nbod * 3;
                const int pair_count = nbod * (nbod - 1) / 2;
                return (body_comp>pair_count) ? body_comp : pair_count;
        }


        template<class T>
        static GENERIC int shmem_per_system(T compile_time_param){
                const int nbod = T::n;
                const int pair_count = nbod * (nbod - 1) / 2;
                return pair_count * 3  * 2 * sizeof(double);
        }


};

        
}
}
}