     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
 }

            d.y += w*c.y;
            d.z += w*c.z;
            d.w += w;

            // store into shmem
            *pix = d;

#ifdef THREADPERVOXEL
            __syncthreads();
#endif

//#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;

