Initial checkin, new error handling macros

chLib/chError.h

@@ -3,11 +3,14 @@
  * chError.h
  *
  * Error handling for CUDA:
- *     CUDA_CHECK() and CUDART_CHECK() macros implement
+ *     cu() and cuda() macros implement goto-based error
- *         goto-based error handling, and
+ *         error handling *, and
  *     chGetErrorString() maps a driver API error to a string.
  *
- * Copyright (c) 2011-2012, Archaea Software, LLC.
+ * * The more-concise formulation of these macros is due to
+ *   Allan MacKinnon.
+ *
+ * Copyright (c) 2011-2016, Archaea Software, LLC.
  * All rights reserved.
  *
  * Redistribution and use in source and binary forms, with or without
@@ -134,6 +137,26 @@ chGetErrorString( CUresult status )
         } \
     } while (0);
 
+#define cuda( fn ) do { \
+        (status) =  (cuda##fn); \
+        if ( cudaSuccess != (status) ) { \
+            fprintf( stderr, "CUDA Runtime Failure (line %d of file %s):\n\t" \
+                "%s returned 0x%x (%s)\n", \
+                __LINE__, __FILE__, #fn, status, chGetErrorString(status) ); \
+            goto Error; \
+        } \
+    } while (0);
+
+#define cu( fn ) do { \
+        (status) =  (cu##fn); \
+        if ( CUDA_SUCCESS != (status) ) { \
+            fprintf( stderr, "CUDA Runtime Failure (line %d of file %s):\n\t%s "\
+                "returned 0x%x (%s)\n", \
+                __LINE__, __FILE__, #fn, status, chGetErrorString(status) ); \
+            goto Error; \
+        } \
+    } while (0);
+
 #define CUDA_CHECK( fn ) do { \
         (status) =  (fn); \
         if ( CUDA_SUCCESS != (status) ) { \
@@ -146,6 +169,7 @@ chGetErrorString( CUresult status )
 
 #else
 
+
 #define CUDART_CHECK( fn ) do { \
     status = (fn); \
     if ( cudaSuccess != (status) ) { \
@@ -153,6 +177,14 @@ chGetErrorString( CUresult status )
 	} \
     } while (0);
 
+#define cuda( fn ) do { \
+    status = (cuda##fn); \
+    if ( cudaSuccess != (status) ) { \
+	    goto Error; \
+	} \
+    } while (0);
+
+
 #define CUDA_CHECK( fn ) do { \
         (status) =  (fn); \
         if ( CUDA_SUCCESS != (status) ) { \
@@ -160,6 +192,13 @@ chGetErrorString( CUresult status )
         } \
     } while (0);
 
+#define cu( fn ) do { \
+        (status) =  (cu##fn); \
+        if ( CUDA_SUCCESS != (status) ) { \
+            goto Error; \
+        } \
+    } while (0);
+
 #endif
 
 #else

concurrency/nullKernelAsync.cu

@@ -54,13 +54,13 @@ usPerLaunch( int cIterations )
     double microseconds, ret;
     chTimerTimestamp start, stop;
 
-    CUDART_CHECK( cudaFree(0) );
+    cuda(Free(0));
 
     chTimerGetTime( &start );
     for ( int i = 0; i < cIterations; i++ ) {
         NullKernel<<<1,1>>>();
     }
-    CUDART_CHECK( cudaThreadSynchronize() );
+    cuda(ThreadSynchronize());
     chTimerGetTime( &stop );
 
     microseconds = 1e6*chTimerElapsedTime( &start, &stop );

driverAPI/chDrv.cpp

@@ -92,14 +92,14 @@ chCUDADevice::Initialize(
     CUdevice device;
     CUcontext ctx = 0;
 
-    CUDA_CHECK( cuDeviceGet( &device, ordinal ) );
+    cu(DeviceGet( &device, ordinal ) );
-    CUDA_CHECK( cuCtxCreate( &ctx, CtxCreateFlags, device ) );
+    cu(CtxCreate( &ctx, CtxCreateFlags, device ) );
     for ( list<string>::iterator it  = moduleList.begin();
                                  it != moduleList.end();
                                  it++ ) {
         CUDA_CHECK( loadModuleFromFile( NULL, *it, numOptions, options, optionValues ) );
     }
-    CUDA_CHECK( cuCtxPopCurrent( &ctx ) );
+    cu(CtxPopCurrent( &ctx ) );
     m_device = device;
     m_context = ctx;
     return CUDA_SUCCESS;
@@ -116,8 +116,8 @@ chCUDAInitialize( list<string>& moduleList )
     int cDevicesInitialized = 0;
     chCUDADevice *newDevice;
     
-    CUDA_CHECK( cuInit( 0 ) );
+    cu(Init( 0 ) );
-    CUDA_CHECK( cuDeviceGetCount( &cDevices ) );
+    cu(DeviceGetCount( &cDevices ) );
     for ( int i = 0; i < cDevices; i++ ) {
         CUdevice device;
         CUcontext ctx = 0;

driverAPI/saxpyDrv.cpp

@@ -55,17 +55,17 @@ TestSAXPY( chCUDADevice *chDevice, size_t N, float alpha )
     float *hostOut = 0;
     float *hostIn = 0;
 
-    CUDA_CHECK( cuCtxPushCurrent( chDevice->context() ) );
+    cu(CtxPushCurrent( chDevice->context() ) );
 
-    CUDA_CHECK( cuMemAlloc( &dptrOut, N*sizeof(float) ) );
+    cu(MemAlloc( &dptrOut, N*sizeof(float) ) );
-    CUDA_CHECK( cuMemsetD32( dptrOut, 0, N ) );
+    cu(MemsetD32( dptrOut, 0, N ) );
-    CUDA_CHECK( cuMemAlloc( &dptrIn, N*sizeof(float) ) );
+    cu(MemAlloc( &dptrIn, N*sizeof(float) ) );
-    CUDA_CHECK( cuMemHostAlloc( (void **) &hostOut, N*sizeof(float), 0 ) );
+    cu(MemHostAlloc( (void **) &hostOut, N*sizeof(float), 0 ) );
-    CUDA_CHECK( cuMemHostAlloc( (void **) &hostIn, N*sizeof(float), 0 ) );
+    cu(MemHostAlloc( (void **) &hostIn, N*sizeof(float), 0 ) );
     for ( size_t i = 0; i < N; i++ ) {
         hostIn[i] = (float) rand() / (float) RAND_MAX;
     }
-    CUDA_CHECK( cuMemcpyHtoDAsync( dptrIn, hostIn, N*sizeof(float ), NULL ) );
+    cu(MemcpyHtoDAsync( dptrIn, hostIn, N*sizeof(float ), NULL ) );
 
     {
         CUmodule moduleSAXPY;
@@ -77,14 +77,14 @@ TestSAXPY( chCUDADevice *chDevice, size_t N, float alpha )
             status = CUDA_ERROR_NOT_FOUND;
             goto Error;
         }
-        CUDA_CHECK( cuModuleGetFunction( &kernelSAXPY, moduleSAXPY, "saxpy" ) );
+        cu(ModuleGetFunction( &kernelSAXPY, moduleSAXPY, "saxpy" ) );
 
-        CUDA_CHECK( cuLaunchKernel( kernelSAXPY, 1500, 1, 1, 512, 1, 1, 0, NULL, params, NULL ) );
+        cu(LaunchKernel( kernelSAXPY, 1500, 1, 1, 512, 1, 1, 0, NULL, params, NULL ) );
 
     }
 
-    CUDA_CHECK( cuMemcpyDtoHAsync( hostOut, dptrOut, N*sizeof(float), NULL ) );
+    cu(MemcpyDtoHAsync( hostOut, dptrOut, N*sizeof(float), NULL ) );
-    CUDA_CHECK( cuCtxSynchronize() );
+    cu(CtxSynchronize() );
     for ( size_t i = 0; i < N; i++ ) {
         if ( fabsf( hostOut[i] - alpha*hostIn[i] ) > 1e-5f ) {
             status = CUDA_ERROR_UNKNOWN;
@@ -117,7 +117,7 @@ main( int argc, char *argv[] )
                                            it++ ) {
         char deviceName[256];
         chCUDADevice *chDevice = *it;
-        CUDA_CHECK( cuDeviceGetName( deviceName, 255, chDevice->device() ) );
+        cu(DeviceGetName( deviceName, 255, chDevice->device() ) );
         printf( "Testing SAXPY on %s (device %d)...", deviceName, chDevice->device() );
         CUDA_CHECK( TestSAXPY( chDevice, 16*1048576, 2.0 ) );
     }