     for (int ix=threadIdx.x; ix<width; ix+=blockDim.x)
     {
         uchar4 *dest = img + line*width+ix;
+#ifdef FLOATIMG
+        *dest = make_uchar4(imgLine[ix].x * 255.f,
+                imgLine[ix].y * 255.f,
+                imgLine[ix].z * 255.f,
+                imgLine[ix].w * 255.f);
+#else
         *dest = imgLine[ix];
+#endif
     }
 #endif
 }

   }
#endif

   // do the computation on the device
#ifdef PITCHED
#define compositeNearest(sliceStep, principal) \
   compositeSlicesNearest<Scalar, 1, sliceStep, principal> <<<to-from, 128, intImg->width*vvSoftImg::PIXEL_SIZE+ vd->vox[principal]*vd->getBPV()*sizeof(Scalar)>>>( \
         d_img, intImg->width, intImg->height, \
         d_voxptr[principal], \
         firstSlice, lastSlice, \
         from, to)
#else
#define compositeNearest(sliceStep, principal) \
   compositeSlicesNearest<Scalar, 1, sliceStep, principal> <<<to-from, 128, intImg->width*vvSoftImg::PIXEL_SIZE+vd->vox[principal]*vd->getBPV()*sizeof(Scalar)>>>( \
         d_img, intImg->width, intImg->height, \
         (Scalar *)(d_voxels+sizeof(Scalar)*vd->getBPV()*principal*(vd->vox[0]*vd->vox[1]*vd->vox[2])), \
         firstSlice, lastSlice, \

//#define NOOP
//#define NOLOAD
//#define NODISPLAY
//#define SHMLOAD
//#define ARRAY
#define PITCHED
#define FLOATDATA
//#define CONSTLOAD
#define THREADPERVOXEL
//#define CONSTDATA
#define FLOATIMG

const int Repetitions = 1;

