In CUDA, Supercomputing for the Masses: Part 16 of this article series, I discussed the CUDA 3.0 release. CUDA 3.0 is a major revision number increment release that adds enhancements valuable to all CUDA developers to make day-to-day development tasks easier, less error prone, and more consistent.

As mentioned in the previous article, expanded consistent coverage appears to have been the thrust behind this major revision number release as it fills in several previous gaps and must-have capabilities. In a nutshell, in this article I discuss runtime and driver API compatibility, the new graphics interoperability API, C++ inheritance plus expanded functionality in CUBLAS and CUFFT. Examples are provided that demonstrate:

• Consistency and interoperability between the CUDA runtime and driver codes through two simple examples that call a runtime kernel and a CUBLAS routine.
• A C++ example that:
• Uses both versions of the OpenGL interoperability APIs in drawing a Mandelbrot set. The C++ source code concisely recreates the simplePBO example from Part 15 of this article series using C++ classes.
• Demonstrates C++ inheritance by deriving a new class from our C++ Mandelbrot example that uses programmable shaders created with Cg. (Please look to the extensive NVIDIA Cg homepage for more information). While Cg compatibility is not new with the CUDA 3.0 release, mixing Cg shaders with CUDA can open up a vast collection of Cg libraries and existing software!

Be aware that the latest NVIDIA driver must be installed to use the CUDA 3.0 toolkit. As always, the latest released driver can be downloaded from CUDA ZONE and installed for a number of systems. Beta drivers and software can be downloaded from nvdevelopers but registration is required. Ubuntu users might wish to follow one of the many available guides, such as the one at Web Upd8, to see how to install the latest released or beta drivers via the Ubuntu tools.

Mixing CUDA Runtime and Driver API Codes

Previous articles in this series have focused on teaching the CUDA with the runtime API (e.g. those methods that start with "cuda" as opposed to "cu") because it is fairly intuitive and not too verbose. Many developers prefer to utilize the driver API because they have more control and can make better use of existing code bases. Now programmers can utilized the best characteristics of both APIs.

The following is the source code for a driver mode CUDA program that calls a kernel via the runtime API. Please put this into a file called vectorAddDrv.cu:

/*
 * Driver APIC code that calls a runtime kernel
 * Vector addition: C = A + B.
 */

// Includes
#include <stdio.h>
#include <cuda.h>
#include <cutil_inline.h>

// Variables
CUdevice cuDevice;
CUcontext cuContext;
CUmodule cuModule;
CUfunction vecAdd;
float* h_A;
float* h_B;
float* h_C;
CUdeviceptr d_A;
CUdeviceptr d_B;
CUdeviceptr d_C;

// Functions

__global__ void kernel(float* d_a, float* d_b, float* d_c, int n)
{
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  
  if(idx < n)
    d_c[idx] = d_a[idx] + d_b[idx];
}

// Allocates an array with random float entries.
void RandomInit(float* data, int n)
{
  for (int i = 0; i < n; ++i)
    data[i] = rand() / (float)RAND_MAX;
}

void errorExit()
{
  printf("Error exit!\n");
  exit(1);
}

// Host code
int main(int argc, char** argv)
{
  int N = 50000;
  unsigned int size = N * sizeof(float);
  CUresult error;
  
  printf("Vector Addition (Driver API)\n");
  // Initialize
  error = cuInit(0);
  if (error != CUDA_SUCCESS) errorExit();
  
  // Get number of devices supporting CUDA
  int deviceCount = 0;
  error = cuDeviceGetCount(&deviceCount);
  if (error != CUDA_SUCCESS) errorExit();
  if (deviceCount == 0) {
    printf("There is no device supporting CUDA.\n");
    exit(1);
  }
  
  // Get handle for device 0
  error = cuDeviceGet(&cuDevice, 0);
  if (error != CUDA_SUCCESS) errorExit();
  
  // Create context
  error = cuCtxCreate(&cuContext, 0, cuDevice);
  if (error != CUDA_SUCCESS) errorExit();
  
  // Allocate input vectors h_A and h_B in host memory
  h_A = (float*)malloc(size);
  if (h_A == 0) errorExit();
  h_B = (float*)malloc(size);
  if (h_B == 0) errorExit();
  h_C = (float*)malloc(size);
  if (h_C == 0) errorExit();
  
  // Initialize input vectors
  RandomInit(h_A, N);
  RandomInit(h_B, N);
  
  // Allocate vectors in device memory
  error = cuMemAlloc(&d_A, size);
  if (error != CUDA_SUCCESS) errorExit();
  error = cuMemAlloc(&d_B, size);
  if (error != CUDA_SUCCESS) errorExit();
  error = cuMemAlloc(&d_C, size);
  if (error != CUDA_SUCCESS) errorExit();
  
  // Copy vectors from host memory to device memory
  error = cuMemcpyHtoD(d_A, h_A, size);
  if (error != CUDA_SUCCESS) errorExit();
  error = cuMemcpyHtoD(d_B, h_B, size);
  if (error != CUDA_SUCCESS) errorExit();
  
  // Invoke kernel (Runtime API)
  int nThreadsPerBlk=128;
  int nBlks = (N/nThreadsPerBlk) + (((N%nThreadsPerBlk)>0)?1:0);
  kernel<<<nBlks,nThreadsPerBlk>>>((float*)d_A,(float*) d_B,(float*) d_C, N);
  
  // Copy result from device memory to host memory
  // h_C contains the result in host memory
  error = cuMemcpyDtoH(h_C, d_C, size);
  if (error != CUDA_SUCCESS) errorExit();
  
  // Verify result
  int i;
  for (i = 0; i < N; ++i) {
    float sum = h_A[i] + h_B[i];
    if (fabs(h_C[i] - sum) > 1e-7f) {
      printf("Mistake index %d %g %g\n",i,h_C[i],sum);
      break;
    }
  }
  printf("Test %s \n", (i == N) ? "PASSED" : "FAILED");
  return(0);
}

A detailed discussion of the driver API is beyond the scope of this article, as I'm focusing on interoperability and the 3.0 release. Even so, much of this code should look familiar to runtime API developers as only GPU setup, memory allocation and data movement is utilized. Many of these calls are similar to the runtime API calls so it should be easy to follow the source code. The specific point made with this example is that the following runtime CUDA call to kernel() works in the 3.0 release:

  kernel<<<nBlks,nThreadsPerBlk>>>((float*)d_A,(float*) d_B,(float*) d_C, N);

Use nvcc to build the executable with debugging enabled as shown below for Linux:

SDK_PATH=$HOME/NVIDIA_GPU_Computing_SDK/C
INC=$SDK_PATH/common/inc
LIB=$SDK_PATH/lib
nvcc -g -G -I$INC -L$LIB vectorAddDrv.cu -lcutil_x86_64 -lcuda -o vectorAddDrv

Now start the program with the command cuda-gdb:

$ cuda-gdb vectorAddDrv
NVIDIA (R) CUDA Debugger
BETA release
Portions Copyright (C) 2008,2009 NVIDIA Corporation
GNU gdb 6.6
Copyright (C) 2006 Free Software Foundation, Inc.
GDB is free software, covered by the GNU General Public License, and you are
welcome to change it and/or distribute copies of it under certain conditions.
Type "show copying" to see the conditions.
There is absolutely no warranty for GDB.  Type "show warranty" for details.
This GDB was configured as "x86_64-unknown-linux-gnu"...
Using host libthread_db library "/lib/libthread_db.so.1".	

Enable memory checking and run the program with the following two commands:

(cuda-gdb) set cuda memcheck on
(cuda-gdb) r

The following illustrative output shows that the code works and apparently used memory correctly. Further the results passed a validation test on the host as indicated by the "Test PASSED" message.

Starting program: PATH_TO_EXECUTABLE/vectorAddDrv 
[Thread debugging using libthread_db enabled]
[New process 31619]
Vector Addition (Driver API)
[New Thread 139794699740944 (LWP 31619)]
Test PASSED 

Program exited normally.

Now set a breakpoint at line 109, which is just after the result vector is moved to the host computer and run the vectorAddDrv again:

.

(cuda-gdb) b 109
Breakpoint 1 at 0x40183e: file vectorAdd.cu, line 109.

(cuda-gdb) r
Starting program: /home/rmfarber/DDJ/Part17/vecAdd/vectorAdd 
[Thread debugging using libthread_db enabled]
[New process 31623]
Vector Addition (Driver API)
[New Thread 140441416316688 (LWP 31623)]
[Switching to Thread 140441416316688 (LWP 31623)]

Breakpoint 1, main (argc=1, argv=0x7fff40b55248) at vectorAdd.cu:109
109	  if (error != CUDA_SUCCESS) errorExit();

Use the CUDA-GDB set command to change index three of the h_C vector on the host to be 1000-times larger:

(cuda-gdb) p h_C[3]
$1 = 1.30907393
(cuda-gdb) set h_C[3] = 1000. * h_C[3]

The program is then allowed to continue. The resulting output shows that the comparison code does find the error caused by our manually modifying the h_C vector:

(cuda-gdb) c
Continuing.
Mistake index 3 1309.07 1.30907
Test FAILED 

Program exited normally.
(cuda-gdb) quit

This example demonstrates that CUDA-GDB in the 3.0 release works with driver API programs. It also shows how straightforward it now is to mix driver and runtime API codes.

Similarly, the following example, blasAddDrv.cu, demonstrates that it is now possible with CUDA 3.0 to call the CUBLAS library routines. In this case, the previous example code for vectorAddDrv.cu was adapted to call the saxpy() routine:

  cublasSaxpy(N, 1.0f,(float*) d_A, 1,(float*) d_B, 1);

The following is the complete source code for vectorAddDrv.cu. Again, I won't discuss the details of the driver API. See the NVIDIA documentation for more information.

/*
 * Vector addition using SAXPY from a driver API example
 */

// Includes
#include <stdio.h>
#include <cuda.h>
#include <cutil_inline.h>
#include <cublas.h>


// Variables
CUdevice cuDevice;
CUcontext cuContext;
CUmodule cuModule;
CUfunction vecAdd;
float* h_A;
float* h_B;
float* h_C;
CUdeviceptr d_A;
CUdeviceptr d_B;

// Functions

// Allocates an array with random float entries.
void RandomInit(float* data, int n)
{
  for (int i = 0; i < n; ++i)
    data[i] = rand() / (float)RAND_MAX;
}

void errorExit()
{
  printf("Error exit!\n");
}

// Host code
int main(int argc, char** argv)
{
  int N = 50000;
  unsigned int size = N * sizeof(float);
  CUresult error;
  int status;
  
  printf("Vector Addition (BLAS and Driver APIs)\n");
  // Initialize
  error = cuInit(0);
  if (error != CUDA_SUCCESS) errorExit();
  
  // Get number of devices supporting CUDA
  int deviceCount = 0;
  error = cuDeviceGetCount(&deviceCount);
  if (error != CUDA_SUCCESS) errorExit();
  if (deviceCount == 0) {
    printf("There is no device supporting CUDA.\n");
    exit(1);
  }
  
  // Get handle for device 0
  error = cuDeviceGet(&cuDevice, 0);
  if (error != CUDA_SUCCESS) errorExit();
  
  // Create context
  error = cuCtxCreate(&cuContext, 0, cuDevice);
  if (error != CUDA_SUCCESS) errorExit();
  
  // Allocate input vectors h_A and h_B in host memory
  h_A = (float*)malloc(size);
  if (h_A == 0) errorExit();
  h_B = (float*)malloc(size);
  if (h_B == 0) errorExit();
  h_C = (float*)malloc(size);
  if (h_C == 0) errorExit();
  
  // Initialize input vectors
  RandomInit(h_A, N);
  RandomInit(h_B, N);
  
  // Allocate vectors in device memory
  error = cuMemAlloc(&d_A, size);
  if (error != CUDA_SUCCESS) errorExit();
  error = cuMemAlloc(&d_B, size);
  if (error != CUDA_SUCCESS) errorExit();
  
  // Copy vectors from host memory to device memory
  error = cuMemcpyHtoD(d_A, h_A, size);
  if (error != CUDA_SUCCESS) errorExit();
  error = cuMemcpyHtoD(d_B, h_B, size);
  if (error != CUDA_SUCCESS) errorExit();
  
  // Invoke kernel
  status = cublasInit();
  if (status != CUBLAS_STATUS_SUCCESS) {
    fprintf (stderr, "!!!! CUBLAS initialization error\n");
    return EXIT_FAILURE;
  }
  
  // Call CUBLAS from a driver API code
  cublasSaxpy(N, 1.0f,(float*) d_A, 1,(float*) d_B, 1);
  
  // Copy result from device memory to host memory
  // h_C contains the result in host memory
  error = cuMemcpyDtoH(h_C, d_B, size);
  if (error != CUDA_SUCCESS) errorExit();
  
  // Verify result
  int i;
  for (i = 0; i < N; ++i) {
    float sum = h_A[i] + h_B[i];
    if (fabs(h_C[i] - sum) > 1e-7f)
      break;
  }
  printf("Test %s \n", (i == N) ? "PASSED" : "FAILED");
  return(0);
}

Use nvcc to build the executable. The following simple script builds the executable on Linux. Please note that this program requires the CUBLAS library. Linking is specified with the -lcublas command-line option.

SDK_PATH=$HOME/NVIDIA_GPU_Computing_SDK/C
INC=$SDK_PATH/common/inc
LIB=$SDK_PATH/lib
nvcc -I$INC -L$LIB blasAddDrv.cu -lcutil_x86_64 -lcuda -lcublas -o blasAddDrv

Running the program demonstrates that it does indeed work correctly as indicated by the "Test PASSED" message:

$ ./blasAddDrv
Vector Addition (BLAS and Driver APIs)
Test PASSED

Unquestionably, the ability to mix and debug driver and runtime API codes and libraries is valuable. For many, this expanded capability alone will make the CUDA 3.0 release an obvious download choice. As the CUDA library and code base expands, this transparent interoperability will continue to pay dividends in ease of use -- although it is likely that in the future most developers will utilize this capability without thought or awareness of the ease in which it occurs.