diff options
Diffstat (limited to 'benchmarks/CUDA/WP/spt.h')
| -rw-r--r-- | benchmarks/CUDA/WP/spt.h | 150 |
1 files changed, 0 insertions, 150 deletions
diff --git a/benchmarks/CUDA/WP/spt.h b/benchmarks/CUDA/WP/spt.h deleted file mode 100644 index 27a1467..0000000 --- a/benchmarks/CUDA/WP/spt.h +++ /dev/null @@ -1,150 +0,0 @@ -// macros and whatnot for translator - -#ifdef CUDA - -#if 0 -Types - -1. Array stored in device memory - - - 1.a fully dimensional - Example 3D - name: qc - how allocated: argument - dimensionality: ims:ime,kms:kme,jms:jme - index as: P3(ti,k,tj) - - Example 2D - name: qc - how allocated: argument - dimensionality: ims:ime,jms:jme - index as: P2(ti,k,tj) - - - 1.b 1 dimensional (vertical only), local storage - Example - name: w3 - how allocated: local - dimensionality: constant (MKX) - index as: w3[k] - -2. Array stored in shared memory - - 2.a fully dimensional - - 3D - Example - name: qc_s - how allocated: as offsets into sm[SM_SIZE] - dimensionality: bx*by*kx - index as: S3(ti,k,tj) - - 2D - Example - name: <none>_s - how allocated: as offsets into sm[SM_SIZE] - dimensionality: bx*by - index as: S2(ti,tj) - -#endif - -#ifndef MKX - -- intentional syntax error -- need a defined constant MKX that is static number of levels -- -#endif - -#ifdef PROMOTE -# define float double -#endif - -#define SM_SIZE (0x1000-0xd4) -#define MAX_THREADS_PER_BLOCK 512 - -#define bi blockIdx.x -#define bj blockIdx.y -#define bx blockDim.x -#define by blockDim.y -#define ti threadIdx.x -#define tj threadIdx.y - -# define ix (ime-ims+1) -# define jx (jme-jms+1) -# define kx (kme-kms+1) - - -// basic indexing macros. indices are always given as global indices -// in undecompsed Domain(ids:ide,jds:jde) -// -// That is, given IJ (global index), the global Index mapped to -// a local index on a Patch(0:nx-1,0:ny-1) in Device Memory as: -// -// I - (ips-ims) + nx * ( J - (jps-jms) ) -// -// where ips is the global index of the start of the patch (the -1 is -// for translating from WRF fortran indices). -// -// The global index I is mapped to a local index on a GPU Block's -// shared memory (0:bx-1, 0:by-1) as: -// -// I - (ips-ims) - bi * bx + by * ( J - (jps-jms) - bj * by ) -// -// Where bi is the index into the GPU Block, and bx is the -// GPU Block Width. - -// global to patch index converter -#define GtoP(i,p,P) ((i)-(p)+(P)) -#define GtoB(i,n,N,p,P) ((i)-(p)+(P)-(n)*(N)) - -// thread index to local memory index = i + bi * bx + ips - ims -#define TtoP(i,a,b,c,d) ((i)+(a)*(b)+(c)-(d)) - -//#define MAX(x,y) ((x)>(y)?(x):(y)) -//#define MIN(x,y) ((x)<(y)?(x):(y)) -#define MAX(x,y) max(x,y) -#define MIN(x,y) min(x,y) - -// basic indexing macros -//#define I2(i,j,m) ((i)+((j)*(m))) -//#define I3(i,j,m,k,n) (I2(i,j,m)+((k)*(m)*(n))) -#define I2(i,j,m) ((i)+(__mul24((j),(m)))) -#define I3(i,j,m,k,n) (I2(i,j,m)+(__mul24((k),__mul24((m),(n))))) - -#if (FLOAT_4 != 4) - -// index into a patch stored on device memory - 1 -# define P2(i,j) I2(TtoP(i,bi,bx,ips,ims),TtoP(j,bj,by,jps,jms),ime-ims+1) -# define P3(i,k,j) I3(TtoP(i,bi,bx,ips,ims),k,ime-ims+1,TtoP(j,bj,by,jps,jms),kme-kms+1) -// index into a block stored on shared memory -# define S2(i,j) I2(i,j,bx) -//# define S3(i,k,j) I3(i,k,bx,j,kme-kms+1) -//# define S3(i,k,j) I3(i,j,bx,k,by) -# define S3(i,k,j) I3(k,i,kx,j,bx) - // Local arrays in device mem -# define LOCDM(a,s) float * a ; cudaMalloc( (void**) & a , (s)*sizeof(float)) ; - // Local scratch arrays in shared memory -# define LOCSM(a,s) __shared__ float * a ; a = &(sm[isize]) ; isize += (s) ; -#define ig (TtoP(ti,bi,bx,ips,ims)) -#define jg (TtoP(tj,bj,by,jps,jms)) - -#else - -// index into a patch stored on device memory - 4 -# define P2(i,j) (I2(TtoP((i)*4,bi,(bx)*4,ips,ims),TtoP((j),bj,by,jps,jms),ime-ims+1)/4) -# define P3(i,k,j) (I3(TtoP((i)*4,bi,(bx)*4,ips,ims),k,ime-ims+1,TtoP((j),bj,by,jps,jms),kme-kms+1)/4) -// index into a block stored on shared memory -# define S2(i,j) (I2(i*4,j,(bx)*4)/4) -# define S3(i,k,j) (I3(i*4,k,(bx)*4,j,kme-kms+1)/4) - // Local arrays in device mem -# define LOCDM(a,s) float4 * a ; cudaMalloc( (void**) & a , (s)*sizeof(float4)) ; - // Local scratch arrays in shared memory -# define LOCSM(a,s) __shared__ Float4 * a ; a = &(sm[isize]) ; isize += (s)*4 ; -#define ig (TtoP((ti)*4,bi,(bx)*4,ips,ims)) -#define jg (TtoP(tj,bj,by,jps,jms)) - -#endif - - -#endif - - - |