   cudaMemcpyToSymbol(c_invViewMatrix, viewM, sizeof(float4) * 4);
   delete[] viewM;
 
-  render<true><<<gridSize, blockSize>>>(d_output, width, height, 2.0f / (float)numSlices);
+  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);

  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;


  float4 dst = make_float4(0.0f);
  float t = tnear;
  float3 pos = eyeRay.o + eyeRay.d * tnear;
  float3 step = eyeRay.d * tstep;

  for (int i=0; i<maxSteps; ++i)
  {
    const float sample = tex3D(volTexture,
                               (pos.x + volSize.x) / (volSize.x * 2),
                               (pos.y + volSize.y) / (volSize.y * 2),
                               (pos.z + volSize.z) / (volSize.z * 2));

    // Post-classification transfer-function lookup.
    float4 src = tex1D(tfTexture, sample);
