     if (line >= to)
         return;
 
+#ifdef NOSHMEM
+    for (int ix=threadIdx.x; ix<width; ix+=blockDim.x)
+    {
+        Pixel d;
+        initPixel(&d);
+        Scalar sf = 0.f;
+
+        // composite slices for this image line
+        for (int slice=firstSlice; slice!=lastSlice; slice += sliceStep)
+        {
+            if(earlyRayTerm && isOpaque(d))
+                break;
+
+            // compute upper left image corner
+            const int iPosY = c_start[slice].y;
+
+            const int iPosX = c_start[slice].x;
+            const int endX = c_stop[slice].x;
+
+            if(line < iPosY)
+                continue;
+            if(line >= c_stop[slice].y)
+                continue;
+
+            // Traverse intermediate image pixels which correspond to the current slice.
+            // 1 is subtracted from each loop counter to remain inside of the volume boundaries:
+            if(ix<iPosX)
+                continue;
+            if(ix>=endX)
+                continue;
+
+            const int vidx = ix - iPosX;
+
+            if(preInt)
+            {
+                const float sb = volume(vidx, line-iPosY, slice, principal);
+                if(slice != firstSlice)
+                {
+                    const float4 c = tex2D(tex_preint, sf, sb);
+
+                    // blend
+                    const float w = d.w*c.w;
+                    d.x += w*c.x;
+                    d.y += w*c.y;
+                    d.z += w*c.z;
+                    d.w -= w;
+                }
+                sf = sb;
+            }
+            else
+            {
+                const float v = volume(vidx, line-iPosY, slice, principal);
+                const float4 c = classify<float4>(v);
+
+                // blend
+                const float w = d.w*c.w;
+                d.x += w*c.x;
+                d.y += w*c.y;
+                d.z += w*c.z;
+                d.w -= w;
+            }
+        }
+
+        // copy pixel to intermediate image
+        uchar4 *dest = img + line*width+ix;
+        blend(dest, d);
+    }
+#else
     // initialise intermediate image line
     extern __shared__ char smem[];
-    // store intermediate image line in shmem
     Pixel *imgLine = (Pixel *)smem;
     Scalar *sf = (Scalar *)&imgLine[width*preInt];
 
-    // clear image line
     for (int ix=threadIdx.x; ix<width; ix+=blockDim.x)
     {
         initPixel(&imgLine[ix]);

        uchar4 *dest = img + line*width+ix;
        blend(dest, imgLine[ix]);
    }
}
#endif



const int MaxCompositeSlices = MAX_SLICES;

const int nthreads = 128;

//#define NOSHMEM
//#define SHMCLASS
//#define NOOP
//#define NOLOAD
