su2hmc
Loading...
Searching...
No Matches
sizes.h
Go to the documentation of this file.
1
28#pragma once
29#ifdef __USE_MKL__
31#define USE_BLAS
32#include <mkl.h>
33#elif (defined AMD_BLAS||defined OPENBLAS)
35#define USE_BLAS
36#include <cblas.h>
37#elif defined GSL_BLAS
39#define USE_BLAS
40#include <gsl/gsl_cblas.h>
41#endif
42#ifdef USE_GPU
43#include <cuda.h>
44#include <cuda_runtime_api.h>
45#include <cublas_v2.h>
46#ifndef __BSIZE__
48#define __BSIZE__ 128
49#endif
51extern cublasHandle_t cublas_handle;
53extern cublasStatus_t cublas_status;
55extern cudaMemPool_t mempool;
57#define cudaDeviceSynchronise() cudaDeviceSynchronize()
58#endif
59#ifdef __CUDACC__
60#include <thrust_complex.h>
61#include <thrust/reduce.h>
62#include <thrust/device_vector.h>
63#else
64#include <complex.h>
66#define Complex_f float complex
68#define Complex double complex
69#endif
70
72#define FILELEN 64
73// Common block definition for parallel variables
74
76#define nx 24
77#if(nx<1)
78#error "nx is expected it to be greater than or equal to 1"
79#endif
80
81// Keep original restriction of single spatial extent
82
84#define ny nx
85#if(ny<1)
86#error "ny is expected it to be greater than or equal to 1"
87#endif
88
90#define nz nx
91#if(nz<1)
92#error "nz is expected it to be greater than or equal to 1"
93#endif
94
96#define nt 24
97#if(nt<1)
98#error "nt is expected it to be greater than or equal to 1"
99#endif
100
102#define gvol (nx*ny*nz*nt)
104#define gvol3 (nx*ny*nz)
105
107#define npx 1
108#if(npx<1)
109#error "npx is expected it to be greater than or equal to 1"
110#elif(nx%npx!=0)
111#error "npx should be a divisor of nx"
112#endif
113
114// Initially restrict to npz = npy = npx
115// This allows us to have a single ksize variable
116
118#define npy npx
119#if(npy<1)
120#error "npy is expected it to be greater than or equal to 1"
121#elif(ny%npy!=0)
122#error "npy should be a divisor of ny"
123#endif
124
126#define npz npx
127#if(npz<1)
128#error "npz is expected it to be greater than or equal to 1"
129#elif(nz%npz!=0)
130#error "npz should be a divisor of nz"
131#endif
132
134#define npt 1
135#if(npt<1)
136#error "npt is expected it to be greater than or equal to 1"
137#elif(nt%npt!=0)
138#error "npt should be a divisor of nt"
139#endif
140
142#define nproc (npx*npy*npz*npt)
143#if (defined USE_GPU && nproc>1)
144#error "Multi-GPU is not yet supported"
145#endif
146
148#define nthreads 8
149
150// Existing parameter definitions.
152#define ksizex (nx/npx)
154#define ksizey (ny/npy)
156#define ksizez (nz/npz)
157
159#define ksize ksizex
160
162#define ksizet (nt/npt)
164#define nf 1
165
167#define kvol (ksizet*ksizez*ksizey*ksizex)
169#define kvol3 (ksizez*ksizey*ksizex)
170
171// integer, parameter :: niterc=2*gvol
172// #define niterc 2*gvol
173// jis: hard limit to avoid runaway trajectories
174#if (nx*ny*nz*nt<=16384)
176#define niterc gvol
177#elif (nx>=(3*nt)/2)
179#define niterc gvol3
180#else
182#define niterc (gvol/4)
183#endif
184// Constants for dimensions.
186#define nc 2
188#define nadj 3
190#define ndirac 4
192#define ndim 4
194#define ngorkov 8
195
197#define kmom (ndim*nadj*kvol)
199#define kferm (nc*ngorkov*kvol)
201#define kferm2 (nc*ndirac*kvol)
211#if(npx>1)
212#define halox (ksizey*ksizez*ksizet)
213#else
214#define halox 0
215#endif
217#if(npy>1)
218#define haloy (ksizex*ksizez*ksizet)
219#else
220#define haloy 0
221#endif
223#if(npz>1)
224#define haloz (ksizex*ksizey*ksizet)
225#else
226#define haloz 0
227#endif
229#if(npt>1)
230#define halot (ksizex*ksizey*ksizez)
231#else
232#define halot 0
233#endif
235#define halo (2*(halox+haloy+haloz+halot))
236
238#define kvolHalo (kvol+halo)
240#define kfermHalo (nc*ngorkov*kvolHalo)
242#define kferm2Halo (nc*ndirac*kvolHalo)
244#define kmomHalo (ndim*nadj*kvolHalo)
245
246// These all used to be multipled by kferm or kferm2 at the start of Congradq or Congradp
247// On 20240516 in Room 2.19 of the Lloyd building of Trinity we copped that doing so means that the residue is larger
248// if running on a smaller number of cores. In the extreme GPU case the subvolume is the entire volume so the residue
249// can be several orders of magnitude larger than in the smallest sublattice case.
250// Instead, we rescale all the default residues here by sqrt(kferm) or sqrt(kferm2). No matter what size sublattice
251// we use now, the residue will match that of a 2^3X4 sublattice used in the earlier FORTRAN runs
253#define respbp 3.2E-5
255#define rescgg 2.26E-5
257#define rescga 2.26E-8
258
259
260#ifdef __AVX512F__
261#ifdef __unix__
262#warning AVX512 detected
263#elif (defined WIN32||_WIN32)
264#pragma message("AVX512 detected")
265#endif
267#define AVX 64
268#elif defined __AVX__
269#ifdef __unix__
270#warning AVX or AVX2 detected
271#elif (defined WIN32||_WIN32)
272#pragma message("AVX or AVX2 detected")
273#endif
275#define AVX 32
276#else
277#ifdef __unix__
278#warning No AVX detected, assuming SSE is present
279#elif (defined WIN32||_WIN32)
280#pragma message("No AVX detected, assuming SSE is present")
281#endif
283#define AVX 16
284#endif
285
286#ifdef USE_GPU
298extern dim3 dimBlock;
301extern dim3 dimGrid;
303extern dim3 dimBlockOne;
305extern dim3 dimGridOne;
307#define USE_BLAS
308#endif
cudaMemPool_t mempool
Memory pool for Async allocations.
Definition main.c:49
cublasStatus_t cublas_status
Status of cuBLAS for error reporting.
Definition main.c:48
dim3 dimBlockOne
block size of one
Definition cusu2hmc.cu:22
dim3 dimGridOne
Grid size of one.
Definition cusu2hmc.cu:23
cublasHandle_t cublas_handle
Handle for cuBLAS.
Definition main.c:47
dim3 dimGrid
Default grid size. First component is normally nt. Second and third depend whatever is needed to get ...
Definition cusu2hmc.cu:27
dim3 dimBlock
Default block size. Usually 128.
Definition cusu2hmc.cu:25
Complex Header for CUDA. Sets macros for C compatability.