 
   cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float4>();
 
-  cudaFreeArray(d_randArray);
-  cudaMallocArray(&d_randArray, &channelDesc, NUM_RAND_VECS, 1);
-  cudaMemcpyToArray(d_randArray, 0, 0, randVecs, NUM_RAND_VECS * sizeof(float4), cudaMemcpyHostToDevice);
+  vvCuda::checkError(&ok, cudaFreeArray(d_randArray),
+                     "vvRayRend::initRandTexture()");
+  vvCuda::checkError(&ok, cudaMallocArray(&d_randArray, &channelDesc, NUM_RAND_VECS, 1),
+                     "vvRayRend::initRandTexture()");
+  vvCuda::checkError(&ok, cudaMemcpyToArray(d_randArray, 0, 0, randVecs, NUM_RAND_VECS * sizeof(float4),
+                                            cudaMemcpyHostToDevice), "vvRayRend::initRandTexture()");
 
   randTexture.filterMode = cudaFilterModeLinear;
   randTexture.addressMode[0] = cudaAddressModeClamp;
 
-  cudaBindTextureToArray(randTexture, d_randArray, channelDesc);
+  vvCuda::checkError(&ok, cudaBindTextureToArray(randTexture, d_randArray, channelDesc),
+                     "vvRayRend::initRandTexture()");
 
   delete[] randVecs;
 }
 
 void vvRayRend::initSpaceSkippingTexture()
 {
+  bool ok;
+
   cudaExtent numBricks = make_cudaExtent(h_numCells[0], h_numCells[1], h_numCells[2]);
 
   cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<bool>();
 
-  cudaFreeArray(d_spaceSkippingArray);
-  cudaMalloc3DArray(&d_spaceSkippingArray, &channelDesc, numBricks);
+  vvCuda::checkError(&ok, cudaFreeArray(d_spaceSkippingArray),
+                     "vvRayRend::initSpaceSkippingTexture()");
+  vvCuda::checkError(&ok, cudaMalloc3DArray(&d_spaceSkippingArray, &channelDesc, numBricks),
+                     "vvRayRend::initSpaceSkippingTexture()");
 
   cudaMemcpy3DParms copyParams = { 0 };
 

        }
        volTexture8.addressMode[0] = cudaAddressModeClamp;
        volTexture8.addressMode[1] = cudaAddressModeClamp;
        cudaBindTextureToArray(volTexture8, d_volumeArrays[0], _channelDesc);
    }
    else if (vd->bpc == 2)
    {


  cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float4>();

  vvCuda::checkError(&ok, cudaFreeArray(d_transferFuncArray),
                     "vvRayRend::updateTransferFunction()");
  vvCuda::checkError(&ok, cudaMallocArray(&d_transferFuncArray, &channelDesc, lutEntries, 1),
                     "vvRayRend::updateTransferFunction()");
  vvCuda::checkError(&ok, cudaMemcpyToArray(d_transferFuncArray, 0, 0, _rgbaTF, lutEntries * 4 * sizeof(float),
                                            cudaMemcpyHostToDevice),
                     "vvRayRend::updateTransferFunction()");


  tfTexture.filterMode = cudaFilterModeLinear;
  tfTexture.normalized = true;    // access with normalized texture coordinates
  tfTexture.addressMode[0] = cudaAddressModeClamp;   // wrap texture coordinates

  vvCuda::checkError(&ok, cudaBindTextureToArray(tfTexture, d_transferFuncArray, channelDesc),
                     "vvRayRend::updateTransferFunction()");
}

void vvRayRend::compositeVolume(int, int)
