Fix cufile (#438)

* Added memory type to helper functions.
* Add code for invalidating buffer regardless of device.
* Support build using nvcc
* Supporting more options. Remove any direct memory access.
* Bugfix IOR: ignoring timestamp (if not set).

configure.ac

@@ -119,8 +119,16 @@ AS_IF([test "$ac_cv_header_cufile_h" = "yes"], [
     ])
 ])
 AM_CONDITIONAL([HAVE_GPU_DIRECT], [test x$with_gpuDirect = xyes])
-AM_COND_IF([HAVE_GPU_DIRECT],[AC_DEFINE([HAVE_GPU_DIRECT], [], [GPUDirect API found])])
-
+AM_COND_IF([HAVE_GPU_DIRECT],[
+  AC_DEFINE([HAVE_GPU_DIRECT], [], [GPUDirect API found])
+  ])
+
+varNVCC=nvcc
+AC_ARG_WITH(nvcc,
+              AS_HELP_STRING([--with-nvcc], [Use the NVCC as specified]),
+              [AS_IF([test "$with_nvcc" != "yes"], varNVCC=$with_nvcc)],
+              [])
+AC_SUBST(NVCC, $varNVCC)
 
 # Check for system capabilities
 AC_SYS_LARGEFILE

src/Makefile.am

@@ -47,6 +47,7 @@ endif
 
 if HAVE_GPU_DIRECT
 extraLDADD    += -lcufile
+extraSOURCES  += utilities-gpu.cu
 endif
 
 if USE_HDF5_AIORI
@@ -177,3 +178,6 @@ all-local: build.conf
 build.conf:
 	@echo LDFLAGS=$(LDFLAGS) $(extraLDFLAGS) $(extraLDADD) $(LIBS) > build.conf
 	@echo CFLAGS=$(CFLAGS) $(extraCPPFLAGS) >> build.conf
+
+.cu.o:
+	$(NVCC) $(CFLAGS) -c -o $@ $<

src/aiori-POSIX.c

@@ -722,9 +722,9 @@ static IOR_offset_t POSIX_Xfer(int access, aiori_fd_t *file, IOR_size_t * buffer
 #ifdef HAVE_GPU_DIRECT
                         }
 #endif
-                        if (rc == -1)
-                                ERRF("write(%d, %p, %lld) failed",
-                                        fd, (void*)ptr, remaining);
+                        if (rc < 0){
+                          WARNF("write(%d, %p, %lld) failed %s", fd, (void*)ptr, remaining, strerror(errno));
+                        }
                         if (hints->fsyncPerWrite == TRUE){
                           POSIX_Fsync((aiori_fd_t*) &fd, param);
                         }
@@ -748,8 +748,8 @@ static IOR_offset_t POSIX_Xfer(int access, aiori_fd_t *file, IOR_size_t * buffer
                           return length - remaining;
                         }
 
-                        if (rc == -1){
-                          WARNF("read(%d, %p, %lld) failed", fd, (void*)ptr, remaining);
+                        if (rc < 0){
+                          WARNF("read(%d, %p, %lld) failed %s", fd, (void*)ptr, remaining, strerror(errno));
                           return length - remaining;
                         }
                 }

src/ior.c

@@ -419,7 +419,7 @@ static size_t
 CompareData(void *expectedBuffer, size_t size, IOR_param_t *test, IOR_offset_t offset, int fillrank, int access)
 {
         assert(access == WRITECHECK || access == READCHECK);
-        return verify_memory_pattern(offset, expectedBuffer, size, test->setTimeStampSignature, fillrank, test->dataPacketType);
+        return verify_memory_pattern(offset, expectedBuffer, size, test->timeStampSignatureValue, fillrank, test->dataPacketType, test->gpuMemoryFlags);
 }
 
 /*
@@ -1185,6 +1185,7 @@ static void TestIoSys(IOR_test_t *test)
         if (params->setTimeStampSignature) { // initialize the buffer properly
                 params->timeStampSignatureValue = (unsigned int) params->setTimeStampSignature;
         }
+
         XferBuffersSetup(&ioBuffers, params, pretendRank);
 
         /* Initial time stamp */
@@ -1212,8 +1213,7 @@ static void TestIoSys(IOR_test_t *test)
                                 if ((currentTime = time(NULL)) == -1) {
                                         ERR("cannot get current time");
                                 }
-                                params->timeStampSignatureValue =
-                                        (unsigned int)currentTime;
+                                params->timeStampSignatureValue = (unsigned int)currentTime;
                         }
                         if (verbose >= VERBOSE_2) {
                                 fprintf(out_logfile,
@@ -1229,7 +1229,7 @@ static void TestIoSys(IOR_test_t *test)
                           (&params->timeStampSignatureValue, 1, MPI_UNSIGNED, 0,
                            testComm), "cannot broadcast start time value");
 
-                generate_memory_pattern((char*) ioBuffers.buffer, params->transferSize, params->setTimeStampSignature, pretendRank, params->dataPacketType);
+                generate_memory_pattern((char*) ioBuffers.buffer, params->transferSize, params->timeStampSignatureValue, pretendRank, params->dataPacketType, params->gpuMemoryFlags);
 
                 /* use repetition count for number of multiple files */
                 if (params->multiFile)
@@ -1650,7 +1650,7 @@ static IOR_offset_t WriteOrReadSingle(IOR_offset_t offset, int pretendRank, IOR_
   if (access == WRITE) {
           /* fills each transfer with a unique pattern
            * containing the offset into the file */
-          update_write_memory_pattern(offset, ioBuffers->buffer, transfer, test->setTimeStampSignature, pretendRank, test->dataPacketType);
+          update_write_memory_pattern(offset, ioBuffers->buffer, transfer, test->setTimeStampSignature, pretendRank, test->dataPacketType, test->gpuMemoryFlags);
           amtXferred = backend->xfer(access, fd, buffer, transfer, offset, test->backend_options);
           if (amtXferred != transfer)
                   ERR("cannot write to file");
@@ -1669,13 +1669,13 @@ static IOR_offset_t WriteOrReadSingle(IOR_offset_t offset, int pretendRank, IOR_
             nanosleep( & wait, NULL);
           }
   } else if (access == WRITECHECK) {
-          ((long long int*) buffer)[0] = ~((long long int*) buffer)[0]; // changes the buffer, no memset to reduce the memory pressure
+          invalidate_buffer_pattern(buffer, transfer, test->gpuMemoryFlags);
           amtXferred = backend->xfer(access, fd, buffer, transfer, offset, test->backend_options);
           if (amtXferred != transfer)
                   ERR("cannot read from file write check");
           *errors += CompareData(buffer, transfer, test, offset, pretendRank, WRITECHECK);
   } else if (access == READCHECK) {
-          ((long long int*) buffer)[0] = ~((long long int*) buffer)[0]; // changes the buffer, no memset to reduce the memory pressure
+          invalidate_buffer_pattern(buffer, transfer, test->gpuMemoryFlags);          
           amtXferred = backend->xfer(access, fd, buffer, transfer, offset, test->backend_options);
           if (amtXferred != transfer){
             ERR("cannot read from file");

src/ior.h

@@ -47,12 +47,6 @@
 
 #define ISPOWEROFTWO(x) ((x != 0) && !(x & (x - 1)))
 
-typedef enum{
-    IOR_MEMORY_TYPE_CPU = 0,
-    IOR_MEMORY_TYPE_GPU_MANAGED = 1,
-    IOR_MEMORY_TYPE_GPU_DEVICE_ONLY = 2,
-} ior_memory_flags;
-
 
 /***************** IOR_BUFFERS *************************************************/
 /* A struct to hold the buffers so we can pass 1 pointer around instead of 3

src/iordef.h

@@ -26,6 +26,12 @@ typedef enum {
   DATA_RANDOM           /* fully scrambled blocks */
 } ior_dataPacketType_e;
 
+typedef enum{
+    IOR_MEMORY_TYPE_CPU = 0,
+    IOR_MEMORY_TYPE_GPU_MANAGED = 1,
+    IOR_MEMORY_TYPE_GPU_DEVICE_ONLY = 2,
+} ior_memory_flags;
+
 #ifdef _WIN32
 #   define _CRT_SECURE_NO_WARNINGS
 #   define _CRT_RAND_S

src/md-workbench.c

@@ -556,7 +556,7 @@ void run_precreate(phase_stat_t * s, int current_index){
   }
 
   char * buf = aligned_buffer_alloc(o.file_size, o.gpu_memory_flags);
-  generate_memory_pattern(buf, o.file_size, o.random_seed, o.rank, o.dataPacketType);
+  generate_memory_pattern(buf, o.file_size, o.random_seed, o.rank, o.dataPacketType, o.gpu_memory_flags);
   double op_timer; // timer for individual operations
   size_t pos = -1; // position inside the individual measurement array
   double op_time;
@@ -572,7 +572,7 @@ void run_precreate(phase_stat_t * s, int current_index){
       if (NULL == aiori_fh){
         FAIL("Unable to open file %s", obj_name);
       }
-      update_write_memory_pattern(f * o.dset_count + d, buf, o.file_size, o.random_seed, o.rank, o.dataPacketType);
+      update_write_memory_pattern(f * o.dset_count + d, buf, o.file_size, o.random_seed, o.rank, o.dataPacketType, o.gpu_memory_flags);
       if ( o.file_size == (int) o.backend->xfer(WRITE, aiori_fh, (IOR_size_t *) buf, o.file_size, 0, o.backend_options)) {
         s->obj_create.suc++;
       }else{
@@ -598,7 +598,7 @@ void run_benchmark(phase_stat_t * s, int * current_index_p){
   char obj_name[MAX_PATHLEN];
   int ret;
   char * buf = aligned_buffer_alloc(o.file_size, o.gpu_memory_flags);
-  memset(buf, o.rank % 256, o.file_size);
+  invalidate_buffer_pattern(buf, o.file_size, o.gpu_memory_flags);
   double op_timer; // timer for individual operations
   size_t pos = -1; // position inside the individual measurement array
   int start_index = *current_index_p;
@@ -653,7 +653,7 @@ void run_benchmark(phase_stat_t * s, int * current_index_p){
       }
       if ( o.file_size == (int) o.backend->xfer(READ, aiori_fh, (IOR_size_t *) buf, o.file_size, 0, o.backend_options) ) {
         if(o.verify_read){
-            if(verify_memory_pattern(prevFile * o.dset_count + d, buf, o.file_size, o.random_seed, readRank, o.dataPacketType) == 0){
+            if(verify_memory_pattern(prevFile * o.dset_count + d, buf, o.file_size, o.random_seed, readRank, o.dataPacketType, o.gpu_memory_flags) == 0){
               s->obj_read.suc++;
             }else{
               s->obj_read.err++;
@@ -694,8 +694,8 @@ void run_benchmark(phase_stat_t * s, int * current_index_p){
       op_timer = GetTimeStamp();
       aiori_fh = o.backend->create(obj_name, IOR_WRONLY | IOR_CREAT, o.backend_options);
       if (NULL != aiori_fh){
-        generate_memory_pattern(buf, o.file_size, o.random_seed, writeRank, o.dataPacketType);
-        update_write_memory_pattern(newFileIndex * o.dset_count + d, buf, o.file_size, o.random_seed, writeRank, o.dataPacketType);
+        generate_memory_pattern(buf, o.file_size, o.random_seed, writeRank, o.dataPacketType, o.gpu_memory_flags);
+        update_write_memory_pattern(newFileIndex * o.dset_count + d, buf, o.file_size, o.random_seed, writeRank, o.dataPacketType, o.gpu_memory_flags);
 
         if ( o.file_size == (int) o.backend->xfer(WRITE, aiori_fh, (IOR_size_t *) buf, o.file_size, 0, o.backend_options)) {
           s->obj_create.suc++;

src/mdtest.c

@@ -401,7 +401,7 @@ static void create_file (const char *path, uint64_t itemNum) {
         VERBOSE(3,5,"create_remove_items_helper: write..." );
 
         o.hints.fsyncPerWrite = o.sync_file;
-        update_write_memory_pattern(itemNum, o.write_buffer, o.write_bytes, o.random_buffer_offset, rank, o.dataPacketType);
+        update_write_memory_pattern(itemNum, o.write_buffer, o.write_bytes, o.random_buffer_offset, rank, o.dataPacketType, o.gpu_memory_flags);
 
         if ( o.write_bytes != (size_t) o.backend->xfer(WRITE, aiori_fh, (IOR_size_t *) o.write_buffer, o.write_bytes, 0, o.backend_options)) {
             WARNF("unable to write file %s", curr_item);
@@ -412,7 +412,7 @@ static void create_file (const char *path, uint64_t itemNum) {
             if (o.write_bytes != (size_t) o.backend->xfer(READ, aiori_fh, (IOR_size_t *) o.write_buffer, o.write_bytes, 0, o.backend_options)) {
                 WARNF("unable to verify write (read/back) file %s", curr_item);
             }
-            int error = verify_memory_pattern(itemNum, o.write_buffer, o.write_bytes, o.random_buffer_offset, rank, o.dataPacketType);
+            int error = verify_memory_pattern(itemNum, o.write_buffer, o.write_bytes, o.random_buffer_offset, rank, o.dataPacketType, o.gpu_memory_flags);
             o.verification_error += error;
             if(error){
                 VERBOSE(1,1,"verification error in file: %s", curr_item);
@@ -651,7 +651,7 @@ void mdtest_read(int random, int dirs, const long dir_iter, char *path) {
     /* allocate read buffer */
     if (o.read_bytes > 0) {
         read_buffer = aligned_buffer_alloc(o.read_bytes, o.gpu_memory_flags);
-        memset(read_buffer, -1, o.read_bytes);
+        invalidate_buffer_pattern(read_buffer, o.read_bytes, o.gpu_memory_flags);
     }
 
     uint64_t stop_items = o.items;
@@ -730,7 +730,7 @@ void mdtest_read(int random, int dirs, const long dir_iter, char *path) {
 
         /* read file */
         if (o.read_bytes > 0) {
-            read_buffer[0] = 42;
+            invalidate_buffer_pattern(read_buffer, o.read_bytes, o.gpu_memory_flags);
             if (o.read_bytes != (size_t) o.backend->xfer(READ, aiori_fh, (IOR_size_t *) read_buffer, o.read_bytes, 0, o.backend_options)) {
                 WARNF("unable to read file %s", item);
                 o.verification_error += 1;
@@ -741,7 +741,7 @@ void mdtest_read(int random, int dirs, const long dir_iter, char *path) {
               if (o.shared_file) {
                 pretend_rank = rank;
               }
-              int error = verify_memory_pattern(item_num, read_buffer, o.read_bytes, o.random_buffer_offset, pretend_rank, o.dataPacketType);
+              int error = verify_memory_pattern(item_num, read_buffer, o.read_bytes, o.random_buffer_offset, pretend_rank, o.dataPacketType, o.gpu_memory_flags);
               o.verification_error += error;
               if(error){
                 VERBOSE(1,1,"verification error in file: %s", item);
@@ -2442,7 +2442,7 @@ mdtest_results_t * mdtest_run(int argc, char **argv, MPI_Comm world_com, FILE *
     /* allocate and initialize write buffer with # */
     if (o.write_bytes > 0) {
         o.write_buffer = aligned_buffer_alloc(o.write_bytes, o.gpu_memory_flags);
-        generate_memory_pattern(o.write_buffer, o.write_bytes, o.random_buffer_offset, rank, o.dataPacketType);
+        generate_memory_pattern(o.write_buffer, o.write_bytes, o.random_buffer_offset, rank, o.dataPacketType, o.gpu_memory_flags);
     }
 
     /* setup directory path to work in */

src/parse_options.c

@@ -66,13 +66,7 @@ static void CheckRunSettings(IOR_test_t *tests)
                 }
 
                 if(params->gpuDirect){
-                  if(params->gpuMemoryFlags == IOR_MEMORY_TYPE_GPU_MANAGED){
-                    ERR("GPUDirect cannot be used with managed memory");
-                  }
                   params->gpuMemoryFlags = IOR_MEMORY_TYPE_GPU_DEVICE_ONLY;
-                  if(params->checkRead || params->checkWrite){
-                    ERR("GPUDirect data cannot yet be checked");
-                  }
                 }
         }
 }

src/utilities-gpu.cu

@@ -0,0 +1,51 @@
+/*
+  This file contains CUDA code for creating and checking memory patterns on the device.
+*/
+#include <cuda_runtime.h>
+
+#ifdef HAVE_CONFIG_H
+#include "config.h"
+#endif
+
+#include <stdint.h>
+
+#include "iordef.h"
+
+#define RANDALGO_GOLDEN_RATIO_PRIME        0x9e37fffffffc0001UL
+
+__global__ 
+void cu_generate_memory_timestamp(uint64_t * buf, size_t length, int rand_seed, uint64_t pretendRank){
+  size_t pos = blockIdx.x * blockDim.x + threadIdx.x;
+  if(pos < length){
+    buf[pos] = pretendRank | rand_seed + pos;
+  }
+}
+
+__global__ 
+void cu_verify_memory_timestamp(uint64_t item, uint64_t * buf, size_t length, int rand_seed, uint64_t pretendRank, int * errors){
+
+}
+
+extern "C" void generate_memory_pattern_gpu(char * buf, size_t bytes, int rand_seed, int pretendRank, ior_dataPacketType_e dataPacketType){    
+  size_t blocks = (bytes+2047)/2048;
+  size_t threads = 256;
+  if(dataPacketType == DATA_TIMESTAMP){    
+    cu_generate_memory_timestamp<<<blocks, threads>>>((uint64_t*) buf, bytes/sizeof(uint64_t), rand_seed, ((uint64_t) pretendRank) << 32);
+  }
+}
+
+extern "C" void update_write_memory_pattern_gpu(uint64_t item, char * buf, size_t bytes, int rand_seed, int rank, ior_dataPacketType_e dataPacketType){
+  // nothing to do for dataPacketType == DATA_TIMESTAMP, i.e., won't be called for this parameter
+  size_t blocks = (bytes+2047)/2048;
+  size_t threads = 256;
+}
+
+extern "C" int verify_memory_pattern_gpu(uint64_t item, char * buffer, size_t bytes, int rand_seed, int pretendRank, ior_dataPacketType_e dataPacketType){
+  int errors = 0;
+  size_t blocks = (bytes+2047)/2048;
+  size_t threads = 256;  
+  if(dataPacketType == DATA_TIMESTAMP){
+    cu_verify_memory_timestamp<<<blocks, threads>>>(item, (uint64_t*) buffer, bytes/sizeof(uint64_t), rand_seed, pretendRank, & errors);
+  }
+  return errors;
+}

src/utilities.c

@@ -91,10 +91,17 @@ enum OutputFormat_t outputFormat;
  * @param pretendRank unique identifier for this process
  * @param dataPacketType identifier to designate pattern to fill buffer
  */
-void update_write_memory_pattern(uint64_t item, char * buf, size_t bytes, int rand_seed, int pretendRank, ior_dataPacketType_e dataPacketType){
+void update_write_memory_pattern(uint64_t item, char * buf, size_t bytes, int rand_seed, int pretendRank, ior_dataPacketType_e dataPacketType, ior_memory_flags type){
   if (dataPacketType == DATA_TIMESTAMP || bytes < 8)
     return;
 
+#ifdef HAVE_GPU_DIRECT
+  if(type == IOR_MEMORY_TYPE_GPU_DEVICE_ONLY){
+    update_write_memory_pattern_gpu(item, buf, bytes, rand_seed,  pretendRank, dataPacketType);
+    return;
+  }
+#endif
+
   size_t size = bytes / sizeof(uint64_t);
   uint64_t * buffi = (uint64_t*) buf;
 
@@ -127,7 +134,13 @@ void update_write_memory_pattern(uint64_t item, char * buf, size_t bytes, int ra
  * @param pretendRank unique identifier for this process
  * @param dataPacketType identifier to designate pattern to fill buffer
  */
-void generate_memory_pattern(char * buf, size_t bytes, int rand_seed, int pretendRank, ior_dataPacketType_e dataPacketType){
+void generate_memory_pattern(char * buf, size_t bytes, int rand_seed, int pretendRank, ior_dataPacketType_e dataPacketType, ior_memory_flags type){
+#ifdef HAVE_GPU_DIRECT
+  if(type == IOR_MEMORY_TYPE_GPU_DEVICE_ONLY){
+    generate_memory_pattern_gpu(buf, bytes, rand_seed,  pretendRank, dataPacketType);
+    return;
+  }
+#endif
   uint64_t * buffi = (uint64_t*) buf;
   // first half of 64 bits use the rank
   const size_t size = bytes / 8;
@@ -156,8 +169,24 @@ void generate_memory_pattern(char * buf, size_t bytes, int rand_seed, int preten
   }
 }
 
-int verify_memory_pattern(uint64_t item, char * buffer, size_t bytes, int rand_seed, int pretendRank, ior_dataPacketType_e dataPacketType){
+void invalidate_buffer_pattern(char * buffer, size_t bytes, ior_memory_flags type){
+  if(type == IOR_MEMORY_TYPE_GPU_DEVICE_ONLY){
+#ifdef HAVE_GPU_DIRECT
+    cudaMemset(buffer, 0x42, bytes > 512 ? 512 : bytes);
+#endif
+  }else{
+    buffer[0] = ~buffer[0]; // changes the buffer, no memset to reduce the memory pressure
+  }
+}
+
+int verify_memory_pattern(uint64_t item, char * buffer, size_t bytes, int rand_seed, int pretendRank, ior_dataPacketType_e dataPacketType, ior_memory_flags type){  
   int error = 0;
+#ifdef HAVE_GPU_DIRECT
+  if(type == IOR_MEMORY_TYPE_GPU_DEVICE_ONLY){
+    error = verify_memory_pattern_gpu(item, buffer, bytes, rand_seed, pretendRank, dataPacketType);
+    return error;
+  }
+#endif
   // always read all data to ensure that performance numbers stay the same
   uint64_t * buffi = (uint64_t*) buffer;