 #endif
 
    bool ok = true;
-#if defined(PITCHED) || defined(VOL3TEX3D)
+#if defined(PITCHED) || defined(VOLTEX3D)
+#ifdef VOLTEX3D
+   for (int i=0; i<VOLTEX3D; ++i)
+#else
    for (int i=0; i<3; ++i)
+#endif
    {
 #ifdef PITCHED
        cudaExtent extent = make_cudaExtent(vd->vox[(i+1)%3]*sizeof(Scalar), vd->vox[(i+2)%3], vd->vox[(i+3)%3]);
        if(!vvCuda::checkError(&ok, cudaMalloc3D(&d_voxptr[i], extent), "cudaMalloc3D vox"))
            break;
 #else
+#if defined(VOLTEX3D) && VOLTEX3D == 1
+       cudaExtent extent = make_cudaExtent(vd->vox[0], vd->vox[1], vd->vox[2]);
+#else
        cudaExtent extent = make_cudaExtent(vd->vox[(i+1)%3], vd->vox[(i+2)%3], vd->vox[(i+3)%3]);
+#endif
        cudaChannelFormatDesc desc = cudaCreateChannelDesc<Scalar>();
        if(!vvCuda::checkError(&ok, cudaMalloc3DArray(&d_voxarr[i], &desc, extent, 0), "cudaMalloc3DArray vox"))
            break;
 #endif
        cudaMemcpy3DParms parms = {0};
+#if defined(VOLTEX3D) && VOLTEX3D == 1
+#ifdef FLOATDATA
+       parms.srcPtr = make_cudaPitchedPtr(fraw[2], sizeof(Scalar)*vd->vox[0], vd->vox[0], vd->vox[1]);
+#else
+       parms.srcPtr = make_cudaPitchedPtr(raw[2], sizeof(Scalar)*vd->vox[0], vd->vox[0], vd->vox[1]);
+#endif
+#else
 #ifdef FLOATDATA
        parms.srcPtr = make_cudaPitchedPtr(fraw[i], sizeof(Scalar)*vd->vox[(i+1)%3], vd->vox[(i+1)%3], vd->vox[(i+2)%3]);
 #else
        parms.srcPtr = make_cudaPitchedPtr(raw[i], sizeof(Scalar)*vd->vox[(i+1)%3], vd->vox[(i+1)%3], vd->vox[(i+2)%3]);
+#endif
 #endif
 
 #ifdef PITCHED

   bool ok = true;
   vvCuda::checkError(&ok, cudaMemcpyToSymbol(c_start, h_start, sizeof(h_start)), "cudaMemcpy start");

#ifdef VOL3TEX3D
   tex_raw.normalized = false;
   tex_raw.filterMode = cudaFilterModePoint;
   tex_raw.addressMode[0] = cudaAddressModeClamp;
   tex_raw.addressMode[1] = cudaAddressModeClamp;
   tex_raw.addressMode[2] = cudaAddressModeClamp;
   cudaChannelFormatDesc desc = cudaCreateChannelDesc<Scalar>();
   cudaBindTextureToArray(tex_raw, d_voxarr[principal], desc);
#endif

#ifndef NODISPLAY

//#define NOLOAD
//#define NODISPLAY
//#define SHMLOAD
#define VOLTEX3D 1 // undef, 1 or 3
//#define PITCHED
#define FLOATDATA
//#define CONSTLOAD
