     src.y *= src.w;
     src.z *= src.w;
 
-    if (frontToBack)
+    if (frontToBack && (mipMode == 0))
     {
       dst = dst + src * (1.0f - dst.w);
     }

  return make_float4(tmp.x, tmp.y, tmp.z, classification.w);
}

template<bool frontToBack, bool lighting>
__global__ void render(uint *d_output, const uint width, const uint height, const float dist,
                       const float3 volSizeHalf, const float3 L, const float3 H)
{
  const int maxSteps = INT_MAX;
  const float tstep = dist;

    tnear = 0.0f;
  }

  // Calc hits with clip sphere.
  float tsnear;
  float tsfar;
  if (clipSphere)
  {
    intersectSphere(ray, sphereCenter, sphereRadius, &tsnear, &tsfar);
  }

  // Calc hits with clip plane.
  float tpnear;
  float nddot;
  if (clipPlane)
  {
    intersectPlane(ray, planeNormal, planeDist, &nddot, &tpnear);
  }

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

  float maxIntensity = 0.0f;
  float minIntensity = FLT_MAX;

  for (int i=0; i<maxSteps; ++i)
  {
    // Test for clipping.
    if ((
        (clipSphere && (t >= tsnear) && (t <= tsfar)) // Sphere.
     || (clipPlane && (((t <= tpnear) && (nddot >= 0.0f))
         || ((t >= tpnear) && (nddot < 0.0f)))) // Plane.
    ))
    {
      t += tstep;
      if (t > tfar)
      {
        break;
      }
      pos += step;
      continue;
    }

    const float3 texCoord = calcTexCoord(pos, volSizeHalf);
    const float sample = volume(texCoord);

    // Post-classification transfer-function lookup.
    float4 src;

    if (mipMode == 0)
    {
      src = tex1D(tfTexture, sample);
    }
    else if (mipMode == 1 && (sample > maxIntensity))
    {
      dst = tex1D(tfTexture, sample);
      maxIntensity = sample;
    }
    else if (mipMode == 2 && (sample  < minIntensity))
    {
      dst = tex1D(tfTexture, sample);
      minIntensity = sample;
    }

    // Local illumination.
    if (lighting && (src.w > 0.1))
