Homec4science
Diffusion Lenstool-HPC (master) Browse

R1448/Benchmarks/GridGradientBenchmarkMixed/grid_gradient_mixed_GPU.cuh
fd9c9e4de80amaster

Last Change
Raw File

grid_gradient_mixed_GPU.cuh
View Options

/*
* grid_gradient_CPU.hpp
*
* Created on: Jan 12, 2017
* Author: cerschae
* Author: gilles.fourestey@epfl.ch
*/
#ifndef GRID_GRADIENT_GPU_HPP_
#define GRID_GRADIENT_GPU_HPP_
#warning "including griad_gradient_GPU"
//
#include <string.h>
//#include <cuda_runtime.h>
#include <math.h>
#include <gradient_avx.hpp>
#include <sys/time.h>
#include <fstream>
#include <iostream>
//
#include "structure_hpc.hpp"
#include "structure_mixed.hpp"
//#include <immintrin.h>
//
//
//
#ifdef __WITH_NVTX
#include "nvToolsExt.h"
const uint32_t colors[] = { 0x0000ff00, 0x000000ff, 0x00ffff00, 0x00ff00ff, 0x0000ffff, 0x00ff0000, 0x00ffffff };
const int num_colors = sizeof(colors)/sizeof(uint32_t);
#define PUSH_RANGE(name,cid) { \
int color_id = cid; \
color_id = color_id%num_colors;\
nvtxEventAttributes_t eventAttrib = {0}; \
eventAttrib.version = NVTX_VERSION; \
eventAttrib.size = NVTX_EVENT_ATTRIB_STRUCT_SIZE; \
eventAttrib.colorType = NVTX_COLOR_ARGB; \
eventAttrib.color = colors[color_id]; \
eventAttrib.messageType = NVTX_MESSAGE_TYPE_ASCII; \
eventAttrib.message.ascii = name; \
nvtxRangePushEx(&eventAttrib); \
}
#define POP_RANGE nvtxRangePop();
#else
#define PUSH_RANGE(name,cid)
#define POP_RANGE
#endif
//
//
//
#define CudaCheckError() __cudaCheckError( __FILE__, __LINE__ )
//
inline void __cudaCheckError( const char *file, const int line )
{
#ifdef CUDA_ERROR_CHECK
cudaError err = cudaGetLastError();
if ( cudaSuccess != err )
{
fprintf( stderr, "cudaCheckError() failed at %s:%i : %s\n",
file, line, cudaGetErrorString( err ) );
exit( -1 );
}
// More careful checking. However, this will affect performance.
// Comment away if needed.
err = cudaDeviceSynchronize();
if( cudaSuccess != err )
{
fprintf( stderr, "cudaCheckError() with sync failed at %s:%i : %s\n",
file, line, cudaGetErrorString( err ) );
exit( -1 );
}
#endif
return;
}
//
//
//
#define CudaSafeCall( err ) __cudaSafeCall( err, __FILE__, __LINE__ )
//
inline void __cudaSafeCall( cudaError err, const char *file, const int line )
{
#ifdef CUDA_ERROR_CHECK
if ( cudaSuccess != err )
{
fprintf( stderr, "cudaSafeCall() failed at %s:%i : %s\n",
file, line, cudaGetErrorString( err ) );
exit( -1 );
}
#endif
return;
}
extern "C"{
static
void cudasafer( cudaError_t error, char* message)
{
if(error!=cudaSuccess) { fprintf(stderr,"ERROR: %s : %i\n",message,error); exit(-1); }
};
}
//
//
/*
struct point_double
{
type_t x;
type_t y;
};
*/
//
//
//
void gradient_grid_GPU(type_t *grid_grad_x, type_t *grid_grad_y, const struct grid_param *frame, const struct Potential_SOA *lens, int Nlens, int nbgridcells, int istart = 0, int jstart = 0);
//
void gradient_grid_GPU(type_t *grid_grad_x, type_t *grid_grad_y, const struct grid_param *frame, const struct Potential_SOA *lens, int Nlens, int nbgridcells_x, int nbgridcells_y, int istart = 0, int jstart = 0);
//
static void gradient_grid_general_CPU_old(type_t *grid_grad_x, type_t *grid_grad_y, const struct grid_param *frame, int Nlens, int nbgridcells, const struct Potential_SOA *lens);
//
void gradient_grid_mixed_GPU(double *grid_grad_x, double *grid_grad_y,
const struct grid_param *frame,
const struct Potential_SOA *lens,
int nhalos,
int nbgridcells_x, int nbgridcells_y,
int istart, int jstart);
//
void gradient_grid_general_mixed_GPU(double *grid_grad_x, double *grid_grad_y, const struct grid_param *frame, int Nlens, int nbgridcells, const struct Potential_SOA *lens, int istart, int jstart);
void
gradient_grid_general_double_GPU(double *grid_grad_x, double *grid_grad_y, const struct grid_param *frame, const struct Potential_SOA *lens, int nhalos, int nbgridcells_x, int nbgridcells_y, int istart, int jstart);
//
void
gradient_grid_general_double_GPU(double *grid_grad_x, double *grid_grad_y, const struct grid_param *frame, const struct Potential_Mixed<double> *lens, int nhalos, int nbgridcells_x, int nbgridcells_y, int istart, int jstart);
//
__global__
void
gradient_grid_general_SP_GPU(double *grid_grad_x, double *grid_grad_y, const struct grid_param *frame, const struct Potential_SOA *lens, int nhalos, int nbgridcells_x, int nbgridcells_y, int istart, int jstart);
//
void
gradient_grid_general_SP_GPU(float* __restrict__ grid_grad_x, float* __restrict__ grid_grad_y, const struct grid_param *frame, const struct Potential_SOA *lens, int Nlens, int nbgridcells_x, int nbgridcells_y, int istart, int jstart);
//
template<typename fp_type = double>
void
gpu_allocate_SOA(Potential_Mixed<fp_type> *lens_gpu, const Potential_SOA *lens, const struct grid_param *frame, int nhalos, int nbgridcells_x, int nbgridcells_y)
{
//
lens_gpu = (Potential_Mixed<fp_type> *) malloc(sizeof(Potential_SOA));
//cudasafer(cudaMalloc( (void**)&(lens_kernel) , sizeof(Potential_SOA)),"Gradientgpu.cu : Alloc Potential_SOA: " );
grid_param *frame_gpu;
int *type_gpu;
fp_type *lens_x_gpu, *lens_y_gpu, *b0_gpu, *angle_gpu, *epot_gpu, *rcore_gpu, *rcut_gpu, *anglecos_gpu, *anglesin_gpu;
//
cudasafer(cudaMalloc( (void**)&(type_gpu) , nhalos*sizeof(int)) ,"Gradientgpu.cu : Alloc type_gpu: " );
cudasafer(cudaMalloc( (void**)&(lens_x_gpu) , nhalos*sizeof(fp_type)),"Gradientgpu.cu : Alloc x_gpu: " );
cudasafer(cudaMalloc( (void**)&(lens_y_gpu) , nhalos*sizeof(fp_type)),"Gradientgpu.cu : Alloc y_gpu: " );
cudasafer(cudaMalloc( (void**)&(b0_gpu) , nhalos*sizeof(fp_type)),"Gradientgpu.cu : Alloc b0_gpu: " );
cudasafer(cudaMalloc( (void**)&(angle_gpu) , nhalos*sizeof(fp_type)),"Gradientgpu.cu : Alloc angle_gpu: " );
cudasafer(cudaMalloc( (void**)&(epot_gpu) , nhalos*sizeof(fp_type)),"Gradientgpu.cu : Alloc epot_gpu: " );
cudasafer(cudaMalloc( (void**)&(rcore_gpu) , nhalos*sizeof(fp_type)),"Gradientgpu.cu : Alloc rcore_gpu: " );
cudasafer(cudaMalloc( (void**)&(rcut_gpu) , nhalos*sizeof(fp_type)),"Gradientgpu.cu : Alloc rcut_gpu: " );
cudasafer(cudaMalloc( (void**)&(anglecos_gpu), nhalos*sizeof(fp_type)),"Gradientgpu.cu : Alloc anglecos_gpu: " );
cudasafer(cudaMalloc( (void**)&(anglesin_gpu), nhalos*sizeof(fp_type)),"Gradientgpu.cu : Alloc anglesin_gpu: " );
cudasafer(cudaMalloc( (void**)&(frame_gpu) , sizeof(grid_param)) ,"Gradientgpu.cu : Alloc frame_gpu: " );
//
//double* grid_grad_temp_x; cudasafer(cudaMallocManaged((void**) &grid_grad_temp_x , nbgridcells_x*nbgridcells_y*sizeof(double)), "grid_grad_temp_x");
//double* grid_grad_temp_y; cudasafer(cudaMallocManaged((void**) &grid_grad_temp_y , nbgridcells_x*nbgridcells_y*sizeof(double)), "grid_grad_temp_y");
//
//double* grid_grad_euler_x; cudasafer(cudaMallocManaged((void**) &grid_grad_euler_x, nbgridcells_x*nbgridcells_y*sizeof(double)), "grid_grad_euler_x");
//double* grid_grad_euler_y; cudasafer(cudaMallocManaged((void**) &grid_grad_euler_y, nbgridcells_x*nbgridcells_y*sizeof(double)), "grid_grad_euler_y");
//
// Copy values to the GPU
//
cudasafer(cudaMemcpy(type_gpu , lens->type , nhalos*sizeof(int) , cudaMemcpyHostToDevice) , "Gradientgpu.cu : Copy type_gpu: " );
cudasafer(cudaMemcpy(lens_x_gpu , lens->position_x , nhalos*sizeof(fp_type), cudaMemcpyHostToDevice), "Gradientgpu.cu : Copy x_gpu: " );
cudasafer(cudaMemcpy(lens_y_gpu , lens->position_y , nhalos*sizeof(fp_type), cudaMemcpyHostToDevice), "Gradientgpu.cu : Copy y_gpu: " );
cudasafer(cudaMemcpy(b0_gpu , lens->b0 , nhalos*sizeof(fp_type), cudaMemcpyHostToDevice), "Gradientgpu.cu : Copy b0_gpu: " );
cudasafer(cudaMemcpy(angle_gpu , lens->ellipticity_angle , nhalos*sizeof(fp_type), cudaMemcpyHostToDevice), "Gradientgpu.cu : Copy angle_gpu: " );
cudasafer(cudaMemcpy(epot_gpu , lens->ellipticity_potential, nhalos*sizeof(fp_type), cudaMemcpyHostToDevice), "Gradientgpu.cu : Copy epot_gpu: " );
cudasafer(cudaMemcpy(rcore_gpu , lens->rcore , nhalos*sizeof(fp_type), cudaMemcpyHostToDevice), "Gradientgpu.cu : Copy rcore_gpu: " );
cudasafer(cudaMemcpy(rcut_gpu , lens->rcut , nhalos*sizeof(fp_type), cudaMemcpyHostToDevice), "Gradientgpu.cu : Copy rcut_gpu: " );
cudasafer(cudaMemcpy(anglecos_gpu , lens->anglecos , nhalos*sizeof(fp_type), cudaMemcpyHostToDevice), "Gradientgpu.cu : Copy anglecos: " );
cudasafer(cudaMemcpy(anglesin_gpu , lens->anglesin , nhalos*sizeof(fp_type), cudaMemcpyHostToDevice), "Gradientgpu.cu : Copy anglesin: " );
cudasafer(cudaMemcpy(frame_gpu , frame , sizeof(grid_param), cudaMemcpyHostToDevice) , "Gradientgpu.cu : Copy fame_gpu: " );
//
lens_gpu->type = type_gpu;
lens_gpu->position_x = lens_x_gpu;
lens_gpu->position_y = lens_y_gpu;
lens_gpu->b0 = b0_gpu;
lens_gpu->ellipticity_angle = angle_gpu;
lens_gpu->ellipticity_potential = epot_gpu;
lens_gpu->rcore = rcore_gpu;
lens_gpu->rcut = rcut_gpu;
lens_gpu->anglecos = anglecos_gpu;
lens_gpu->anglesin = anglesin_gpu;
//
//cudaMemcpy(lens_kernel, lens_gpu, sizeof(Potential_SOA), cudaMemcpyHostToDevice);
//
};
#endif /* GRID_GRADIENT_CPU_HPP_ */
c4science · Help