helpers.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 "device_settings.hpp"
#include "utilities.hpp"


template<int Begin, int End, int Step = 1>
struct Unroller {
        template<typename Action>
        __device__      static void step(const Action& action) {
                action(Begin);
                Unroller<Begin+Step, End, Step>::step(action);
        }
};

template<int End, int Step>
struct Unroller<End, End, Step> {
        template<typename Action>
        __device__      static void step(const Action& action) { }
};


template<template<int> class T,int N,int MAXN,class B,class P>
struct choose {
B operator ()(const int& n, const P& x){
        if(n == N)
                return T<N>::choose(x);
        else if(n <= MAXN && n > N)
                return choose<T,(N<MAXN ? N+1 : MAXN),MAXN,B,P>()(n,x);
        else
                return B();
}
};

namespace swarm {


template<int i>
struct compile_time_params_t {
        const static int n = i;
};


template<class implementation,class T>
__global__ void generic_kernel(implementation* integ,T compile_time_param) {
        integ->kernel(compile_time_param);
}

template< class implementation, class T>
void launch_template(implementation* integ, implementation* gpu_integ, T compile_time_param)
{
        if(integ->get_ensemble().nbod() == T::n) 
                generic_kernel<<<integ->gridDim(), integ->threadDim(), integ->shmemSize() >>>(gpu_integ,compile_time_param);
        else
          ERROR("Error launching kernel.  Active ensemble has " + inttostr(integ->get_ensemble().nbod()) + " bodies per system.\n");
                
}


template<int N>
struct launch_template_choose {
        template<class integ_pair>
        static void choose(integ_pair p){
                compile_time_params_t<N> ctp;
                typename integ_pair::first_type integ = p.first;

                int sys_p_block = integ->override_system_per_block();
                const int nsys = integ->get_ensemble().nsys();
                const int tps = integ->thread_per_system(ctp);
                const int shm = integ->shmem_per_system(ctp);
                if(sys_p_block == 0){
                        sys_p_block = optimized_system_per_block(SHMEM_CHUNK_SIZE, tps, shm);
                }


                const int nblocks = ( nsys + sys_p_block - 1 ) / sys_p_block;
                const int shmemSize = sys_p_block * shm;

                dim3 gridDim;
                find_best_factorization(gridDim.x,gridDim.y,nblocks);

                dim3 threadDim;
                threadDim.x = sys_p_block;
                threadDim.y = tps;

                int blocksize = threadDim.x * threadDim.y;
                if(!check_cuda_limits(blocksize, shmemSize )){
                        throw runtime_error("The block size settings exceed CUDA requirements");
                }

                generic_kernel<<<gridDim, threadDim, shmemSize>>>(p.second, ctp);
        }
};


template<class implementation>
void launch_templatized_integrator(implementation* integ){

        if(integ->get_ensemble().nbod() <= MAX_NBODIES){
                implementation* gpu_integ;
                cudaErrCheck ( cudaMalloc(&gpu_integ,sizeof(implementation)) );
                cudaErrCheck ( cudaMemcpy(gpu_integ,integ,sizeof(implementation),cudaMemcpyHostToDevice) );

                typedef std::pair<implementation*,implementation*> integ_pair ;
                integ_pair p ( integ, gpu_integ );
                int nbod = integ->get_ensemble().nbod();

                choose< launch_template_choose, 3, MAX_NBODIES, void, integ_pair > c;
                        c( nbod, p );

                cudaFree(gpu_integ);
        } else {
                char b[100];
                snprintf(b,100,"Invalid number of bodies. (Swarm-NG was compiled with MAX_NBODIES = %d bodies per system.)",MAX_NBODIES);
                ERROR(b);
        }

}


        
}