 typedef struct
 {
   float4 m[4];
-} float4x4;
-
-__constant__ float4x4 c_invViewMatrix;  // inverse view matrix
+} matrix4x4;
+
+__constant__ matrix4x4 c_invViewMatrix;  // inverse view matrix
 
 struct Ray
 {

  float3 d;
};

__device__ bool intersectBox(Ray r, float3 boxmin, float3 boxmax, float *tnear, float *tfar)
{
    // compute intersection of ray with all six bbox planes
    float3 invR = make_float3(1.0f, 1.0f, 1.0f) / r.d;

  cudaMemcpyToSymbol(c_invViewMatrix, viewM, sizeof(float4) * 4);
  delete[] viewM;

  float3 volSize;

  if ((vd->vox[0] >= vd->vox[1]) && (vd->vox[0] >= vd->vox[2]))
  {
    const float inv = 1.0f / vd->vox[0];
    volSize = make_float3(1.0f, vd->vox[1] * inv, vd->vox[2] * inv);
  }
  else if ((vd->vox[1] >= vd->vox[0]) && (vd->vox[1] >= vd->vox[2]))
  {
    const float inv = 1.0f / vd->vox[1];
    volSize = make_float3(vd->vox[0] * inv, 1.0f, vd->vox[2] * inv);
  }
  else
  {
    const float inv = 1.0f / vd->vox[2];
    volSize = make_float3(vd->vox[0] * inv, vd->vox[1] * inv, 1.0f);
  }

  render<true><<<gridSize, blockSize>>>(d_output, width, height, 2.0f / (float)numSlices, volSize);
  cudaGLUnmapBufferObject(pbo);

  glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, pbo);
