diff options
Diffstat (limited to 'benchmarks/CUDA/WP/wsm5.cu')
| -rw-r--r-- | benchmarks/CUDA/WP/wsm5.cu | 500 |
1 files changed, 0 insertions, 500 deletions
diff --git a/benchmarks/CUDA/WP/wsm5.cu b/benchmarks/CUDA/WP/wsm5.cu deleted file mode 100644 index 894e937..0000000 --- a/benchmarks/CUDA/WP/wsm5.cu +++ /dev/null @@ -1,500 +0,0 @@ -#include <stdio.h> -#include <stdlib.h> -#include <math.h> -#include "cublas.h" - -#define IDEBUG 12 -#define JDEBUG 0 - -#ifndef CRAY -# ifdef NOUNDERSCORE -# define WSM5_HOST wsm5_host -# define WSM5_HOST_2 wsm5_host_2 -# define WSM5_GPU_INIT wsm5_gpu_init -# define GET_WSM5_GPU_LEVELS get_wsm5_gpu_levels -# else -# ifdef F2CSTYLE -# define WSM5_HOST wsm5_host__ -# define WSM5_HOST_2 wsm5_host_2__ -# define WSM5_GPU_INIT wsm5_gpu_init__ -# define GET_WSM5_GPU_LEVELS get_wsm5_gpu_levels__ -# else -# define WSM5_HOST wsm5_host_ -# define WSM5_HOST_2 wsm5_host_2_ -# define WSM5_GPU_INIT wsm5_gpu_init_ -# define GET_WSM5_GPU_LEVELS get_wsm5_gpu_levels_ -# endif -# endif -#endif - -#define I2(i,j,m) ((i)+((j)*(m))) -#define I3(i,j,m,k,n) (I2(i,j,m)+((k)*(m)*(n))) - -#if 1 -# define TODEV(A,s) float *A##_d;cudaMalloc((void**)&A##_d,((s))*sizeof(float));cudaMemcpy(A##_d,A,(s)*sizeof(float),cudaMemcpyHostToDevice); -# define FROMDEV(A,s) cudaMemcpy(A,A##_d,(s)*sizeof(float),cudaMemcpyDeviceToHost); -# define CLNUP(A) cudaFree(A##_d) -#else -# define TODEV(A,s) s1=rsl_internal_microclock_() ; float *A##_d;cudaMalloc((void**)&A##_d,((s))*sizeof(float));cudaMemcpy(A##_d,A,(s)*sizeof(float),cudaMemcpyHostToDevice); e1=rsl_internal_microclock_() ; fprintf(stderr,"TODEV %d\n",e1-s1) -# define FROMDEV(A,s) s1=rsl_internal_microclock_() ; cudaMemcpy(A,A##_d,(s)*sizeof(float),cudaMemcpyDeviceToHost); e1=rsl_internal_microclock_() ; fprintf(stderr,"FROMDEV %d\n",e1-s1) -# define CLNUP(A) s1=rsl_internal_microclock_() ; cudaFree(A##_d) ; e1=rsl_internal_microclock_() ; fprintf(stderr,"Free %d\n",e1-s1) -#endif - -#if FLOAT_4==4 -#define TODEV2(A)s1=rsl_internal_microclock_();\ -float*A##_d;\ -cudaMalloc((void**)&A##_d,(dipe*djpe*sizeof(float)));\ -for(j=*jps-1;j<=*jpe-1;j++){\ - for(i=*ips-1;i<=*ipe-1;i++){\ - bigbuf[I2(i-*ips+1,j-*jps+1,dipe)]=\ - A[I2(i-*ims+1,j-*jms+1,(*ime-*ims+1))];\ -}}\ -cudaMemcpy(A##_d,bigbuf,(dipe*djpe)*sizeof(float),cudaMemcpyHostToDevice);\ -e1=rsl_internal_microclock_();fprintf(stderr,"TODEV2 %d\n",e1-s1); - -#define TODEV3(A)s1=rsl_internal_microclock_();\ -float*A##_d;\ -cudaMalloc((void**)&A##_d,(dipe*djpe*dkpe*sizeof(float)));\ -for(j=*jps-1;j<=*jpe-1;j++){\ - for(k=*kps-1;k<=*kpe-1;k++){\ - for(i=*ips-1;i<=*ipe-1;i++){\ - bigbuf[I3(i-*ips+1,k-*kps+1,dipe,j-*jps+1,dkpe)]=\ - A[I3(i-*ims+1,k-*kms+1,*ime-*ims+1,j-*jms+1,*kme-*kms+1)];\ -}}}\ -cudaMemcpy(A##_d,bigbuf,(dipe*djpe*dkpe)*sizeof(float),cudaMemcpyHostToDevice);\ -e1=rsl_internal_microclock_();fprintf(stderr,"TODEV3 %d\n",e1-s1); - -// for debugging only -#define TODEV3a(A)s1=rsl_internal_microclock_();\ -float*A##_d;\ -cudaMalloc((void**)&A##_d,(dipe*djpe*dkpe*sizeof(float)));\ -for(j=*jps-1;j<=*jpe-1;j++){\ - for(k=*kps-1;k<=*kpe-1;k++){\ - for(i=*ips-1;i<=*ipe-1;i++){\ - bigbuf[I3(i-*ips+1,k-*kps+1,dipe,j-*jps+1,dkpe)]=\ - A[I3(i-*ims+1,k-*kms+1,*ime-*ims+1,j-*jms+1,*kme-*kms+1)];\ -if (i==*ips-1){\ - fprintf(stderr,"There %d %d %d (%d)| %d %d %d (%d)| %f\n",\ - i-*ips+1,k-*kps+1,j-*jps+1,I3(i-*ips+1,k-*kps+1,dipe,j-*jps+1,dkpe),\ - i-*ims+1,k-*kms+1,j-*jms+1,I3(i-*ims+1,k-*kms+1,*ime-*ims+1,j-*jms+1,*kme-*kms+1),\ - bigbuf[I3(i-*ips+1,k-*kps+1,dipe,j-*jps+1,dkpe)]);\ - A[I3(i-*ims+1,k-*kms+1,*ime-*ims+1,j-*jms+1,*kme-*kms+1)]=199.;\ -}\ -}}}\ -cudaMemcpy(A##_d,bigbuf,(dipe*djpe*dkpe)*sizeof(float),cudaMemcpyHostToDevice);\ -e1=rsl_internal_microclock_();fprintf(stderr,"TODEV3 %d\n",e1-s1); - - -#define FROMDEV2(A) s1=rsl_internal_microclock_();\ -cudaMemcpy(bigbuf,A##_d,dipe*djpe*sizeof(float),cudaMemcpyDeviceToHost);\ -for(j=*jps-1;j<=*jpe-1;j++){\ - for(i=*ips-1;i<=*ipe-1;i++){\ - A[I2(i-*ims+1,j-*jms+1,(*ime-*ims+1))]=\ - bigbuf[I2(i-*ips+1,j-*jps+1,dipe)];\ -}}\ -e1=rsl_internal_microclock_() ; fprintf(stderr,"FROMDEV2 %d\n",e1-s1); - -#define FROMDEV3(A) s1=rsl_internal_microclock_();\ -cudaMemcpy(bigbuf,A##_d,dipe*djpe*dkpe*sizeof(float),cudaMemcpyDeviceToHost);\ -for(j=*jps-1;j<=*jpe-1;j++){\ - for(k=*kps-1;k<=*kpe-1;k++){\ - for(i=*ips-1;i<=*ipe-1;i++){\ - A[I3(i-*ims+1,k-*kms+1,*ime-*ims+1,j-*jms+1,*kme-*kms+1)]=\ - bigbuf[I3(i-*ips+1,k-*kps+1,dipe,j-*jps+1,dkpe)];\ -}}}\ -e1=rsl_internal_microclock_();fprintf(stderr,"FROMDEV3 %d\n",e1-s1); -#else -# define TODEV3(A) TODEV(A,d3) -# define TODEV2(A) TODEV(A,d2) -# define FROMDEV3(A) FROMDEV(A,d3) -# define FROMDEV2(A) FROMDEV(A,d2) -#endif - -extern "C" int rsl_internal_microclock_() ; - -extern __global__ void wsm5_gpu ( - float *th, float *pii //_def_ arg ikj:th,pii - ,float *q //_def_ arg ikj:q - ,float *qc,float *qi,float *qr,float *qs //_def_ arg ikj:qc,qi,qr,qs - ,float *den, float *p, float *delz //_def_ arg ikj:den,p,delz -#ifdef DEBUGGAL_ARRAY -,float *debuggal //_def_ arg ikj:debuggal -#endif - ,float *rain,float *rainncv //_def_ arg ij:rain,rainncv - ,float *sr //_def_ arg ij:sr - ,float *snow,float *snowncv //_def_ arg ij:snow,snowncv - ,float delt -,float* retvals - ,int ids, int ide, int jds, int jde, int kds, int kde - ,int ims, int ime, int jms, int jme, int kms, int kme - ,int ips, int ipe, int jps, int jpe, int kps, int kpe - ) ; - -extern "C" { - -int gethostname(char *name, size_t len); -void bzero(void *s, size_t n); -char *strcpy(char *dest, const char *src); - -#define MAXDEVICES 4 -#define MAXNODES 16 -int -WSM5_GPU_INIT ( int * myproc , int * nproc, int * mydevice ) -{ - float x, *x_d ; - int s, e ; - int i, dc, m ; - cudaError_t cerr ; - char hostname[64] ; - struct cudaDeviceProp dp ; -// manage devices if multiheaded - cudaGetDeviceCount( &dc ) ; - if ( dc > MAXDEVICES ) - { fprintf(stderr, "warning: more than %d devices on node (%d)\n", MAXDEVICES, dc ) ; dc = MAXDEVICES ; } - fprintf(stderr,"Number of devices on this node: %d\n", dc) ; - - // i = *myproc % dc ; - - i = *mydevice ; - if ( dc > 0 ) - { - if ( cerr = cudaSetDevice( i ) ) { - fprintf(stderr," non-zero cerr %d\n",cerr) ; - } - } - gethostname( hostname, 64 ) ; - fprintf(stderr,"Setting device %02d for task %03d on host %s\n",i,*myproc,hostname) ; - - if ( cerr = cudaGetDeviceProperties( &dp, i ) ) { - fprintf(stderr,"Device %02d: cerr = %d\n", cerr) ; - } else { - fprintf(stderr,"Device %02d: name %s\n",i,dp.name) ; - fprintf(stderr,"Device %02d: mem %d\n",i,dp.totalGlobalMem) ; - fprintf(stderr,"Device %02d: smem %d\n",i,dp.sharedMemPerBlock) ; - fprintf(stderr,"Device %02d: nreg %d\n",i,dp.regsPerBlock) ; - fprintf(stderr,"Device %02d: warp %d\n",i,dp.warpSize) ; - fprintf(stderr,"Device %02d: pitch %d\n",i,dp.memPitch) ; - fprintf(stderr,"Device %02d: maxthrds %d\n",i,dp.maxThreadsPerBlock) ; - fprintf(stderr,"Device %02d: maxtdim %d %d %d\n",i,dp.maxThreadsDim[0] - ,dp.maxThreadsDim[1] - ,dp.maxThreadsDim[2]) ; - fprintf(stderr,"Device %02d: maxgdim %d %d %d\n",i,dp.maxGridSize[0] - ,dp.maxGridSize[1] - ,dp.maxGridSize[2]) ; - fprintf(stderr,"Device %02d: clock %d\n",i,dp.clockRate) ; - fprintf(stderr,"Device %02d: talign %d\n",i,dp.textureAlignment) ; - } - -// do a dummy init to get things going - s=rsl_internal_microclock_() ; - cudaMalloc((void **)&x_d,sizeof(float)) ; - cudaMemcpy(x_d,&x,sizeof(float),cudaMemcpyHostToDevice) ; - cudaFree(x_d) ; - e=rsl_internal_microclock_() ; - fprintf(stderr,"wsm5_init: %d\n",e-s) ; - return(0) ; -} - -int -WSM5_HOST ( - float *th, float *pii - ,float *q - ,float *qc, float *qi, float *qr, float *qs - ,float *den, float *p, float *delz -#ifdef DEBUGGAL_ARRAY -,float *debuggal -#endif - ,float *delt - ,float *rain,float *rainncv - ,float *sr - ,float *snow,float *snowncv - ,int *ids, int *ide, int *jds, int *jde, int *kds, int *kde - ,int *ims, int *ime, int *jms, int *jme, int *kms, int *kme - ,int *ips, int *ipe, int *jps, int *jpe, int *kps, int *kpe - ) -{ - int i, j, k ; - float *bigbuf ; - int s, e, s1, e1, s2, e2 ; - int d3 = (*ime-*ims+1) * (*jme-*jms+1) * (*kme-*kms+1) ; - int d2 = (*ime-*ims+1) * (*jme-*jms+1) ; - - -//fprintf(stderr,"d3 = %d\n",d3) ; -//fprintf(stderr,"d2 = %d\n",d2) ; - -#if FLOAT_4 == 4 - int dips = 0 ; int dipe = (((*ipe-*ips+1+3)/4)*4) ; // round up four -#else - int dips = 0 ; int dipe = (*ipe-*ips+1) ; -#endif - int djps = 0 ; int djpe = (*jpe-*jps+1) ; - int dkps = 0 ; int dkpe = (*kpe-*kps+1) ; - - bigbuf = (float *)malloc( dipe * djpe * dkpe * sizeof(float) ) ; - -//fprintf(stderr,"ids %d ide %d jds %d jde %d kds %d kde %d\n",*ids,*ide,*jds,*jde,*kds,*kde) ; -//fprintf(stderr,"ims %d ime %d jms %d jme %d kms %d kme %d\n",*ims,*ime,*jms,*jme,*kms,*kme) ; -//fprintf(stderr,"ips %d ipe %d jps %d jpe %d kps %d kpe %d\n",*ips,*ipe,*jps,*jpe,*kps,*kpe) ; -//fprintf(stderr,"dipe %d djpe %d dkpe %d\n",dipe,djpe,dkpe) ; - - s = rsl_internal_microclock_() ; - TODEV3(th) ; - TODEV3(pii) ; - TODEV3(q) ; - TODEV3(qc) ; - TODEV3(qi) ; - TODEV3(qr) ; - TODEV3(qs) ; - TODEV3(den) ; - TODEV3(p) ; - TODEV3(delz) ; -#ifdef DEBUGGAL_ARRAY -//TODEV3(debuggal) ; -#endif - TODEV2(rain) ; - TODEV2(rainncv) ; - TODEV2(sr) ; - TODEV2(snow) ; - TODEV2(snowncv) ; -float retvals[100] ; -{ int k ; -for (k=0 ;k<*kme-*kms+1;k++) {retvals[k] = 0.; } -} -TODEV(retvals,(*kme-*kms+1)) ; - - int remx, remy ; // remainder? - - remx = (*ipe-*ips+1) % XXX != 0 ? 1 : 0 ; - remy = (*jpe-*jps+1) % YYY != 0 ? 1 : 0 ; - - dim3 dimBlock( XXX , YYY ) ; -// fprintf(stderr,"ipe ips remx jpe jps remy %d %d %d %d %d %d\n",*ipe,*ips,remx,*jpe,*jps,remy) ; - dim3 dimGrid ( (*ipe-*ips+1) / XXX + remx , (*jpe-*jps+1) / YYY + remy ) ; - - fprintf(stderr,"Call to wsm5_gpu: block dims %d %d\n",dimBlock.x,dimBlock.y) ; - fprintf(stderr,"Call to wsm5_gpu: grid dims %d %d\n",dimGrid.x,dimGrid.y) ; - -#if 1 -//fprintf(stderr,"calling wsm5_gpu \n") ; -//fprintf(stderr,"d %d %d %d %d %d %d\n",dips+1 , (*ipe-*ips+1) , djps+1 , (*jpe-*jps+1) , dkps+1 , (*kpe-*kps+1)) ; -//fprintf(stderr,"m %d %d %d %d %d %d\n",dips+1 , dipe , djps+1 , djpe , dkps+1 , dkpe ) ; -//fprintf(stderr,"p %d %d %d %d %d %d\n",dips+1 , dipe , djps+1 , djpe , dkps+1 , dkpe ) ; - - s2 = rsl_internal_microclock_() ; - wsm5_gpu <<< dimGrid, dimBlock >>> ( - th_d, pii_d, q_d, qc_d, qi_d, qr_d, qs_d, den_d, p_d, delz_d -#ifdef DEBUGGAL_ARRAY -,debuggal_d -#endif - ,rain_d,rainncv_d - ,sr_d - ,snow_d,snowncv_d - ,*delt -,retvals_d - ,dips+1 , (*ipe-*ips+1) , djps+1 , (*jpe-*jps+1) , dkps+1 , (*kpe-*kps+1) - ,dips+1 , dipe , djps+1 , djpe , dkps+1 , dkpe - ,dips+1 , dipe , djps+1 , djpe , dkps+1 , dkpe - ) ; - cudaThreadSynchronize() ; - e2 = rsl_internal_microclock_() ; - fprintf(stderr,"Call to wsm5_gpu (not including data xfer): %d microseconds\n",e2-s2) ; -#endif - - FROMDEV3(th) ; - FROMDEV3(pii) ; - FROMDEV3(q) ; - FROMDEV3(qc) ; - FROMDEV3(qi) ; - FROMDEV3(qr) ; - FROMDEV3(qs) ; -#ifdef DEBUGGAL_ARRAY -FROMDEV3(debuggal) ; -#endif - FROMDEV2(rain) ; - FROMDEV2(rainncv) ; - FROMDEV2(sr) ; - FROMDEV2(snow) ; - FROMDEV2(snowncv) ; - e = rsl_internal_microclock_() ; -//fprintf(stderr,"retrieving retvals %d\n",*kme-*kms+1) ; -FROMDEV(retvals,(*kme-*kms+1)) ; - fprintf(stderr,"Call to wsm5_gpu (including data xfer): %d microseconds\n",e-s) ; - -{ int k ; -//for (k=0 ;k<*kme-*kms+1;k++) {fprintf(stderr,"retvals %d %f\n",k,retvals[k]) ;} -//for (k=0 ;k<5;k++) {fprintf(stderr,"retvals %d %f\n",k,retvals[k]) ;} -} - - CLNUP(th) ; - CLNUP(pii) ; - CLNUP(q) ; - CLNUP(qc) ; - CLNUP(qi) ; - CLNUP(qr) ; - CLNUP(qs) ; - CLNUP(den) ; - CLNUP(p) ; - CLNUP(delz) ; -#ifdef DEBUGGAL_ARRAY -CLNUP(debuggal) ; -#endif - CLNUP(rain) ; - CLNUP(rainncv) ; - CLNUP(sr) ; - CLNUP(snow) ; - CLNUP(snowncv) ; -CLNUP(retvals) ; - - return(0) ; -} - -#if 0 -static int first_wsm5_host_2=1 ; -// 3d -static float * th_h ; -static float * pii_h ; -static float * q_h ; -static float * qc_h ; -static float * qi_h ; -static float * qr_h ; -static float * qs_h ; -static float * den_h ; -static float * p_h ; -static float * delz_h ; -// 2d -static float * rain_h ; -static float * rainncv_h ; -static float * sr_h ; -static float * snow_h ; -static float * snowncv_h ; - -// idea here is to copy the data into pinned (paged-locked) mem for faster xfer -int -WSM5_HOST_2 ( - float *th, float *pii - ,float *q - ,float *qc, float *qi, float *qr, float *qs - ,float *den, float *p, float *delz - ,float *delt - ,float *rain,float *rainncv - ,float *sr - ,float *snow,float *snowncv - ,int *ids, int *ide, int *jds, int *jde, int *kds, int *kde - ,int *ims, int *ime, int *jms, int *jme, int *kms, int *kme - ,int *ips, int *ipe, int *jps, int *jpe, int *kps, int *kpe - ) -{ - int i,j,k ; - float *ptr ; - int d3 = (*ipe-*ips+1) * (*jpe-*jps+1) * (*kpe-*kps+1) ; - int d2 = (*ipe-*ips+1) * (*jpe-*jps+1) ; - - if ( first_wsm5_host_2 == 1 ) { - cudaMallocHost( (void **)&th_h , d3*sizeof(float) ) ; //3d - cudaMallocHost( (void **)&pii_h , d3*sizeof(float) ) ; - cudaMallocHost( (void **)&q_h , d3*sizeof(float) ) ; - cudaMallocHost( (void **)&qc_h , d3*sizeof(float) ) ; - cudaMallocHost( (void **)&qi_h , d3*sizeof(float) ) ; - cudaMallocHost( (void **)&qr_h , d3*sizeof(float) ) ; - cudaMallocHost( (void **)&qs_h , d3*sizeof(float) ) ; - cudaMallocHost( (void **)&den_h , d3*sizeof(float) ) ; - cudaMallocHost( (void **)&p_h , d3*sizeof(float) ) ; - cudaMallocHost( (void **)&delz_h , d3*sizeof(float) ) ; - cudaMallocHost( (void **)&rain_h , d2*sizeof(float) ) ; //2d - cudaMallocHost( (void **)&rainncv_h , d2*sizeof(float) ) ; - cudaMallocHost( (void **)&sr_h , d2*sizeof(float) ) ; - cudaMallocHost( (void **)&snow_h , d2*sizeof(float) ) ; - cudaMallocHost( (void **)&snowncv_h , d2*sizeof(float) ) ; - first_wsm5_host_2 = 0 ; - } - -#define PIN3(A) ptr=A##_h;for(j=*jps;j<=*jpe;j++){for(k=*kps;k<=*kpe;k++){for(i=*ips;i<=*ipe;i++){*ptr++=A [I3(i-*ims,k-*kms,*ime-*ims+1,j-*jms,*kme-*kms+1)];}}}; -#define PIN2(A) ptr=A##_h;for(j=*jps;j<=*jpe;j++);for(i=*ips;i<=*ipe;i++){*ptr++=A [I2(i-*ims,j-*jms,*ime-*ims+1)];}; -#define UNPIN3(A) ptr=A##_h;for(j=*jps;j<=*jpe;j++);for(k=*kps;k<=*kpe;k++);for(i=*ips;i<=*ipe;i++){A [I3(i-*ims,k-*kms,*ime-*ims+1,j-*jms,*kme-*kms+1)]=*ptr++;}; -#define UNPIN2(A) ptr=A##_h;for(j=*jps;j<=*jpe;j++);for(i=*ips;i<=*ipe;i++){A [I2(i-*ims,j-*jms,*ime-*ims+1)]=*ptr++;}; - - PIN3(th) ; - PIN3(th) ; - PIN3(pii) ; - PIN3(q) ; - PIN3(qc) ; - PIN3(qi) ; - PIN3(qr) ; - PIN3(qs) ; - PIN3(den) ; - PIN3(p) ; - PIN3(delz) ; - PIN2(rain) ; - PIN2(rainncv) ; - PIN2(sr) ; - PIN2(snow) ; - PIN2(snowncv) ; - - WSM5_HOST ( - th_h, pii_h - ,q_h - ,qc_h, qi_h, qr_h, qs_h - ,den_h, p_h, delz_h - ,delt - ,rain_h,rainncv_h - ,sr_h - ,snow_h,snowncv_h - ,ids, ide, jds, jde, kds, kde - ,ips, ipe, jps, jpe, kps, kpe - ,ips, ipe, jps, jpe, kps, kpe - ) ; - - - UNPIN3(th) ; - UNPIN3(th) ; - UNPIN3(pii) ; - UNPIN3(q) ; - UNPIN3(qc) ; - UNPIN3(qi) ; - UNPIN3(qr) ; - UNPIN3(qs) ; - UNPIN3(den) ; - UNPIN3(p) ; - UNPIN3(delz) ; - UNPIN2(rain) ; - UNPIN2(rainncv) ; - UNPIN2(sr) ; - UNPIN2(snow) ; - UNPIN2(snowncv) ; - -} -#endif - -int -GET_WSM5_GPU_LEVELS ( int * retval ) -{ - *retval = MKX ; /* MKX is hard coded value set in the makefile */ -} -} - -#if 0 -main( int argc, char **argv ) -{ - float *th ; float *pii ; float *q ; - float *qc; float *qi; float *qr; float *qs ; - float *den; float *p; float *delz ; - float *delt ; - float *rain;float *rainncv ; - float *sr ; - float *snow;float *snowncv ; - int *ids; int *ide; int *jds; int *jde; int *kds; int *kde ; - int *ims; int *ime; int *jms; int *jme; int *kms; int *kme ; - int *ips; int *ipe; int *jps; int *jpe; int *kps; int *kpe ; - WSM5_HOST ( - th, pii, q, qc, qi, qr, qs, den, p, delz - ,rain,rainncv - ,sr - ,snow,snowncv - ,delt - ,ids, ide, jds, jde, kds, kde - ,ims, ime, jms, jme, kms, kme - ,ips, ipe, jps, jpe, kps, kpe - ) ; -} -#endif |