portelli/Grid
1
0
Fork 0
mirror of https://github.com/paboyle/Grid.git synced 2025-04-12 07:00:45 +01:00
Code Releases Activity

MADWF 5d source option for hadrons - look at Grid of source

Browse Source 
Abort on GPU error
This commit is contained in:
Peter Boyle 2021-02-08 10:47:22 -05:00
parent cd99edcc5f
commit eda9ab487b
3 changed files with 15 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

14
Grid/qcd/action/fermion/MADWF.h
View File

@ -85,7 +85,7 @@ class MADWF
maxiter =_maxiter; maxiter =_maxiter;
}; };
void operator() (const FermionFieldo &src4,FermionFieldo &sol5) void operator() (const FermionFieldo &src,FermionFieldo &sol5)
{ {
std::cout << GridLogMessage<< " ************************************************" << std::endl; std::cout << GridLogMessage<< " ************************************************" << std::endl;
std::cout << GridLogMessage<< " MADWF-like algorithm " << std::endl; std::cout << GridLogMessage<< " MADWF-like algorithm " << std::endl;
@ -114,8 +114,16 @@ class MADWF
/////////////////////////////////////// ///////////////////////////////////////
//Import source, include Dminus factors //Import source, include Dminus factors
/////////////////////////////////////// ///////////////////////////////////////
Mato.ImportPhysicalFermionSource(src4,b); GridBase *src_grid = src.Grid();
std::cout << GridLogMessage << " src4 " <<norm2(src4)<<std::endl;
assert( (src_grid == Mato.GaugeGrid()) || (src_grid == Mato.FermionGrid()));
if ( src_grid == Mato.GaugeGrid() ) {
Mato.ImportPhysicalFermionSource(src,b);
} else {
b=src;
}
std::cout << GridLogMessage << " src " <<norm2(src)<<std::endl;
std::cout << GridLogMessage << " b " <<norm2(b)<<std::endl; std::cout << GridLogMessage << " b " <<norm2(b)<<std::endl;
defect = b; defect = b;

1
Grid/threads/Accelerator.cc
View File

@ -1,6 +1,7 @@
#include <Grid/GridCore.h> #include <Grid/GridCore.h>
NAMESPACE_BEGIN(Grid); NAMESPACE_BEGIN(Grid);
int acceleratorAbortOnGpuError=1;
uint32_t accelerator_threads=2; uint32_t accelerator_threads=2;
uint32_t acceleratorThreads(void) {return accelerator_threads;}; uint32_t acceleratorThreads(void) {return accelerator_threads;};
void acceleratorThreads(uint32_t t) {accelerator_threads = t;}; void acceleratorThreads(uint32_t t) {accelerator_threads = t;};

3
Grid/threads/Accelerator.h
View File

@ -100,6 +100,8 @@ void acceleratorInit(void);
#define accelerator __host__ __device__ #define accelerator __host__ __device__
#define accelerator_inline __host__ __device__ inline #define accelerator_inline __host__ __device__ inline
extern int acceleratorAbortOnGpuError;
accelerator_inline int acceleratorSIMTlane(int Nsimd) { accelerator_inline int acceleratorSIMTlane(int Nsimd) {
#ifdef GRID_SIMT #ifdef GRID_SIMT
return threadIdx.z; return threadIdx.z;
@ -140,6 +142,7 @@ void LambdaApply(uint64_t num1, uint64_t num2, uint64_t num3, lambda Lambda)
printf("Cuda error %s \n", cudaGetErrorString( err )); \ printf("Cuda error %s \n", cudaGetErrorString( err )); \
puts(__FILE__); \ puts(__FILE__); \
printf("Line %d\n",__LINE__); \ printf("Line %d\n",__LINE__); \
if (acceleratorAbortOnGpuError) assert(err==cudaSuccess); \
} \ } \
} }