             d.w += w;
 
             // store into shmem
-            *pix = d;
+            *pix = d;//make_uchar4(*v, ix&0xff, ix&0xff, 255);
         }
     }
 

        return;

    // initialise intermediate image line
    __shared__ extern uchar4 imgLine[];
    for (int ix=threadIdx.x; ix<width; ix+=blockDim.x)
    {
        imgLine[ix] = make_uchar4(0,0,0,0);

        // the voxel row of the current slice corresponding to this image line
        const uchar *voxLine = voxels +BPV * (MUL24(slice+1,MUL24(c_vox[principal+1],c_vox[principal+0]))
                + MUL24((iPosY-line-1),c_vox[principal+0]));

#ifdef SHMLOAD
        for (int ix=threadIdx.x; ix<c_vox[principal+0]; ix+=blockDim.x)
        {
            //((uchar4 *)voxel)[ix] = ((uchar4 *)voxels)[ix];
            voxel[ix] = voxLine[ix];
        }
        __syncthreads();
#endif

        // Traverse intermediate image pixels which correspond to the current slice.
        // 1 is subtracted from each loop counter to remain inside of the volume boundaries:
