        parms.extent = make_cudaExtent(vd->vox[(i+1)%3]*sizeof(Scalar), vd->vox[(i+2)%3], vd->vox[(i+3)%3]);
 #else
        parms.dstArray = d_voxarr[i];
+#if defined(VOLTEX3D) && VOLTEX3D == 1
+       parms.extent = make_cudaExtent(vd->vox[0], vd->vox[1], vd->vox[2]);
+#else
        parms.extent = make_cudaExtent(vd->vox[(i+1)%3], vd->vox[(i+2)%3], vd->vox[(i+3)%3]);
+#endif
 #endif
        parms.kind = cudaMemcpyHostToDevice;
        if(!vvCuda::checkError(&ok, cudaMemcpy3D(&parms), "cudaMemcpy3D vox"))

   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
