Main Page   Namespace List   Class Hierarchy   Alphabetical List   Compound List   File List   Namespace Members   Compound Members   File Members   Related Pages  

CUDABench.cu

Go to the documentation of this file.
00001 /***************************************************************************
00002  *cr
00003  *cr            (C) Copyright 1995-2019 The Board of Trustees of the
00004  *cr                        University of Illinois
00005  *cr                         All Rights Reserved
00006  *cr
00007  ***************************************************************************/
00008 /***************************************************************************
00009  * RCS INFORMATION:
00010  *
00011  *      $RCSfile: CUDABench.cu,v $
00012  *      $Author: johns $        $Locker:  $             $State: Exp $
00013  *      $Revision: 1.42 $      $Date: 2022/02/09 04:03:19 $
00014  *
00015  ***************************************************************************/
00021 #include <stdio.h>
00022 #include <stdlib.h>
00023 #include <string.h>
00024 #include <cuda.h>
00025 
00026 #include "Inform.h"
00027 #include "WKFThreads.h"
00028 #include "WKFUtils.h"
00029 #include "CUDAKernels.h" 
00030 #include "Measure.h"
00031 
00032 
00033 //
00034 // Restrict macro to make it easy to do perf tuning tests
00035 //
00036 #if 1
00037 #define RESTRICT __restrict__
00038 #else
00039 #define RESTRICT
00040 #endif
00041 
00042 
00043 //#define VMDUSECUDAGDS 1
00044 #if defined(VMDUSECUDAGDS)
00045 #include </usr/local/gds-beta-0.7.1/lib/cufile.h>        // GPU-Direct Storage
00046 
00047 // direct calls to JS plugin for devel/testing until the plugin manager
00048 // and headers incorporate the new out-of-core GPU-direct I/O hooks.
00049 #define VMDJSPLUGININCLUDESRC 1
00050 #include "/home/johns/plugins/molfile_plugin/src/jsplugin.c"
00051 #endif
00052 
00053 
00054 #define CUERR { cudaError_t err; \
00055   if ((err = cudaGetLastError()) != cudaSuccess) { \
00056   printf("CUDA error: %s, %s line %d\n", cudaGetErrorString(err), __FILE__, __LINE__); \
00057   return -1; }}
00058 
00059 
00060 //
00061 // Benchmark peak Multiply-Add instruction performance, in GFLOPS
00062 //
00063 
00064 // FMADD16 macro contains a sequence of operations that the compiler
00065 // won't optimize out, and will translate into a densely packed block
00066 // of multiply-add instructions with no intervening register copies/moves
00067 // or other instructions. 
00068 #define FMADD16 \
00069     tmp0  = tmp0*tmp4+tmp7;     \
00070     tmp1  = tmp1*tmp5+tmp0;     \
00071     tmp2  = tmp2*tmp6+tmp1;     \
00072     tmp3  = tmp3*tmp7+tmp2;     \
00073     tmp4  = tmp4*tmp0+tmp3;     \
00074     tmp5  = tmp5*tmp1+tmp4;     \
00075     tmp6  = tmp6*tmp2+tmp5;     \
00076     tmp7  = tmp7*tmp3+tmp6;     \
00077     tmp8  = tmp8*tmp12+tmp15;   \
00078     tmp9  = tmp9*tmp13+tmp8;    \
00079     tmp10 = tmp10*tmp14+tmp9;   \
00080     tmp11 = tmp11*tmp15+tmp10;  \
00081     tmp12 = tmp12*tmp8+tmp11;   \
00082     tmp13 = tmp13*tmp9+tmp12;   \
00083     tmp14 = tmp14*tmp10+tmp13;  \
00084     tmp15 = tmp15*tmp11+tmp14;
00085 
00086 // CUDA grid, thread block, loop, and MADD operation counts
00087 #define GRIDSIZEX       6144  // number of 1-D thread blocks
00088 #define BLOCKSIZEX      64    // number of threads per 1-D block
00089 #define GLOOPS          2000  // iteration count (all threads)
00090 #define FMADD16COUNT    32    // 32 reps
00091 #define FLOPSPERFMADD16 32    // 16 MULs and 16 ADDs
00092 
00093 // FLOP counting
00094 #define FLOPSPERLOOP (FMADD16COUNT * FLOPSPERFMADD16)
00095 
00096 __global__ static void madd_kernel(float *doutput) {
00097   int tid = blockIdx.x * blockDim.x + threadIdx.x;
00098   float tmp0,tmp1,tmp2,tmp3,tmp4,tmp5,tmp6,tmp7;
00099   float tmp8,tmp9,tmp10,tmp11,tmp12,tmp13,tmp14,tmp15;
00100   tmp0=tmp1=tmp2=tmp3=tmp4=tmp5=tmp6=tmp7=0.0f;
00101   tmp8=tmp9=tmp10=tmp11=tmp12=tmp13=tmp14=tmp15 = 0.0f;
00102 
00103   tmp15=tmp7 = blockIdx.x * 0.001f; // prevent compiler from optimizing out
00104   tmp1 = blockIdx.y * 0.001f;       // the body of the loop...
00105 
00106   int loop;
00107   for(loop=0; loop<GLOOPS; loop++){
00108     FMADD16
00109     FMADD16
00110     FMADD16
00111     FMADD16
00112     FMADD16
00113     FMADD16
00114     FMADD16
00115     FMADD16
00116     FMADD16
00117     FMADD16
00118     FMADD16
00119     FMADD16
00120     FMADD16
00121     FMADD16
00122     FMADD16
00123     FMADD16
00124     FMADD16
00125     FMADD16
00126     FMADD16
00127     FMADD16
00128     FMADD16
00129     FMADD16
00130     FMADD16
00131     FMADD16
00132     FMADD16
00133     FMADD16
00134     FMADD16
00135     FMADD16
00136     FMADD16
00137     FMADD16
00138     FMADD16
00139     FMADD16
00140   }
00141 
00142   doutput[tid] = tmp0+tmp1+tmp2+tmp3+tmp4+tmp5+tmp6+tmp7
00143                  +tmp8+tmp9+tmp10+tmp11+tmp12+tmp13+tmp14+tmp15;
00144 }
00145 
00146 
00147 static int cudamaddgflops(int cudadev, double *gflops, int testloops) {
00148   float *doutput = NULL;
00149   dim3 Bsz, Gsz;
00150   wkf_timerhandle timer;
00151   int i;
00152 
00153   cudaError_t rc;
00154   rc = cudaSetDevice(cudadev);
00155   if (rc != cudaSuccess) {
00156 #if CUDART_VERSION >= 2010
00157     rc = cudaGetLastError(); // query last error and reset error state
00158     if (rc != cudaErrorSetOnActiveProcess)
00159       return -1; // abort and return an error
00160 #else
00161     cudaGetLastError(); // just ignore and reset error state, since older CUDA
00162                         // revs don't have a cudaErrorSetOnActiveProcess enum
00163 #endif
00164   }
00165 
00166 
00167   // setup CUDA grid and block sizes
00168   Bsz.x = BLOCKSIZEX;
00169   Bsz.y = 1;
00170   Bsz.z = 1;
00171   Gsz.x = GRIDSIZEX;
00172   Gsz.y = 1;
00173   Gsz.z = 1;
00174 
00175   // allocate output array
00176   cudaMalloc((void**)&doutput, BLOCKSIZEX * GRIDSIZEX * sizeof(float));
00177   CUERR // check and clear any existing errors
00178 
00179   // warmup run
00180   madd_kernel<<<Gsz, Bsz>>>(doutput);
00181   cudaDeviceSynchronize(); // wait for kernel to finish
00182 
00183   // benchmark run
00184   timer=wkf_timer_create();
00185   wkf_timer_start(timer);
00186   for (i=0; i<testloops; i++) { 
00187     madd_kernel<<<Gsz, Bsz>>>(doutput);
00188   }
00189   cudaDeviceSynchronize(); // wait for kernel to finish
00190   CUERR // check and clear any existing errors
00191   wkf_timer_stop(timer);
00192 
00193   double runtime = wkf_timer_time(timer);
00194   double gflop = ((double) GLOOPS) * ((double) FLOPSPERLOOP) *
00195                   ((double) BLOCKSIZEX) * ((double) GRIDSIZEX) * (1.0e-9) * testloops;
00196   
00197   *gflops = gflop / runtime;
00198 
00199   cudaFree(doutput);
00200   CUERR // check and clear any existing errors
00201 
00202   wkf_timer_destroy(timer);
00203 
00204   return 0;
00205 }
00206 
00207 typedef struct {
00208   int deviceid;
00209   int testloops;
00210   double gflops;
00211 } maddthrparms;
00212 
00213 static void * cudamaddthread(void *voidparms) {
00214   maddthrparms *parms = (maddthrparms *) voidparms;
00215   cudamaddgflops(parms->deviceid, &parms->gflops, parms->testloops);
00216   return NULL;
00217 }
00218 
00219 int vmd_cuda_madd_gflops(int numdevs, int *devlist, double *gflops,
00220                          int testloops) {
00221   maddthrparms *parms;
00222   wkf_thread_t * threads;
00223   int i;
00224 
00225   /* allocate array of threads */
00226   threads = (wkf_thread_t *) calloc(numdevs * sizeof(wkf_thread_t), 1);
00227 
00228   /* allocate and initialize array of thread parameters */
00229   parms = (maddthrparms *) malloc(numdevs * sizeof(maddthrparms));
00230   for (i=0; i<numdevs; i++) {
00231     if (devlist != NULL)
00232       parms[i].deviceid = devlist[i];
00233     else
00234       parms[i].deviceid = i;
00235 
00236     parms[i].testloops = testloops;
00237     parms[i].gflops = 0.0;
00238   }
00239 
00240 #if defined(VMDTHREADS)
00241   /* spawn child threads to do the work */
00242   /* thread 0 must also be processed this way otherwise    */
00243   /* we'll permanently bind the main thread to some device */
00244   for (i=0; i<numdevs; i++) {
00245     wkf_thread_create(&threads[i], cudamaddthread, &parms[i]);
00246   }
00247 
00248   /* join the threads after work is done */
00249   for (i=0; i<numdevs; i++) {
00250     wkf_thread_join(threads[i], NULL);
00251   }
00252 #else
00253   /* single thread does all of the work */
00254   cudamaddthread((void *) &parms[0]);
00255 #endif
00256 
00257   for (i=0; i<numdevs; i++) {
00258     gflops[i] = parms[i].gflops; 
00259   }
00260 
00261   /* free thread parms */
00262   free(parms);
00263   free(threads);
00264 
00265   return 0;
00266 }
00267 
00268 
00269 
00270 
00271 
00272 
00273 //
00274 // Host-GPU memcpy I/O bandwidth benchmark
00275 //
00276 
00277 #define BWITER 500
00278 #define LATENCYITER 50000
00279 
00280 static int cudabusbw(int cudadev, 
00281                      double *hdmbsec, double *hdlatusec, 
00282                      double *phdmbsec, double *phdlatusec, 
00283                      double *dhmbsec, double *dhlatusec,
00284                      double *pdhmbsec, double *pdhlatusec) {
00285   float *hdata = NULL;   // non-pinned DMA buffer
00286   float *phdata = NULL;  // pinned DMA buffer
00287   float *ddata = NULL;
00288   int i;
00289   double runtime;
00290   wkf_timerhandle timer;
00291   int memsz = 1024 * 1024 * sizeof(float);
00292 
00293   *hdmbsec = 0.0;
00294   *hdlatusec = 0.0;
00295   *dhmbsec = 0.0;
00296   *dhlatusec = 0.0;
00297   *phdmbsec = 0.0;
00298   *phdlatusec = 0.0;
00299   *pdhmbsec = 0.0;
00300   *pdhlatusec = 0.0;
00301 
00302   // attach to the selected device
00303   cudaError_t rc;
00304   rc = cudaSetDevice(cudadev);
00305   if (rc != cudaSuccess) {
00306 #if CUDART_VERSION >= 2010
00307     rc = cudaGetLastError(); // query last error and reset error state
00308     if (rc != cudaErrorSetOnActiveProcess)
00309       return -1; // abort and return an error
00310 #else
00311     cudaGetLastError(); // just ignore and reset error state, since older CUDA
00312                         // revs don't have a cudaErrorSetOnActiveProcess enum
00313 #endif
00314   }
00315 
00316   // allocate non-pinned output array
00317   hdata = (float *) malloc(memsz); 
00318 
00319   // allocate pinned output array
00320   cudaMallocHost((void**) &phdata, memsz);
00321   CUERR // check and clear any existing errors
00322 
00323   // allocate device memory
00324   cudaMalloc((void**) &ddata, memsz);
00325   CUERR // check and clear any existing errors
00326 
00327   // create timer
00328   timer=wkf_timer_create();
00329 
00330   //
00331   // Host to device timings
00332   //
00333 
00334   // non-pinned bandwidth
00335   wkf_timer_start(timer);
00336   for (i=0; i<BWITER; i++) {
00337     cudaMemcpy(ddata, hdata, memsz,  cudaMemcpyHostToDevice);
00338   }
00339   wkf_timer_stop(timer);
00340   CUERR // check and clear any existing errors
00341   runtime = wkf_timer_time(timer);
00342   *hdmbsec = ((double) BWITER) * ((double) memsz) / runtime / (1024.0 * 1024.0);
00343 
00344   // non-pinned latency
00345   wkf_timer_start(timer);
00346   for (i=0; i<LATENCYITER; i++) {
00347     cudaMemcpy(ddata, hdata, 1,  cudaMemcpyHostToDevice);
00348   }
00349   wkf_timer_stop(timer);
00350   CUERR // check and clear any existing errors
00351   runtime = wkf_timer_time(timer);
00352   *hdlatusec = runtime * 1.0e6 / ((double) LATENCYITER);
00353 
00354 
00355   // pinned bandwidth
00356   wkf_timer_start(timer);
00357   for (i=0; i<BWITER; i++) {
00358     cudaMemcpy(ddata, phdata, memsz,  cudaMemcpyHostToDevice);
00359   }
00360   wkf_timer_stop(timer);
00361   CUERR // check and clear any existing errors
00362   runtime = wkf_timer_time(timer);
00363   *phdmbsec = ((double) BWITER) * ((double) memsz) / runtime / (1024.0 * 1024.0);
00364 
00365   // pinned latency
00366   wkf_timer_start(timer);
00367   for (i=0; i<LATENCYITER; i++) {
00368     cudaMemcpy(ddata, phdata, 1,  cudaMemcpyHostToDevice);
00369   }
00370   wkf_timer_stop(timer);
00371   CUERR // check and clear any existing errors
00372   runtime = wkf_timer_time(timer);
00373   *phdlatusec = runtime * 1.0e6 / ((double) LATENCYITER);
00374 
00375  
00376   //
00377   // Device to host timings
00378   //
00379 
00380   // non-pinned bandwidth
00381   wkf_timer_start(timer);
00382   for (i=0; i<BWITER; i++) {
00383     cudaMemcpy(hdata, ddata, memsz,  cudaMemcpyDeviceToHost);
00384   }
00385   wkf_timer_stop(timer);
00386   CUERR // check and clear any existing errors
00387   runtime = wkf_timer_time(timer);
00388   *dhmbsec = ((double) BWITER) * ((double) memsz) / runtime / (1024.0 * 1024.0);
00389 
00390   // non-pinned latency
00391   wkf_timer_start(timer);
00392   for (i=0; i<LATENCYITER; i++) {
00393     cudaMemcpy(hdata, ddata, 1,  cudaMemcpyDeviceToHost);
00394   }
00395   wkf_timer_stop(timer);
00396   CUERR // check and clear any existing errors
00397   runtime = wkf_timer_time(timer);
00398   *dhlatusec = runtime * 1.0e6 / ((double) LATENCYITER);
00399 
00400 
00401   // pinned bandwidth
00402   wkf_timer_start(timer);
00403   for (i=0; i<BWITER; i++) {
00404     cudaMemcpy(phdata, ddata, memsz,  cudaMemcpyDeviceToHost);
00405   }
00406   wkf_timer_stop(timer);
00407   CUERR // check and clear any existing errors
00408   runtime = wkf_timer_time(timer);
00409   *pdhmbsec = ((double) BWITER) * ((double) memsz) / runtime / (1024.0 * 1024.0);
00410 
00411   // pinned latency
00412   wkf_timer_start(timer);
00413   for (i=0; i<LATENCYITER; i++) {
00414     cudaMemcpy(phdata, ddata, 1,  cudaMemcpyDeviceToHost);
00415   }
00416   wkf_timer_stop(timer);
00417   CUERR // check and clear any existing errors
00418   runtime = wkf_timer_time(timer);
00419   *pdhlatusec = runtime * 1.0e6 / ((double) LATENCYITER);
00420  
00421  
00422   cudaFree(ddata);
00423   CUERR // check and clear any existing errors
00424   cudaFreeHost(phdata);
00425   CUERR // check and clear any existing errors
00426   free(hdata);
00427 
00428   wkf_timer_destroy(timer);
00429 
00430   return 0;
00431 }
00432 
00433 typedef struct {
00434   int deviceid;
00435   double hdmbsec;
00436   double hdlatusec;
00437   double phdmbsec;
00438   double phdlatusec;
00439   double dhmbsec;
00440   double dhlatusec;
00441   double pdhmbsec;
00442   double pdhlatusec;
00443 } busbwthrparms;
00444 
00445 static void * cudabusbwthread(void *voidparms) {
00446   busbwthrparms *parms = (busbwthrparms *) voidparms;
00447   cudabusbw(parms->deviceid, 
00448             &parms->hdmbsec, &parms->hdlatusec,
00449             &parms->phdmbsec, &parms->phdlatusec,
00450             &parms->dhmbsec, &parms->dhlatusec,
00451             &parms->pdhmbsec, &parms->pdhlatusec);
00452   return NULL;
00453 }
00454 
00455 int vmd_cuda_bus_bw(int numdevs, int *devlist, 
00456                     double *hdmbsec, double *hdlatusec,
00457                     double *phdmbsec,double *phdlatusec,
00458                     double *dhmbsec, double *dhlatusec,
00459                     double *pdhmbsec, double *pdhlatusec) {
00460   busbwthrparms *parms;
00461   wkf_thread_t * threads;
00462   int i;
00463 
00464   /* allocate array of threads */
00465   threads = (wkf_thread_t *) calloc(numdevs * sizeof(wkf_thread_t), 1);
00466 
00467   /* allocate and initialize array of thread parameters */
00468   parms = (busbwthrparms *) malloc(numdevs * sizeof(busbwthrparms));
00469   for (i=0; i<numdevs; i++) {
00470     if (devlist != NULL)
00471       parms[i].deviceid = devlist[i];
00472     else
00473       parms[i].deviceid = i;
00474     parms[i].hdmbsec = 0.0;
00475     parms[i].hdlatusec = 0.0;
00476     parms[i].phdmbsec = 0.0;
00477     parms[i].phdlatusec = 0.0;
00478     parms[i].dhmbsec = 0.0;
00479     parms[i].dhlatusec = 0.0;
00480     parms[i].pdhmbsec = 0.0;
00481     parms[i].pdhlatusec = 0.0;
00482   }
00483 
00484 #if defined(VMDTHREADS)
00485   /* spawn child threads to do the work */
00486   /* thread 0 must also be processed this way otherwise    */
00487   /* we'll permanently bind the main thread to some device */
00488   for (i=0; i<numdevs; i++) {
00489     wkf_thread_create(&threads[i], cudabusbwthread, &parms[i]);
00490   }
00491 
00492   /* join the threads after work is done */
00493   for (i=0; i<numdevs; i++) {
00494     wkf_thread_join(threads[i], NULL);
00495   }
00496 #else
00497   /* single thread does all of the work */
00498   cudabusbwthread((void *) &parms[0]);
00499 #endif
00500 
00501   for (i=0; i<numdevs; i++) {
00502     hdmbsec[i] = parms[i].hdmbsec; 
00503     hdlatusec[i] = parms[i].hdlatusec; 
00504     phdmbsec[i] = parms[i].phdmbsec; 
00505     phdlatusec[i] = parms[i].phdlatusec; 
00506     dhmbsec[i] = parms[i].dhmbsec; 
00507     dhlatusec[i] = parms[i].dhlatusec; 
00508     pdhmbsec[i] = parms[i].pdhmbsec; 
00509     pdhlatusec[i] = parms[i].pdhlatusec; 
00510   }
00511 
00512   /* free thread parms */
00513   free(parms);
00514   free(threads);
00515 
00516   return 0;
00517 }
00518 
00519 
00520 
00521 //
00522 // GPU device global memory bandwidth benchmark
00523 //
00524 template <class T>
00525 __global__ void gpuglobmemcpybw(T *dest, const T *src) {
00526   const unsigned int idx = threadIdx.x + blockIdx.x * blockDim.x;
00527   dest[idx] = src[idx];
00528 }
00529 
00530 template <class T>
00531 __global__ void gpuglobmemsetbw(T *dest, const T val) {
00532   int idx = threadIdx.x + blockIdx.x * blockDim.x;
00533   dest[idx] = val;
00534 }
00535 
00536 typedef float4 datatype;
00537 
00538 static int cudaglobmembw(int cudadev, double *gpumemsetgbsec, double *gpumemcpygbsec) {
00539   int i;
00540   int len = 1 << 22; // one thread per data element
00541   int loops = 500;
00542   datatype *src, *dest;
00543   datatype val=make_float4(1.0f, 1.0f, 1.0f, 1.0f);
00544 
00545   // initialize to zero for starters
00546   float memsettime = 0.0f;
00547   float memcpytime = 0.0f;
00548   *gpumemsetgbsec = 0.0;
00549   *gpumemcpygbsec = 0.0;
00550 
00551   // attach to the selected device
00552   cudaError_t rc;
00553   rc = cudaSetDevice(cudadev);
00554   if (rc != cudaSuccess) {
00555 #if CUDART_VERSION >= 2010
00556     rc = cudaGetLastError(); // query last error and reset error state
00557     if (rc != cudaErrorSetOnActiveProcess)
00558       return -1; // abort and return an error
00559 #else
00560     cudaGetLastError(); // just ignore and reset error state, since older CUDA
00561                         // revs don't have a cudaErrorSetOnActiveProcess enum
00562 #endif
00563   }
00564 
00565   cudaMalloc((void **) &src, sizeof(datatype)*len);
00566   CUERR
00567   cudaMalloc((void **) &dest, sizeof(datatype)*len);
00568   CUERR
00569 
00570   dim3 BSz(256, 1, 1);
00571   dim3 GSz(len / (BSz.x * BSz.y * BSz.z), 1, 1); 
00572 
00573   // do a warm-up pass
00574   gpuglobmemsetbw<datatype><<< GSz, BSz >>>(src, val);
00575   CUERR
00576   gpuglobmemsetbw<datatype><<< GSz, BSz >>>(dest, val);
00577   CUERR
00578   gpuglobmemcpybw<datatype><<< GSz, BSz >>>(dest, src);
00579   CUERR
00580 
00581   cudaEvent_t start, end;
00582   cudaEventCreate(&start);
00583   cudaEventCreate(&end);
00584 
00585   // execute the memset kernel
00586   cudaEventRecord(start, 0);
00587   for (i=0; i<loops; i++) {
00588     gpuglobmemsetbw<datatype><<< GSz, BSz >>>(dest, val);
00589   }
00590   CUERR
00591   cudaEventRecord(end, 0);
00592   CUERR
00593   cudaEventSynchronize(start);
00594   CUERR
00595   cudaEventSynchronize(end);
00596   CUERR
00597   cudaEventElapsedTime(&memsettime, start, end);
00598   CUERR
00599 
00600   // execute the memcpy kernel
00601   cudaEventRecord(start, 0);
00602   for (i=0; i<loops; i++) {
00603     gpuglobmemcpybw<datatype><<< GSz, BSz >>>(dest, src);
00604   }
00605   cudaEventRecord(end, 0);
00606   CUERR
00607   cudaEventSynchronize(start);
00608   CUERR
00609   cudaEventSynchronize(end);
00610   CUERR
00611   cudaEventElapsedTime(&memcpytime, start, end);
00612   CUERR
00613 
00614   cudaEventDestroy(start);
00615   CUERR
00616   cudaEventDestroy(end);
00617   CUERR
00618 
00619   *gpumemsetgbsec = (len * sizeof(datatype) / (1024.0 * 1024.0)) / (memsettime / loops);
00620   *gpumemcpygbsec = (2 * len * sizeof(datatype) / (1024.0 * 1024.0)) / (memcpytime / loops);
00621   cudaFree(dest);
00622   cudaFree(src);
00623   CUERR
00624 
00625   return 0;
00626 }
00627 
00628 typedef struct {
00629   int deviceid;
00630   double memsetgbsec;
00631   double memcpygbsec;
00632 } globmembwthrparms;
00633 
00634 static void * cudaglobmembwthread(void *voidparms) {
00635   globmembwthrparms *parms = (globmembwthrparms *) voidparms;
00636   cudaglobmembw(parms->deviceid, &parms->memsetgbsec, &parms->memcpygbsec);
00637   return NULL;
00638 }
00639 
00640 int vmd_cuda_globmem_bw(int numdevs, int *devlist, 
00641                         double *memsetgbsec, double *memcpygbsec) {
00642   globmembwthrparms *parms;
00643   wkf_thread_t * threads;
00644   int i;
00645 
00646   /* allocate array of threads */
00647   threads = (wkf_thread_t *) calloc(numdevs * sizeof(wkf_thread_t), 1);
00648 
00649   /* allocate and initialize array of thread parameters */
00650   parms = (globmembwthrparms *) malloc(numdevs * sizeof(globmembwthrparms));
00651   for (i=0; i<numdevs; i++) {
00652     if (devlist != NULL)
00653       parms[i].deviceid = devlist[i];
00654     else
00655       parms[i].deviceid = i;
00656     parms[i].memsetgbsec = 0.0;
00657     parms[i].memcpygbsec = 0.0;
00658   }
00659 
00660 #if defined(VMDTHREADS)
00661   /* spawn child threads to do the work */
00662   /* thread 0 must also be processed this way otherwise    */
00663   /* we'll permanently bind the main thread to some device */
00664   for (i=0; i<numdevs; i++) {
00665     wkf_thread_create(&threads[i], cudaglobmembwthread, &parms[i]);
00666   }
00667 
00668   /* join the threads after work is done */
00669   for (i=0; i<numdevs; i++) {
00670     wkf_thread_join(threads[i], NULL);
00671   }
00672 #else
00673   /* single thread does all of the work */
00674   cudaglobmembwthread((void *) &parms[0]);
00675 #endif
00676 
00677   for (i=0; i<numdevs; i++) {
00678     memsetgbsec[i] = parms[i].memsetgbsec;
00679     memcpygbsec[i] = parms[i].memcpygbsec;
00680   }
00681 
00682   /* free thread parms */
00683   free(parms);
00684   free(threads);
00685 
00686   return 0;
00687 }
00688 
00689 
00690 //
00691 // Benchmark latency for complete threadpool barrier wakeup/run/sleep cycle
00692 //
00693 static void * vmddevpoollatencythread(void *voidparms) {
00694   return NULL;
00695 }
00696 
00697 static void * vmddevpooltilelatencythread(void *voidparms) {
00698   int threadid=-1;
00699   int tilesize=1;
00700   void *parms=NULL;
00701   wkf_threadpool_worker_getid(voidparms, &threadid, NULL);
00702   wkf_threadpool_worker_getdata(voidparms, (void **) &parms);
00703 
00704   // grind through task tiles until none are left
00705   wkf_tasktile_t tile;
00706   while (wkf_threadpool_next_tile(voidparms, tilesize, &tile) != WKF_SCHED_DONE) {
00707     // do nothing but eat work units...
00708   }
00709 
00710   return NULL;
00711 }
00712 
00713 
00714 // no-op kernel for timing kernel launches
00715 __global__ static void nopkernel(float * ddata) {
00716   unsigned int xindex  = blockIdx.x * blockDim.x + threadIdx.x;
00717   unsigned int yindex  = blockIdx.y * blockDim.y + threadIdx.y;
00718   unsigned int outaddr = gridDim.x * blockDim.x * yindex + xindex;
00719 
00720   if (ddata != NULL)
00721     ddata[outaddr] = outaddr;
00722 }
00723 
00724 // empty kernel for timing kernel launches
00725 __global__ static void voidkernel(void) {
00726   return;
00727 }
00728 
00729 static void * vmddevpoolcudatilelatencythread(void *voidparms) {
00730   int threadid=-1;
00731   int tilesize=1;
00732   float *parms=NULL;
00733   wkf_threadpool_worker_getid(voidparms, &threadid, NULL);
00734 
00735   // XXX Note that we expect parms to be set to NULL or a valid CUDA
00736   //     global memory pointer for correct operation of the NOP kernel below
00737   wkf_threadpool_worker_getdata(voidparms, (void **) &parms);
00738 
00739 #if 0
00740   // scale tile size by device performance
00741   tilesize=4; // GTX 280, Tesla C1060 starting point tile size
00742   wkf_threadpool_worker_devscaletile(voidparms, &tilesize);
00743 #endif
00744 
00745   // grind through task tiles until none are left
00746   wkf_tasktile_t tile;
00747   dim3 Gsz(1,1,0);
00748   dim3 Bsz(8,8,1);
00749   while (wkf_threadpool_next_tile(voidparms, tilesize, &tile) != WKF_SCHED_DONE) {
00750     // launch a no-op CUDA kernel
00751     nopkernel<<<Gsz, Bsz, 0>>>(parms);
00752   }
00753 
00754   // wait for all GPU kernels to complete
00755   cudaDeviceSynchronize();
00756 
00757   return NULL;
00758 }
00759 
00760 
00761 int vmd_cuda_devpool_latency(wkf_threadpool_t *devpool, int tilesize,
00762                              double *kernlaunchlatency,
00763                              double *barlatency,
00764                              double *cyclelatency, 
00765                              double *tilelatency,
00766                              double *kernellatency) {
00767   int i;
00768   wkf_tasktile_t tile;
00769   wkf_timerhandle timer;
00770   int loopcount;
00771 
00772   timer=wkf_timer_create();
00773 
00774   // execute just a CUDA kernel launch and measure latency on whatever
00775   // GPU we get.
00776   loopcount = 15000;
00777   dim3 VGsz(1,1,0);
00778   dim3 VBsz(8,8,1);
00779   wkf_timer_start(timer);
00780   for (i=0; i<loopcount; i++) {
00781     voidkernel<<<VGsz, VBsz, 0>>>();
00782   }
00783   // wait for GPU kernels to complete
00784   cudaDeviceSynchronize();
00785   wkf_timer_stop(timer);
00786   *kernlaunchlatency = wkf_timer_time(timer) / ((double) loopcount);
00787 
00788   // execute just a raw barrier sync and measure latency
00789   loopcount = 15000;
00790   wkf_timer_start(timer);
00791   for (i=0; i<loopcount; i++) {
00792     wkf_threadpool_wait(devpool);
00793   }
00794   wkf_timer_stop(timer);
00795   *barlatency = wkf_timer_time(timer) / ((double) loopcount);
00796 
00797   // time wake-up, launch, and sleep/join of device pool doing a no-op
00798   loopcount = 5000;
00799   wkf_timer_start(timer);
00800   for (i=0; i<loopcount; i++) {
00801     tile.start=0;
00802     tile.end=0;
00803     wkf_threadpool_sched_dynamic(devpool, &tile);
00804     wkf_threadpool_launch(devpool, vmddevpoollatencythread, NULL, 1);
00805   }
00806   wkf_timer_stop(timer);
00807   *cyclelatency = wkf_timer_time(timer) / ((double) loopcount);
00808 
00809   // time wake-up, launch, and sleep/join of device pool eating tiles
00810   loopcount = 5000;
00811   wkf_timer_start(timer);
00812   for (i=0; i<loopcount; i++) {
00813     tile.start=0;
00814     tile.end=tilesize;
00815     wkf_threadpool_sched_dynamic(devpool, &tile);
00816     wkf_threadpool_launch(devpool, vmddevpooltilelatencythread, NULL, 1);
00817   }
00818   wkf_timer_stop(timer);
00819   *tilelatency = wkf_timer_time(timer) / ((double) loopcount);
00820 
00821   // time wake-up, launch, and sleep/join of device pool eating tiles
00822   loopcount = 2000;
00823   wkf_timer_start(timer);
00824   for (i=0; i<loopcount; i++) {
00825     tile.start=0;
00826     tile.end=tilesize;
00827     wkf_threadpool_sched_dynamic(devpool, &tile);
00828     wkf_threadpool_launch(devpool, vmddevpoolcudatilelatencythread, NULL, 1);
00829   }
00830   wkf_timer_stop(timer);
00831   *kernellatency = wkf_timer_time(timer) / ((double) loopcount);
00832 
00833   wkf_timer_destroy(timer);
00834 
00835 #if 1
00836   vmd_cuda_measure_latencies(devpool);
00837 #endif
00838 
00839   return 0;
00840 }
00841 
00842 
00843 //
00844 // Benchmark CUDA kernel launch and memory copy latencies in isolation
00845 //
00846 typedef struct {
00847   int deviceid;
00848   int testloops;
00849   double kernlatency;
00850   double bcopylatency;
00851   double kbseqlatency;
00852 } latthrparms;
00853 
00854 static void * vmddevpoolcudalatencythread(void *voidparms) {
00855   int threadid=-1;
00856   latthrparms *parms=NULL;
00857 
00858   wkf_threadpool_worker_getid(voidparms, &threadid, NULL);
00859   wkf_threadpool_worker_getdata(voidparms, (void **) &parms);
00860   if (parms->deviceid == threadid) { 
00861     wkf_timerhandle timer;
00862     timer=wkf_timer_create();
00863     printf("Thread/device %d running...\n", threadid);
00864     cudaStream_t devstream;
00865     cudaStreamCreate(&devstream);
00866 
00867     char *hostbuf = (char *) calloc(1, 65536 * sizeof(char));
00868     char  *gpubuf = NULL;
00869     cudaMalloc((void**)&gpubuf, 65536 * sizeof(char));
00870 
00871     dim3 Gsz(1,1,0);
00872     dim3 Bsz(8,8,1);
00873 
00874     // measure back-to-back NULL kernel launches
00875     wkf_timer_start(timer);
00876     int i;
00877     for (i=0; i<parms->testloops; i++) {
00878       // launch a no-op CUDA kernel
00879       nopkernel<<<Gsz, Bsz, 0, devstream>>>(NULL);
00880     }
00881     // wait for all GPU kernels to complete
00882     cudaStreamSynchronize(devstream);
00883     wkf_timer_stop(timer);
00884     parms->kernlatency =  1000000 * wkf_timer_time(timer) / ((double) parms->testloops);
00885 
00886     // measure back-to-back round-trip 1-byte memcpy latencies
00887     wkf_timer_start(timer);
00888     for (i=0; i<parms->testloops; i++) {
00889       cudaMemcpyAsync(gpubuf, hostbuf, 1, cudaMemcpyHostToDevice, devstream);
00890       cudaMemcpyAsync(hostbuf, gpubuf, 1, cudaMemcpyDeviceToHost, devstream);
00891     }
00892     // wait for all GPU kernels to complete
00893     cudaStreamSynchronize(devstream);
00894     wkf_timer_stop(timer);
00895     parms->kernlatency =  1000000 * wkf_timer_time(timer) / ((double) parms->testloops);
00896 
00897     printf("NULL kernel launch latency (usec): %.2f\n", parms->kernlatency);
00898 
00899     cudaStreamDestroy(devstream);
00900     cudaFree(gpubuf);
00901     free(hostbuf);
00902     wkf_timer_destroy(timer);
00903   }
00904 
00905   return NULL;
00906 }
00907 
00908 
00909 int vmd_cuda_measure_latencies(wkf_threadpool_t *devpool) {
00910   latthrparms thrparms;
00911   int workers = wkf_threadpool_get_workercount(devpool);
00912   int i;
00913 printf("vmd_cuda_measure_latencies()...\n");
00914   for (i=0; i<workers; i++) {
00915     memset(&thrparms, 0, sizeof(thrparms));
00916     thrparms.deviceid = i;
00917     thrparms.testloops = 2500;
00918     wkf_threadpool_launch(devpool, vmddevpoolcudalatencythread, &thrparms, 1);
00919   }
00920 
00921   return 0;
00922 }
00923 
00924 
00925 #if defined(VMDUSECUDAGDS)
00926 typedef struct {
00927   int nfiles;
00928   const char **trjfileset;
00929   jshandle **jshandles;
00930   CUfileHandle_t *cfh;
00931   int devcount;
00932   int natoms;
00933   const AtomSel *sel;
00934   int first;
00935   int last;
00936   int step;
00937 } gpuoocbenchthreadparms;
00938 
00939 #define VMDGDSMAXFRAMEBUF     8
00940 static void * gpu_ooc_bench_thread(void *voidparms) {
00941   int threadid, numthreads;
00942   gpuoocbenchthreadparms *parms = NULL;
00943   wkf_threadpool_worker_getdata(voidparms, (void **) &parms);
00944   wkf_threadpool_worker_getid(voidparms, &threadid, &numthreads);
00945 
00946   //
00947   // copy in per-thread parameters
00948   //
00949   int nfiles = parms->nfiles;
00950   int natoms = parms->natoms;
00951   const AtomSel *sel = parms->sel;
00952   int first  = parms->first;
00953   int last   = parms->last;
00954   int step   = parms->step;
00955 
00956   int usecufile = 1;
00957   fio_fd *hostfds = NULL;
00958   int pinhostiobuffer = 1;
00959   if (getenv("VMDGDSHOSTNOPIN")) {
00960     pinhostiobuffer=0;
00961   }
00962 
00963   if (getenv("VMDGDSUSEHOST")) {
00964     usecufile=0;
00965     hostfds = (fio_fd *) calloc(1, nfiles * sizeof(fio_fd));
00966 
00967     int hostusedirectio = 1;
00968     if (getenv("VMDGDSHOSTBUFFEREDIO") != NULL)
00969       hostusedirectio = 0;
00970 
00971     int openmode = FIO_READ;
00972     if (hostusedirectio)
00973       openmode |= FIO_DIRECT;
00974 
00975     int i;
00976     for (i=0; i<nfiles; i++) {
00977       if (fio_open(parms->trjfileset[i], openmode, &hostfds[i]) < 0) {
00978         if (hostusedirectio) {
00979           printf("Thr[%d] direct I/O unavailable or can't open file '%s'\n",
00980                  threadid, parms->trjfileset[i]);
00981         } else {
00982           printf("Thr[%d] can't open file '%s'\n",
00983                  threadid, parms->trjfileset[i]);
00984         }
00985         return NULL;
00986       }
00987     }
00988   }
00989 
00990   if (hostfds && usecufile) {
00991     printf("Inconsistent cufile/hostfds state, aborting!\n");
00992     return NULL;
00993   }
00994 
00995   /* ensure we have a large enough allocation so we can align */
00996   /* the starting pointer to a blocksz page boundary          */
00997   long blocksz = MOLFILE_DIRECTIO_MIN_BLOCK_SIZE;
00998   long sz = 3L*sizeof(float)*natoms + blocksz;
00999 
01000   /* pad the allocation to an even multiple of the block size */
01001   size_t blockpadsz = (sz + (blocksz - 1)) & (~(blocksz - 1));
01002 
01003   int framecount = (last - first + 1) / step;
01004   int framesperfile = framecount / nfiles;
01005 
01006   if (threadid == 0) {
01007     printf("Thr[%2d] %d frames total, natoms: %d selected: %d\n",
01008            threadid, framecount, natoms, sel->selected);
01009     printf("Thr[%2d] %d frames/file\n", threadid, framesperfile);
01010   }
01011 
01012   cudaError_t crc;
01013   cudaStream_t oocstream;
01014   float *devptr=NULL;
01015   float *hostptr=NULL;
01016   float *hostptr_unaligned=NULL;
01017 
01018   float *crdx1=NULL, *crdy1=NULL, *crdz1=NULL;
01019   float *crdx2=NULL, *crdy2=NULL, *crdz2=NULL;
01020   int multiframeio = 0;
01021   if (getenv("VMDGDSMULTIFRAME"))
01022     multiframeio = atoi(getenv("VMDGDSMULTIFRAME"));
01023   if (multiframeio > VMDGDSMAXFRAMEBUF)
01024     multiframeio = VMDGDSMAXFRAMEBUF;
01025 
01026   // set block sizes and counts for IO bench calcs
01027   dim3 IOBsz = dim3(256, 1, 1);
01028   dim3 IOGsz = dim3((natoms + IOBsz.x - 1) / IOBsz.x, 1, 1);
01029 
01030   if (parms->devcount > 0) {
01031     long gpuallocsz = (VMDGDSMAXFRAMEBUF+1) * blockpadsz;
01032 
01033     if (threadid == 0) {
01034       printf("Thr[%2d] Allocating GPU timestep I/O buf: %ld \n",
01035              threadid, gpuallocsz);
01036     }
01037     crc = cudaMalloc((void**) &devptr, gpuallocsz);
01038 
01039     if (hostfds != NULL) {
01040       if (pinhostiobuffer) {
01041         crc = cudaMallocHost((void**) &hostptr, gpuallocsz);
01042       } else {
01043         hostptr = (float *) alloc_aligned_ptr(gpuallocsz, 4096,
01044                                               (void**) &hostptr_unaligned);
01045         if (!hostptr) {
01046           printf("Thr[%d]: Failed allocation!\n", threadid);
01047           return NULL;
01048         }
01049       }
01050     }
01051 
01052     long crdsz = sel->selected * sizeof(float);
01053 
01054     // atomic coord buffers
01055     crc = cudaMalloc((void**) &crdx1, crdsz);
01056     crc = cudaMalloc((void**) &crdy1, crdsz);
01057     crc = cudaMalloc((void**) &crdz1, crdsz);
01058     crc = cudaMalloc((void**) &crdx2, crdsz);
01059     crc = cudaMalloc((void**) &crdy2, crdsz);
01060     crc = cudaMalloc((void**) &crdz2, crdsz);
01061     if (crc != cudaSuccess) {
01062       printf("Thr[%2d], Failed to allocate GPU buffer!\n", threadid);
01063       return NULL; // XXX error handling needs to be done here
01064     }
01065 
01066     cudaStreamCreate(&oocstream);
01067 
01068 #if defined(VMDUSECUDAGDS)
01069     cuFileBufRegister(devptr, gpuallocsz, 0);
01070 #endif
01071   }
01072 
01073   int verbose = (getenv("VMDGDSVERBOSE") != NULL) ? 1 : 0;
01074   
01075   int filestrategy = 0;
01076   if (getenv("VMDGDSFILESTRATEGY")) {
01077     filestrategy = atoi(getenv("VMDGDSFILESTRATEGY"));
01078   }
01079   if (threadid == 0) {
01080     printf("Thr[%2d] file strategy set to: %d\n", threadid, filestrategy);
01081   }
01082 
01083   wkf_tasktile_t tile;
01084   while (wkf_threadlaunch_next_tile(voidparms, VMDGDSMAXFRAMEBUF * 1, &tile) != WKF_SCHED_DONE) {
01085     //
01086     // simple I/O + compute benchmarking...
01087     //
01088     int idx;
01089     int threadspergpu;
01090     if (parms->devcount > 0)
01091       threadspergpu = numthreads / parms->devcount;
01092     else 
01093       threadspergpu = 1;
01094 
01095     for (idx=tile.start; idx<tile.end; idx++) {
01096       int myfileidx, fileframeidx;
01097 
01098       switch (filestrategy) {
01099         case 1:
01100           myfileidx = (idx / multiframeio) % nfiles;
01101           fileframeidx = idx % framesperfile;
01102           break;
01103 
01104         case 2:
01105           myfileidx = (idx / (multiframeio * threadspergpu)) % nfiles;
01106           fileframeidx = idx % framesperfile;
01107           break;
01108 
01109         case 3:
01110           myfileidx = (threadid / 4) % nfiles;
01111           fileframeidx = idx % framesperfile;
01112           break;
01113 
01114         case 0:
01115         default:
01116           myfileidx = (threadid / threadspergpu) % nfiles;
01117           fileframeidx = idx % framesperfile;
01118           break;
01119       }
01120 
01121       //
01122       // compute multi-frame or single-frame I/O offsets and sizes
01123       //
01124       long startoffset, foffset, readlen;
01125       read_js_timestep_index_offsets(parms->jshandles[myfileidx],
01126                                      natoms, fileframeidx, 0, natoms, NULL,
01127                                      &startoffset, &foffset, &readlen);
01128       if (multiframeio) {
01129         // multi-frame reads use the same starting offset, but the 
01130         // read length is computed from the first and last frames 
01131         // in the group.
01132         long multistartoffset, multifoffset, multireadlen;
01133         read_js_timestep_index_offsets(parms->jshandles[myfileidx], natoms,
01134                                        fileframeidx+multiframeio-1,
01135                                        0, natoms, NULL,
01136                                        &multistartoffset, &multifoffset,
01137                                        &multireadlen);
01138 
01139         multireadlen = (multifoffset + multireadlen) - foffset;
01140 
01141         //printf("** readlen: %ld  multireadlen: %ld\n", readlen, multireadlen);
01142         readlen = multireadlen;
01143         idx+=multiframeio-1; // add in the required increment...
01144       }
01145 
01146       //
01147       // perform the required I/O via GDS or by host kernel I/O
01148       //
01149       long ret=0;
01150       if (usecufile) {
01151         ret = cuFileRead(parms->cfh[myfileidx], (char *) devptr, readlen, foffset, 0);
01152       } else if (hostfds) {
01153         foffset=0;
01154         ret=fio_fseek(hostfds[myfileidx], foffset, FIO_SEEK_SET);
01155         if (ret<0) { printf("fio_fseek() error!\n"); return NULL; }
01156         ret=fio_fread(hostptr, readlen, 1, hostfds[myfileidx]);
01157         if (ret<0) { printf("fio_fseek() error!\n"); return NULL; }
01158         cudaMemcpy(devptr, hostptr, readlen, cudaMemcpyHostToDevice);
01159       } else {
01160         printf("Inconsistent cufile/hostfds state, aborting!\n");
01161         return NULL;
01162       }
01163 
01164       // handle errors if they have occured
01165       if (ret < 0) {
01166         printf("Thr[%2d] Error: cuFileRead(): %ld\n", threadid, ret);
01167         return NULL; // XXX error handling needs to be done here
01168       }
01169 
01170       if (verbose) {
01171         printf("Thr[%2d]F[%d][tile: %d to %d] frame: %d cuFile len: %ld off: %ld\n",
01172                threadid, myfileidx, tile.start, tile.end, idx, 
01173                readlen, foffset);
01174       }
01175     }
01176   }
01177 
01178   cudaFree(crdx1);
01179   cudaFree(crdy1);
01180   cudaFree(crdz1);
01181   cudaFree(crdx2);
01182   cudaFree(crdy2);
01183   cudaFree(crdz2);
01184 
01185 #if defined(VMDUSECUDAGDS)
01186   if (usecufile) {
01187     cuFileBufDeregister(devptr);
01188   }
01189 #endif
01190 
01191   if (hostfds != NULL) {
01192     int i;
01193     for (i=0; i<nfiles; i++) {
01194       fio_fclose(hostfds[i]);
01195     }
01196     free(hostfds);
01197   }
01198 
01199   if (hostptr != NULL) {
01200     if (pinhostiobuffer) {
01201       cudaFreeHost(hostptr);
01202     } else {
01203       free(hostptr_unaligned);
01204     }
01205   }
01206 
01207   return NULL;
01208 }
01209 
01210 #endif
01211 
01212 
01213 int gpu_ooc_bench(wkf_threadpool_t *devpool, // VMD GPU worker thread pool
01214                   int nfiles, const char **trjfileset, const AtomSel *sel,
01215                   int first, int last, int step) {
01216   printf("gpu_ooc_bench()\n");
01217   wkf_threadpool_t *bigpool = NULL;
01218 
01219 #if defined(VMDUSECUDAGDS)
01220   int devcount;
01221   cudaError_t crc = cudaGetDeviceCount(&devcount);
01222   printf("gpu_ooc_bench) GPU device count: %d\n", devcount);
01223   if (devcount==0)
01224     printf("gpu_ooc_bench)   No GPU devices, continuing with host only...\n");
01225 
01226   CUfileHandle_t * cfh = (CUfileHandle_t *) calloc(1, nfiles * sizeof(CUfileHandle_t));
01227   CUfileDescr_t * cfhdesc = (CUfileDescr_t *) calloc(1, nfiles * sizeof(CUfileDescr_t));
01228   memset(&cfh[0], 0, sizeof(cfh));
01229   memset(&cfhdesc[0], 0, sizeof(cfhdesc));
01230 
01231   int natomschk = 0;
01232   jshandle **jshandles = (jshandle **) calloc(1, nfiles * sizeof(jshandle *));
01233   fio_fd *directio_fds = (fio_fd *) calloc(1, nfiles * sizeof(fio_fd));
01234 
01235   int i;
01236   for (i=0; i<nfiles; i++) {
01237     const char *filename = trjfileset[i];
01238     printf("gpu_ooc_bench) File[%d] GDS setup, opening '%s'\n", i, filename);
01239     jshandles[i] = (jshandle *) open_js_read(filename, "js", &natomschk);
01240     if (!jshandles[i]) {
01241       printf("gpu_ooc_bench) File[%d] open_js_read failed for file %s\n", i, filename);
01242       return -1; // deal with error handling later
01243     }
01244 
01245 #if vmdplugin_ABIVERSION > 17
01246     long blocksz = MOLFILE_DIRECTIO_MIN_BLOCK_SIZE;
01247     int filepgalignsz = 1;
01248     read_js_timestep_pagealign_size(jshandles[i], &filepgalignsz);
01249     if (filepgalignsz != blocksz) {
01250       printf("gpu_ooc_bench) File[%d] Plugin-returned page alignment size mismatch!\n", i);
01251     } else {
01252       printf("gpu_ooc_bench) File[%d] Page alignment size: %d\n", i, filepgalignsz);
01253     }
01254 #endif
01255 
01256     read_js_timestep_index_offsets(jshandles[i], natomschk, 0, 0, 0,
01257                                    &directio_fds[i], NULL, NULL, NULL);
01258 
01259     cfhdesc[i].handle.fd = directio_fds[i]; // typedef of Unix FD
01260     cfhdesc[i].type = CU_FILE_HANDLE_TYPE_OPAQUE_FD;
01261     CUfileError_t cferr = cuFileHandleRegister(&cfh[i], &cfhdesc[i]);
01262 
01263     if (cferr.err != CU_FILE_SUCCESS) {
01264       printf("gpu_ooc_bench) File[%d] cuFileImportExternalFile on fd %d failed!\n",
01265              i, cfhdesc[i].handle.fd);
01266       return -1; // XXX error handling needs to be done here
01267     }
01268   }
01269 
01270 
01271   //
01272   // copy in per-thread parameters
01273   //
01274   gpuoocbenchthreadparms parms;
01275   memset(&parms, 0, sizeof(parms));
01276   parms.devcount = devcount;
01277   parms.nfiles = nfiles;
01278   parms.trjfileset = trjfileset;
01279   parms.jshandles = jshandles;
01280   parms.cfh = cfh;
01281   parms.natoms = sel->num_atoms;
01282   parms.sel = sel;
01283   parms.first = first;
01284   parms.last = last;
01285   parms.step = step;
01286 
01287   int framecount = nfiles * (last / step);
01288 
01289   // create timers
01290   wkf_timerhandle timer;
01291   timer=wkf_timer_create();
01292 
01293   // spawn child threads to do the work
01294   wkf_tasktile_t tile;
01295   tile.start=0;
01296   tile.end=framecount - 1; // first row only
01297 
01298   printf("gpu_ooc_bench) tile start: %d  end: %d\n", tile.start, tile.end);
01299 
01300   int gdsthreadspergpu = 1;
01301   if (getenv("VMDGDSTHREADSPERGPU") != NULL)
01302     gdsthreadspergpu = atoi(getenv("VMDGDSTHREADSPERGPU"));
01303 
01304   printf("gpu_ooc_bench) gdsthreadspergpu: %d\n", gdsthreadspergpu);
01305 
01306   if (gdsthreadspergpu > 1) {
01307     // XXX extra-large GPU device thread pool
01308     int workercount = devcount * gdsthreadspergpu;
01309 
01310     int *devlist = new int[workercount];
01311     int k;
01312     for (k=0; k<workercount; k++) {
01313       devlist[k] = k / gdsthreadspergpu; // XXX ignores VMD CUDA device masks
01314     }
01315 
01316     msgInfo << "Creating Multi-worker ("
01317             << gdsthreadspergpu << " per-GPU) CUDA device pool..." << sendmsg;
01318     bigpool=wkf_threadpool_create(workercount, devlist);
01319     delete [] devlist;
01320 
01321     // associate each worker thread with a specific GPU
01322     if (getenv("VMDCUDAVERBOSE") != NULL)
01323       wkf_threadpool_launch(bigpool, vmd_cuda_devpool_setdeviceonly, (void*)"VMD CUDA Dev Init", 1);
01324     else
01325       wkf_threadpool_launch(bigpool, vmd_cuda_devpool_setdeviceonly, NULL, 1);
01326 
01327     // clear all available device memory on each of the GPUs
01328     wkf_threadpool_launch(bigpool, vmd_cuda_devpool_clear_device_mem, NULL, 1);
01329 
01330     // XXX override which GPU device pool we're going to use
01331     devpool = bigpool;
01332   }
01333 
01334   // XXX affinitize GPU worker threads for best perf
01335   wkf_threadpool_launch(devpool, vmd_cuda_affinitize_threads, NULL, 1);
01336 
01337   wkf_threadpool_sched_dynamic(devpool, &tile);
01338   wkf_timer_start(timer);
01339   wkf_threadpool_launch(devpool, gpu_ooc_bench_thread, &parms, 1);
01340   wkf_timer_stop(timer);
01341 
01342   double runtime = wkf_timer_time(timer);
01343   double gbytes = sel->num_atoms * 12L * (tile.end+1) / (1024.0 * 1024.0 * 1024.0);
01344 
01345   printf("gpu_ooc_bench) natoms: %d, fsz: %ld, tsz: %ld\n",
01346          sel->num_atoms, sel->num_atoms * 12L,
01347          sel->num_atoms * 12L * (tile.end+1));
01348 
01349   int pinhostiobuffer = 1;
01350   if (getenv("VMDGDSHOSTNOPIN"))
01351     pinhostiobuffer=0;
01352 
01353   int hostusedirectio = 1;
01354   if (getenv("VMDGDSHOSTBUFFEREDIO") != NULL)
01355     hostusedirectio = 0;
01356 
01357   int usecufile=1;
01358   if (getenv("VMDGDSUSEHOST"))
01359     usecufile=0;
01360 
01361   if (usecufile) {
01362     printf("OOC I/O via GDS + cuFile\n");
01363   } else {
01364     printf("OOC I/O via host, %s APIs, %s memory buffers\n",
01365            (hostusedirectio) ? "Direct I/O" : "Buffered I/O",
01366            (pinhostiobuffer) ? "pinned" : "unpinned");
01367   }
01368 
01369   int multiframeio = 0;
01370   if (getenv("VMDGDSMULTIFRAME"))
01371     multiframeio = atoi(getenv("VMDGDSMULTIFRAME"));
01372   if (multiframeio > VMDGDSMAXFRAMEBUF)
01373     multiframeio = VMDGDSMAXFRAMEBUF;
01374   if (multiframeio) {
01375     printf("GDS multi-frame read opt: %d frames per call, %ld bytes\n",
01376            multiframeio,
01377            multiframeio * sel->num_atoms * 12L);
01378   }
01379 
01380   printf("OOC runtime: %.1f, %.2fGB/sec\n", runtime, gbytes/runtime);
01381 
01382   for (i=0; i<nfiles; i++) {
01383 #if defined(VMDUSECUDAGDS)
01384     cuFileHandleDeregister(cfh[i]);
01385 #endif
01386     close_js_read(jshandles[i]);
01387   }
01388 #endif
01389 
01390 #if defined(VMDUSECUDAGDS)
01391   if (cfh != NULL)
01392     free(cfh);
01393 
01394   if (cfhdesc != NULL)
01395      free(cfhdesc);
01396 
01397   if (jshandles != NULL)
01398     free(jshandles);
01399 
01400   if (directio_fds != NULL)
01401     free(directio_fds);
01402 #endif
01403 
01404   // if we created an extra large-thread-count-per-GPU thread pool, we
01405   // need to destroy it here...
01406   if (bigpool != NULL)
01407     wkf_threadpool_destroy(bigpool);
01408 
01409   return 0;
01410 } 
01411 
01412 
01413 

Generated on Fri Oct 4 02:43:49 2024 for VMD (current) by doxygen1.2.14 written by Dimitri van Heesch, © 1997-2002