Unification of GPU kernels
git-svn-id: svn://scm.gforge.inria.fr/svnroot/paradiseo@2245 331e1502-861f-0410-8da2-ba01fb791d7f
This commit is contained in:
boufaras
parent
89f469a3f4
commit
054cfac03b
4 changed files with 288 additions and 0 deletions
71
branches/ParadisEO-GPU/src/eval/moGPUKernelEvalByCpy.h
Normal file
71
branches/ParadisEO-GPU/src/eval/moGPUKernelEvalByCpy.h
Normal file
|
|
@ -0,0 +1,71 @@
|
|||
/*
|
||||
<moGPUKernelEvalByCpy.h>
|
||||
Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
|
||||
|
||||
Karima Boufaras, Thé Van LUONG
|
||||
|
||||
This software is governed by the CeCILL license under French law and
|
||||
abiding by the rules of distribution of free software. You can use,
|
||||
modify and/ or redistribute the software under the terms of the CeCILL
|
||||
license as circulated by CEA, CNRS and INRIA at the following URL
|
||||
"http://www.cecill.info".
|
||||
|
||||
As a counterpart to the access to the source code and rights to copy,
|
||||
modify and redistribute granted by the license, users are provided only
|
||||
with a limited warranty and the software's author, the holder of the
|
||||
economic rights, and the successive licensors have only limited liability.
|
||||
|
||||
In this respect, the user's attention is drawn to the risks associated
|
||||
with loading, using, modifying and/or developing or reproducing the
|
||||
software by the user in light of its specific status of free software,
|
||||
that may mean that it is complicated to manipulate, and that also
|
||||
therefore means that it is reserved for developers and experienced
|
||||
professionals having in-depth computer knowledge. Users are therefore
|
||||
encouraged to load and test the software's suitability as regards their
|
||||
requirements in conditions enabling the security of their systems and/or
|
||||
data to be ensured and, more generally, to use and operate it in the
|
||||
same conditions as regards security.
|
||||
The fact that you are presently reading this means that you have had
|
||||
knowledge of the CeCILL license and that you accept its terms.
|
||||
|
||||
ParadisEO WebSite : http://paradiseo.gforge.inria.fr
|
||||
Contact: paradiseo-help@lists.gforge.inria.fr
|
||||
*/
|
||||
|
||||
#ifndef __moGPUKernelEvalByCpy_H
|
||||
#define __moGPUKernelEvalByCpy_H
|
||||
///////////////////////////////////////////////////////////////////////////////////////////////////////////////
|
||||
|
||||
/**
|
||||
* The kernel function called from the host and executed in device to compute all neighbors fitness at one time
|
||||
* without mapping, each thread id compute one fitness by copying localy the solution
|
||||
* @param _eval how to evaluate each neighbor
|
||||
* @param _solution the representation of solution( vector of int,float....)
|
||||
* @param _allFitness the array of Fitness to save all neighbors fitness
|
||||
* @param _fitness the current solution fitness
|
||||
* @param _neighborhoodsize the size of the neighborhood
|
||||
*/
|
||||
|
||||
template<class T, class Fitness, class Eval>
|
||||
|
||||
__global__ void moGPUKernelEvalByCpy(Eval _eval, T * _solution, Fitness* _allFitness,
|
||||
Fitness _fitness, unsigned _neighborhoodsize) {
|
||||
|
||||
// The thread identifier within a grid block's
|
||||
int id = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
// array to save index to be changed
|
||||
unsigned index[2];
|
||||
T sol_tmp[SIZE];
|
||||
// In this representation each id identify one and only one neighbor in neighborhood
|
||||
if (id < _neighborhoodsize) {
|
||||
for(unsigned i=0;i<SIZE;i++)
|
||||
sol_tmp[i]=_solution[i];
|
||||
//Change the id'th element of solution
|
||||
index[0]=id;
|
||||
index[1]=_neighborhoodsize;
|
||||
//Compute fitness for id'th neighbor
|
||||
_allFitness[id] = _eval(sol_tmp,_fitness,index);
|
||||
}
|
||||
}
|
||||
|
||||
#endif
|
||||
68
branches/ParadisEO-GPU/src/eval/moGPUKernelEvalByModif.h
Normal file
68
branches/ParadisEO-GPU/src/eval/moGPUKernelEvalByModif.h
Normal file
|
|
@ -0,0 +1,68 @@
|
|||
/*
|
||||
<moGPUKernelEvalByModif.h>
|
||||
Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
|
||||
|
||||
Karima Boufaras, Thé Van LUONG
|
||||
|
||||
This software is governed by the CeCILL license under French law and
|
||||
abiding by the rules of distribution of free software. You can use,
|
||||
modify and/ or redistribute the software under the terms of the CeCILL
|
||||
license as circulated by CEA, CNRS and INRIA at the following URL
|
||||
"http://www.cecill.info".
|
||||
|
||||
As a counterpart to the access to the source code and rights to copy,
|
||||
modify and redistribute granted by the license, users are provided only
|
||||
with a limited warranty and the software's author, the holder of the
|
||||
economic rights, and the successive licensors have only limited liability.
|
||||
|
||||
In this respect, the user's attention is drawn to the risks associated
|
||||
with loading, using, modifying and/or developing or reproducing the
|
||||
software by the user in light of its specific status of free software,
|
||||
that may mean that it is complicated to manipulate, and that also
|
||||
therefore means that it is reserved for developers and experienced
|
||||
professionals having in-depth computer knowledge. Users are therefore
|
||||
encouraged to load and test the software's suitability as regards their
|
||||
requirements in conditions enabling the security of their systems and/or
|
||||
data to be ensured and, more generally, to use and operate it in the
|
||||
same conditions as regards security.
|
||||
The fact that you are presently reading this means that you have had
|
||||
knowledge of the CeCILL license and that you accept its terms.
|
||||
|
||||
ParadisEO WebSite : http://paradiseo.gforge.inria.fr
|
||||
Contact: paradiseo-help@lists.gforge.inria.fr
|
||||
*/
|
||||
|
||||
#ifndef __moGPUKernelEvalByModif_H
|
||||
#define __moGPUKernelEvalByModif_H
|
||||
///////////////////////////////////////////////////////////////////////////////////////////////////////////////
|
||||
|
||||
/**
|
||||
* The kernel function called from the host and executed in device to compute all neighbors fitness at one time
|
||||
* without mapping, each thread id compute one fitness by modif of solution
|
||||
* @param _eval how to evaluate each neighbor
|
||||
* @param _solution the representation of solution( vector of int,float....)
|
||||
* @param _allFitness the array of Fitness to save all neighbors fitness
|
||||
* @param _fitness the current solution fitness
|
||||
* @param _neighborhoodsize the size of the neighborhood
|
||||
*/
|
||||
|
||||
template<class T, class Fitness, class Eval>
|
||||
|
||||
__global__ void moGPUKernelEvalByModif(Eval _eval, T * _solution, Fitness* _allFitness,
|
||||
Fitness _fitness, unsigned _neighborhoodsize) {
|
||||
|
||||
// The thread identifier within a grid block's
|
||||
int id = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
// array to save index to be changed
|
||||
unsigned index[2];
|
||||
// In this representation each id identify one and only one neighbor in neighborhood
|
||||
if (id < _neighborhoodsize) {
|
||||
//Change the id'th element of solution
|
||||
index[0]=id;
|
||||
index[1]=_neighborhoodsize;
|
||||
//Compute fitness for id'th neighbor
|
||||
_allFitness[id] = _eval(_solution,_fitness,index);
|
||||
}
|
||||
}
|
||||
|
||||
#endif
|
||||
|
|
@ -0,0 +1,76 @@
|
|||
/*
|
||||
<moGPUMappingKernelEvalByCpy.h>
|
||||
Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
|
||||
|
||||
Karima Boufaras, Thé Van LUONG
|
||||
|
||||
This software is governed by the CeCILL license under French law and
|
||||
abiding by the rules of distribution of free software. You can use,
|
||||
modify and/ or redistribute the software under the terms of the CeCILL
|
||||
license as circulated by CEA, CNRS and INRIA at the following URL
|
||||
"http://www.cecill.info".
|
||||
|
||||
As a counterpart to the access to the source code and rights to copy,
|
||||
modify and redistribute granted by the license, users are provided only
|
||||
with a limited warranty and the software's author, the holder of the
|
||||
economic rights, and the successive licensors have only limited liability.
|
||||
|
||||
In this respect, the user's attention is drawn to the risks associated
|
||||
with loading, using, modifying and/or developing or reproducing the
|
||||
software by the user in light of its specific status of free software,
|
||||
that may mean that it is complicated to manipulate, and that also
|
||||
therefore means that it is reserved for developers and experienced
|
||||
professionals having in-depth computer knowledge. Users are therefore
|
||||
encouraged to load and test the software's suitability as regards their
|
||||
requirements in conditions enabling the security of their systems and/or
|
||||
data to be ensured and, more generally, to use and operate it in the
|
||||
same conditions as regards security.
|
||||
The fact that you are presently reading this means that you have had
|
||||
knowledge of the CeCILL license and that you accept its terms.
|
||||
|
||||
ParadisEO WebSite : http://paradiseo.gforge.inria.fr
|
||||
Contact: paradiseo-help@lists.gforge.inria.fr
|
||||
*/
|
||||
|
||||
#ifndef __moGPUMappingKernelEvalByCpy_H
|
||||
#define __moGPUMappingKernelEvalByCpy_H
|
||||
///////////////////////////////////////////////////////////////////////////////////////////////////////////////
|
||||
|
||||
/**
|
||||
* The kernel function called from the host and executed in device to compute all neighbors fitness at one time
|
||||
* without mapping, each thread id compute one fitness by modif of solution
|
||||
* @param _eval how to evaluate each neighbor
|
||||
* @param _solution the representation of solution( vector of int,float....)
|
||||
* @param _allFitness the array of Fitness to save all neighbors fitness
|
||||
* @param _fitness the current solution fitness
|
||||
* @param _mapping associate to each threadID a set of correspondent indexes
|
||||
* @param _neighborhoodsize the size of the neighborhood
|
||||
*/
|
||||
|
||||
template<class T, class Fitness, class Eval>
|
||||
|
||||
__global__ void moGPUMappingKernelEvalByCpy(Eval _eval, T * _solution, Fitness* _allFitness,
|
||||
Fitness _fitness,unsigned * _mapping,unsigned _neighborhoodsize) {
|
||||
|
||||
// The thread identifier within a grid block's
|
||||
int id = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
//counter of number of x-change
|
||||
unsigned i;
|
||||
// array to save set a set of indexes corresponding to the current thread identifier
|
||||
unsigned index[NB_POS+2];
|
||||
T sol_tmp[SIZE];
|
||||
// In this representation each id identify one and only one neighbor in neighborhood
|
||||
if (id < _neighborhoodsize) {
|
||||
for(i=0;i<SIZE;i++)
|
||||
sol_tmp[i]=_solution[i];
|
||||
for(i=0;i<NB_POS;i++)
|
||||
index[i]=_mapping[id + i * _neighborhoodsize];
|
||||
index[NB_POS]=_neighborhoodsize;
|
||||
index[NB_POS+1]=id;
|
||||
//Evaluate by Modif Id'th neighbor with index mapping
|
||||
_allFitness[id]=_eval(sol_tmp,_fitness, index);
|
||||
|
||||
}
|
||||
}
|
||||
|
||||
#endif
|
||||
|
|
@ -0,0 +1,73 @@
|
|||
/*
|
||||
<moGPUMappingKernelEvalByModif.h>
|
||||
Copyright (C) DOLPHIN Project-Team, INRIA Lille - Nord Europe, 2006-2010
|
||||
|
||||
Karima Boufaras, Thé Van LUONG
|
||||
|
||||
This software is governed by the CeCILL license under French law and
|
||||
abiding by the rules of distribution of free software. You can use,
|
||||
modify and/ or redistribute the software under the terms of the CeCILL
|
||||
license as circulated by CEA, CNRS and INRIA at the following URL
|
||||
"http://www.cecill.info".
|
||||
|
||||
As a counterpart to the access to the source code and rights to copy,
|
||||
modify and redistribute granted by the license, users are provided only
|
||||
with a limited warranty and the software's author, the holder of the
|
||||
economic rights, and the successive licensors have only limited liability.
|
||||
|
||||
In this respect, the user's attention is drawn to the risks associated
|
||||
with loading, using, modifying and/or developing or reproducing the
|
||||
software by the user in light of its specific status of free software,
|
||||
that may mean that it is complicated to manipulate, and that also
|
||||
therefore means that it is reserved for developers and experienced
|
||||
professionals having in-depth computer knowledge. Users are therefore
|
||||
encouraged to load and test the software's suitability as regards their
|
||||
requirements in conditions enabling the security of their systems and/or
|
||||
data to be ensured and, more generally, to use and operate it in the
|
||||
same conditions as regards security.
|
||||
The fact that you are presently reading this means that you have had
|
||||
knowledge of the CeCILL license and that you accept its terms.
|
||||
|
||||
ParadisEO WebSite : http://paradiseo.gforge.inria.fr
|
||||
Contact: paradiseo-help@lists.gforge.inria.fr
|
||||
*/
|
||||
|
||||
#ifndef __moGPUMappingKernelEvalByModif_H
|
||||
#define __moGPUMappingKernelEvalByModif_H
|
||||
///////////////////////////////////////////////////////////////////////////////////////////////////////////////
|
||||
|
||||
/**
|
||||
* The kernel function called from the host and executed in device to compute all neighbors fitness at one time
|
||||
* without mapping, each thread id compute one fitness by modif of solution
|
||||
* @param _eval how to evaluate each neighbor
|
||||
* @param _solution the representation of solution( vector of int,float....)
|
||||
* @param _allFitness the array of Fitness to save all neighbors fitness
|
||||
* @param _fitness the current solution fitness
|
||||
* @param _mapping associate to each threadID a set of correspondent indexes
|
||||
* @param _neighborhoodsize the size of the neighborhood
|
||||
*/
|
||||
|
||||
template<class T, class Fitness, class Eval>
|
||||
|
||||
__global__ void moGPUMappingKernelEvalByModif(Eval _eval, T * _solution, Fitness* _allFitness,
|
||||
Fitness _fitness,unsigned int * _mapping,unsigned _neighborhoodsize) {
|
||||
|
||||
// The thread identifier within a grid block's
|
||||
int id = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
//counter of number of x-change
|
||||
unsigned i;
|
||||
// array to save set a set of indexes corresponding to the current thread identifier
|
||||
unsigned index[NB_POS+2];
|
||||
// In this representation each id identify one and only one neighbor in neighborhood
|
||||
if (id < _neighborhoodsize) {
|
||||
for(i=0;i<NB_POS;i++)
|
||||
index[i]=_mapping[id + i * _neighborhoodsize];
|
||||
index[NB_POS]=_neighborhoodsize;
|
||||
index[NB_POS+1]=id;
|
||||
//Evaluate by Modif Id'th neighbor with index mapping
|
||||
_allFitness[id]=_eval(_solution,_fitness, index);
|
||||
|
||||
}
|
||||
}
|
||||
|
||||
#endif
|
||||
Loading…
Add table
Add a link
Reference in a new issue