diff --git a/Data/Array/Accelerate/CUDA.hs b/Data/Array/Accelerate/CUDA.hs
--- a/Data/Array/Accelerate/CUDA.hs
+++ b/Data/Array/Accelerate/CUDA.hs
@@ -13,9 +13,34 @@
 -- Accelerate. Expressions are on-line translated into CUDA code, compiled, and
 -- executed in parallel on the GPU.
 --
--- The accelerate-cuda library is hosted at <https://github.com/tmcdonell/accelerate-cuda>.
--- Comments, bug reports, and patches are always welcome.
+-- The accelerate-cuda library is hosted at: <https://github.com/AccelerateHS/accelerate-cuda>.
+-- Comments, bug reports, and patches, are always welcome.
 --
+--
+-- /NOTES:/
+--
+-- CUDA devices are categorised into different \'compute capabilities\',
+-- indicating what operations are supported by the hardware. For example, double
+-- precision arithmetic is only supported on devices of compute capability 1.3
+-- or higher.
+--
+-- Devices generally perform best when dealing with (tuples of) 32-bit types, so
+-- be cautious when introducing 8-, 16-, or 64-bit elements. Keep in mind the
+-- size of 'Int' and 'Data.Word.Word' changes depending on the architecture GHC
+-- runs on.
+--
+-- Additional notes:
+--
+--  * 'Double' precision requires compute-1.3.
+--
+--  * 'Bool' is represented internally using 'Data.Word.Word8', 'Char' by
+--    'Data.Word.Word32'.
+--
+--  * If the permutation function to 'Data.Array.Accelerate.permute' resolves to
+--    non-unique indices, the combination function requires compute-1.1 to
+--    combine 32-bit types, or compute-1.2 for 64-bit types. Tuple components
+--    are resolved separately.
+--
 
 module Data.Array.Accelerate.CUDA (
 
@@ -57,10 +82,18 @@
 -- This will select the fastest device available on which to execute
 -- computations, based on compute capability and estimated maximum GFLOPS.
 --
+-- /NOTE:/
+--   GPUs typically have their own attached memory, which is separate from the
+--   computer's main memory. Hence, every 'Data.Array.Accelerate.use' operation
+--   implies copying data to the device, and every 'run' operation must copy the
+--   results of a computation back to the host. Thus, it is best to keep all
+--   computations in the 'Acc' meta-language form and only 'run' the computation
+--   once at the end, to avoid transferring (unused) intermediate results.
+--
 run :: Arrays a => Acc a -> a
 run a
   = unsafePerformIO
-  $ withMVar defaultContext $ \ctx -> evaluate (runIn ctx a)
+  $ evaluate (runIn defaultContext a)
 
 -- | As 'run', but allow the computation to continue running in a thread and
 -- return immediately without waiting for the result. The status of the
@@ -72,7 +105,7 @@
 runAsync :: Arrays a => Acc a -> Async a
 runAsync a
   = unsafePerformIO
-  $ withMVar defaultContext $ \ctx -> evaluate (runAsyncIn ctx a)
+  $ evaluate (runAsyncIn defaultContext a)
 
 -- | As 'run', but execute using the specified device context rather than using
 -- the default, automatically selected device.
@@ -111,15 +144,21 @@
 -- This function can be used to improve performance in cases where the array
 -- program is constant between invocations, because it allows us to bypass all
 -- front-end conversion stages and move directly to the execution phase. If you
--- have a computation applied repeatedly to different input data, use this.
+-- have a computation applied repeatedly to different input data, use this. If
+-- the function is only evaluated once, this is equivalent to 'run'.
 --
+-- >  let step :: Vector a -> Vector b
+-- >      step = run1 f
+-- >  in
+-- >  simulate step ...
+--
 -- See the Crystal demo, part of the 'accelerate-examples' package, for an
 -- example.
 --
 run1 :: (Arrays a, Arrays b) => (Acc a -> Acc b) -> a -> b
 run1 f
   = unsafePerformIO
-  $ withMVar defaultContext $ \ctx -> evaluate (run1In ctx f)
+  $ evaluate (run1In defaultContext f)
 
 
 -- | As 'run1', but the computation is executed asynchronously.
@@ -127,7 +166,7 @@
 run1Async :: (Arrays a, Arrays b) => (Acc a -> Acc b) -> a -> Async b
 run1Async f
   = unsafePerformIO
-  $ withMVar defaultContext $ \ctx -> evaluate (run1AsyncIn ctx f)
+  $ evaluate (run1AsyncIn defaultContext f)
 
 -- | As 'run1', but execute in the specified context.
 --
@@ -156,7 +195,7 @@
 stream :: (Arrays a, Arrays b) => (Acc a -> Acc b) -> [a] -> [b]
 stream f arrs
   = unsafePerformIO
-  $ withMVar defaultContext $ \ctx -> evaluate (streamIn ctx f arrs)
+  $ evaluate (streamIn defaultContext f arrs)
 
 -- | As 'stream', but execute in the specified context.
 --
diff --git a/Data/Array/Accelerate/CUDA/Array/Data.hs b/Data/Array/Accelerate/CUDA/Array/Data.hs
--- a/Data/Array/Accelerate/CUDA/Array/Data.hs
+++ b/Data/Array/Accelerate/CUDA/Array/Data.hs
@@ -1,4 +1,7 @@
-{-# LANGUAGE CPP, GADTs, TypeFamilies, ScopedTypeVariables #-}
+{-# LANGUAGE CPP                 #-}
+{-# LANGUAGE GADTs               #-}
+{-# LANGUAGE ScopedTypeVariables #-}
+{-# LANGUAGE TypeFamilies        #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.Array.Data
 -- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
@@ -26,15 +29,15 @@
 ) where
 
 -- libraries
-import Prelude                                          hiding (fst, snd)
+import Prelude                                          hiding ( fst, snd )
 import Data.Label.PureM
 import Control.Applicative
 import Control.Monad.Trans
 
 -- friends
 import Data.Array.Accelerate.Array.Data
-import Data.Array.Accelerate.Array.Sugar                (Array(..), Shape, Elt, fromElt, toElt)
-import Data.Array.Accelerate.Array.Representation       (size, index)
+import Data.Array.Accelerate.Array.Sugar                ( Array(..), Shape, Elt, fromElt, toElt )
+import Data.Array.Accelerate.Array.Representation       ( size, index )
 import Data.Array.Accelerate.CUDA.State
 import Data.Array.Accelerate.CUDA.Array.Table
 import qualified Data.Array.Accelerate.CUDA.Array.Prim  as Prim
@@ -61,6 +64,14 @@
 snd :: ArrayData (a,b) -> ArrayData b
 snd = sndArrayData
 
+-- Extract the state information to pass along to the primitive data handlers
+--
+run :: (Context -> MemoryTable -> IO a) -> CIO a
+run f = do
+  ctx   <- gets activeContext
+  mt    <- gets memoryTable
+  liftIO $ f ctx mt
+
 -- CPP hackery to generate the cases where we dispatch to the worker function handling
 -- elementary types.
 --
@@ -77,53 +88,53 @@
 ; dispatcher ArrayEltRword64 = worker                                       \
 ; dispatcher ArrayEltRfloat  = worker                                       \
 ; dispatcher ArrayEltRdouble = worker                                       \
-; dispatcher ArrayEltRbool   = error "mkPrimDispatcher: ArrayEltRbool"      \
-; dispatcher ArrayEltRchar   = error "mkPrimDispatcher: ArrayEltRchar"      \
+; dispatcher ArrayEltRbool   = worker                                       \
+; dispatcher ArrayEltRchar   = worker                                       \
 ; dispatcher _               = error "mkPrimDispatcher: not primitive"
 
 
 -- |Allocate a new device array to accompany the given host-side array.
 --
 mallocArray :: (Shape dim, Elt e) => Array dim e -> CIO ()
-mallocArray (Array sh adata) = doMalloc =<< gets memoryTable
+mallocArray (Array sh adata) = run doMalloc
   where
-    doMalloc mt = liftIO $ mallocR arrayElt adata
+    doMalloc ctx mt = mallocR arrayElt adata
       where
         mallocR :: ArrayEltR e -> ArrayData e -> IO ()
         mallocR ArrayEltRunit             _  = return ()
         mallocR (ArrayEltRpair aeR1 aeR2) ad = mallocR aeR1 (fst ad) >> mallocR aeR2 (snd ad)
-        mallocR aer                       ad = mallocPrim aer mt ad (size sh)
+        mallocR aer                       ad = mallocPrim aer ctx mt ad (size sh)
         --
-        mallocPrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> Int -> IO ()
+        mallocPrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> Int -> IO ()
         mkPrimDispatch(mallocPrim,Prim.mallocArray)
 
 
 -- |Upload an existing array to the device
 --
 useArray :: (Shape dim, Elt e) => Array dim e -> CIO ()
-useArray (Array sh adata) = doUse =<< gets memoryTable
+useArray (Array sh adata) = run doUse
   where
-    doUse mt = liftIO $ useR arrayElt adata
+    doUse ctx mt = useR arrayElt adata
       where
         useR :: ArrayEltR e -> ArrayData e -> IO ()
         useR ArrayEltRunit             _  = return ()
         useR (ArrayEltRpair aeR1 aeR2) ad = useR aeR1 (fst ad) >> useR aeR2 (snd ad)
-        useR aer                       ad = usePrim aer mt ad (size sh)
+        useR aer                       ad = usePrim aer ctx mt ad (size sh)
         --
-        usePrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> Int -> IO ()
+        usePrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> Int -> IO ()
         mkPrimDispatch(usePrim,Prim.useArray)
 
 useArrayAsync :: (Shape dim, Elt e) => Array dim e -> Maybe CUDA.Stream -> CIO ()
-useArrayAsync (Array sh adata) ms = doUse =<< gets memoryTable
+useArrayAsync (Array sh adata) ms = run doUse
   where
-    doUse mt = liftIO $ useR arrayElt adata
+    doUse ctx mt = useR arrayElt adata
       where
         useR :: ArrayEltR e -> ArrayData e -> IO ()
         useR ArrayEltRunit             _  = return ()
         useR (ArrayEltRpair aeR1 aeR2) ad = useR aeR1 (fst ad) >> useR aeR2 (snd ad)
-        useR aer                       ad = usePrim aer mt ad (size sh) ms
+        useR aer                       ad = usePrim aer ctx mt ad (size sh) ms
         --
-        usePrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> Int -> Maybe CUDA.Stream -> IO ()
+        usePrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> Int -> Maybe CUDA.Stream -> IO ()
         mkPrimDispatch(usePrim,Prim.useArrayAsync)
 
 
@@ -131,19 +142,33 @@
 -- synchronous operation.
 --
 indexArray :: (Shape dim, Elt e) => Array dim e -> dim -> CIO e
-indexArray (Array sh adata) ix = doIndex =<< gets memoryTable
+indexArray (Array sh adata) ix = run doIndex
   where
-    i          = index sh (fromElt ix)
-    doIndex mt = toElt <$> (liftIO $ indexR arrayElt adata)
+    i              = index sh (fromElt ix)
+    doIndex ctx mt = toElt <$> indexR arrayElt adata
       where
         indexR :: ArrayEltR e -> ArrayData e -> IO e
         indexR ArrayEltRunit             _  = return ()
         indexR (ArrayEltRpair aeR1 aeR2) ad = (,) <$> indexR aeR1 (fst ad)
                                                   <*> indexR aeR2 (snd ad)
-        indexR aer                       ad = indexPrim aer mt ad i
         --
-        indexPrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> Int -> IO e
-        mkPrimDispatch(indexPrim,Prim.indexArray)
+        indexR ArrayEltRbool             ad = toBool <$> Prim.indexArray ctx mt ad i
+          where toBool 0 = False
+                toBool _ = True
+        --
+        indexR ArrayEltRint              ad = Prim.indexArray ctx mt ad i
+        indexR ArrayEltRint8             ad = Prim.indexArray ctx mt ad i
+        indexR ArrayEltRint16            ad = Prim.indexArray ctx mt ad i
+        indexR ArrayEltRint32            ad = Prim.indexArray ctx mt ad i
+        indexR ArrayEltRint64            ad = Prim.indexArray ctx mt ad i
+        indexR ArrayEltRword             ad = Prim.indexArray ctx mt ad i
+        indexR ArrayEltRword8            ad = Prim.indexArray ctx mt ad i
+        indexR ArrayEltRword16           ad = Prim.indexArray ctx mt ad i
+        indexR ArrayEltRword32           ad = Prim.indexArray ctx mt ad i
+        indexR ArrayEltRword64           ad = Prim.indexArray ctx mt ad i
+        indexR ArrayEltRfloat            ad = Prim.indexArray ctx mt ad i
+        indexR ArrayEltRdouble           ad = Prim.indexArray ctx mt ad i
+        indexR ArrayEltRchar             ad = Prim.indexArray ctx mt ad i
 
 
 -- |Copy data between two device arrays. The operation is asynchronous with
@@ -152,75 +177,75 @@
 copyArray :: (Shape dim, Elt e) => Array dim e -> Array dim e -> CIO ()
 copyArray (Array sh1 adata1) (Array sh2 adata2)
   = BOUNDS_CHECK(check) "copyArray" "shape mismatch" (sh1 == sh2)
-  $ doCopy =<< gets memoryTable
+  $ run doCopy
   where
-    doCopy mt = liftIO $ copyR arrayElt adata1 adata2
+    doCopy ctx mt = copyR arrayElt adata1 adata2
       where
         copyR :: ArrayEltR e -> ArrayData e -> ArrayData e -> IO ()
         copyR ArrayEltRunit             _   _   = return ()
         copyR (ArrayEltRpair aeR1 aeR2) ad1 ad2 = copyR aeR1 (fst ad1) (fst ad2) >>
                                                   copyR aeR2 (snd ad1) (snd ad2)
-        copyR aer                       ad1 ad2 = copyPrim aer mt ad1 ad2 (size sh1)
+        copyR aer                       ad1 ad2 = copyPrim aer ctx mt ad1 ad2 (size sh1)
         --
-        copyPrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> ArrayData e -> Int -> IO ()
+        copyPrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> ArrayData e -> Int -> IO ()
         mkPrimDispatch(copyPrim,Prim.copyArray)
 
 
 -- Copy data from the device into the associated Accelerate host-side array
 --
 peekArray :: (Shape dim, Elt e) => Array dim e -> CIO ()
-peekArray (Array sh adata) = doPeek =<< gets memoryTable
+peekArray (Array sh adata) = run doPeek
   where
-    doPeek mt = liftIO $ peekR arrayElt adata
+    doPeek ctx mt = peekR arrayElt adata
       where
         peekR :: ArrayEltR e -> ArrayData e -> IO ()
         peekR ArrayEltRunit             _  = return ()
         peekR (ArrayEltRpair aeR1 aeR2) ad = peekR aeR1 (fst ad) >> peekR aeR2 (snd ad)
-        peekR aer                       ad = peekPrim aer mt ad (size sh)
+        peekR aer                       ad = peekPrim aer ctx mt ad (size sh)
         --
-        peekPrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> Int -> IO ()
+        peekPrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> Int -> IO ()
         mkPrimDispatch(peekPrim,Prim.peekArray)
 
 peekArrayAsync :: (Shape dim, Elt e) => Array dim e -> Maybe CUDA.Stream -> CIO ()
-peekArrayAsync (Array sh adata) ms = doPeek =<< gets memoryTable
+peekArrayAsync (Array sh adata) ms = run doPeek
   where
-    doPeek mt = liftIO $ peekR arrayElt adata
+    doPeek ctx mt = peekR arrayElt adata
       where
         peekR :: ArrayEltR e -> ArrayData e -> IO ()
         peekR ArrayEltRunit             _  = return ()
         peekR (ArrayEltRpair aeR1 aeR2) ad = peekR aeR1 (fst ad) >> peekR aeR2 (snd ad)
-        peekR aer                       ad = peekPrim aer mt ad (size sh) ms
+        peekR aer                       ad = peekPrim aer ctx mt ad (size sh) ms
         --
-        peekPrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> Int -> Maybe CUDA.Stream -> IO ()
+        peekPrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> Int -> Maybe CUDA.Stream -> IO ()
         mkPrimDispatch(peekPrim,Prim.peekArrayAsync)
 
 
 -- Copy data from an Accelerate array into the associated device array
 --
 pokeArray :: (Shape dim, Elt e) => Array dim e -> CIO ()
-pokeArray (Array sh adata) = doPoke =<< gets memoryTable
+pokeArray (Array sh adata) = run doPoke
   where
-    doPoke mt = liftIO $ pokeR arrayElt adata
+    doPoke ctx mt = pokeR arrayElt adata
       where
         pokeR :: ArrayEltR e -> ArrayData e -> IO ()
         pokeR ArrayEltRunit             _  = return ()
         pokeR (ArrayEltRpair aeR1 aeR2) ad = pokeR aeR1 (fst ad) >> pokeR aeR2 (snd ad)
-        pokeR aer                       ad = pokePrim aer mt ad (size sh)
+        pokeR aer                       ad = pokePrim aer ctx mt ad (size sh)
         --
-        pokePrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> Int -> IO ()
+        pokePrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> Int -> IO ()
         mkPrimDispatch(pokePrim,Prim.pokeArray)
 
 pokeArrayAsync :: (Shape dim, Elt e) => Array dim e -> Maybe CUDA.Stream -> CIO ()
-pokeArrayAsync (Array sh adata) ms = doPoke =<< gets memoryTable
+pokeArrayAsync (Array sh adata) ms = run doPoke
   where
-    doPoke mt = liftIO $ pokeR arrayElt adata
+    doPoke ctx mt = pokeR arrayElt adata
       where
         pokeR :: ArrayEltR e -> ArrayData e -> IO ()
         pokeR ArrayEltRunit             _  = return ()
         pokeR (ArrayEltRpair aeR1 aeR2) ad = pokeR aeR1 (fst ad) >> pokeR aeR2 (snd ad)
-        pokeR aer                       ad = pokePrim aer mt ad (size sh) ms
+        pokeR aer                       ad = pokePrim aer ctx mt ad (size sh) ms
         --
-        pokePrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> Int -> Maybe CUDA.Stream -> IO ()
+        pokePrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> Int -> Maybe CUDA.Stream -> IO ()
         mkPrimDispatch(pokePrim,Prim.pokeArrayAsync)
 
 
@@ -244,17 +269,17 @@
 -- that can be passed to a kernel upon invocation.
 --
 marshalArrayData :: ArrayElt e => ArrayData e -> CIO [CUDA.FunParam]
-marshalArrayData adata = doMarshal =<< gets memoryTable
+marshalArrayData adata = run doMarshal
   where
-    doMarshal mt = liftIO $ marshalR arrayElt adata
+    doMarshal ctx mt = marshalR arrayElt adata
       where
         marshalR :: ArrayEltR e -> ArrayData e -> IO [CUDA.FunParam]
         marshalR ArrayEltRunit             _  = return []
         marshalR (ArrayEltRpair aeR1 aeR2) ad = (++) <$> marshalR aeR1 (fst ad)
                                                      <*> marshalR aeR2 (snd ad)
-        marshalR aer                       ad = return <$> marshalPrim aer mt ad
+        marshalR aer                       ad = return <$> marshalPrim aer ctx mt ad
         --
-        marshalPrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> IO CUDA.FunParam
+        marshalPrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> IO CUDA.FunParam
         mkPrimDispatch(marshalPrim,Prim.marshalArrayData)
 
 
@@ -263,9 +288,9 @@
 -- consumed, in projection index order --- i.e. right-to-left
 --
 marshalTextureData :: ArrayElt e => ArrayData e -> Int -> [CUDA.Texture] -> CIO ()
-marshalTextureData adata n texs = doMarshal =<< gets memoryTable
+marshalTextureData adata n texs = run doMarshal
   where
-    doMarshal mt = liftIO $ marshalR arrayElt adata texs >> return ()
+    doMarshal ctx mt = marshalR arrayElt adata texs >> return ()
       where
         marshalR :: ArrayEltR e -> ArrayData e -> [CUDA.Texture] -> IO Int
         marshalR ArrayEltRunit             _  _ = return 0
@@ -274,27 +299,27 @@
                l <- marshalR aeR1 (fst ad) (drop r t)
                return (l + r)
         marshalR aer                       ad t
-          = do marshalPrim aer mt ad n (head t)
+          = do marshalPrim aer ctx mt ad n (head t)
                return 1
         --
-        marshalPrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> Int -> CUDA.Texture -> IO ()
+        marshalPrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> Int -> CUDA.Texture -> IO ()
         mkPrimDispatch(marshalPrim,Prim.marshalTextureData)
 
 
 -- |Raw device pointers associated with a host-side array
 --
 devicePtrsOfArrayData :: ArrayElt e => ArrayData e -> CIO (Prim.DevicePtrs e)
-devicePtrsOfArrayData adata = ptrs =<< gets memoryTable
+devicePtrsOfArrayData adata = run ptrs
   where
-    ptrs mt = liftIO $ ptrsR arrayElt adata
+    ptrs ctx mt = ptrsR arrayElt adata
       where
         ptrsR :: ArrayEltR e -> ArrayData e -> IO (Prim.DevicePtrs e)
         ptrsR ArrayEltRunit             _  = return ()
         ptrsR (ArrayEltRpair aeR1 aeR2) ad = (,) <$> ptrsR aeR1 (fst ad)
                                                  <*> ptrsR aeR2 (snd ad)
-        ptrsR aer                       ad = ptrsPrim aer mt ad
+        ptrsR aer                       ad = ptrsPrim aer ctx mt ad
         --
-        ptrsPrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> IO (Prim.DevicePtrs e)
+        ptrsPrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> IO (Prim.DevicePtrs e)
         mkPrimDispatch(ptrsPrim,Prim.devicePtrsOfArrayData)
 
 
diff --git a/Data/Array/Accelerate/CUDA/Array/Prim.hs b/Data/Array/Accelerate/CUDA/Array/Prim.hs
--- a/Data/Array/Accelerate/CUDA/Array/Prim.hs
+++ b/Data/Array/Accelerate/CUDA/Array/Prim.hs
@@ -1,4 +1,8 @@
-{-# LANGUAGE BangPatterns, CPP, GADTs, TypeFamilies, ScopedTypeVariables #-}
+{-# LANGUAGE BangPatterns        #-}
+{-# LANGUAGE CPP                 #-}
+{-# LANGUAGE GADTs               #-}
+{-# LANGUAGE ScopedTypeVariables #-}
+{-# LANGUAGE TypeFamilies        #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.Array.Prim
 -- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
@@ -88,14 +92,10 @@
 -- CFloat
 -- CDouble
 
--- FIXME:
--- No concrete implementation in Data.Array.Accelerate.Array.Data
---
-type instance HostPtrs   Bool = ()
-type instance DevicePtrs Bool = ()
+type instance HostPtrs   Bool = CUDA.HostPtr   Word8
+type instance DevicePtrs Bool = CUDA.DevicePtr Word8
 
-type instance HostPtrs   Char = ()
-type instance DevicePtrs Char = ()
+primArrayElt(Char)
 
 -- FIXME:
 -- CChar
@@ -126,19 +126,20 @@
 instance TextureData Word64 where format _ = (CUDA.Word32, 2)
 instance TextureData Float  where format _ = (CUDA.Float,  1)
 instance TextureData Double where format _ = (CUDA.Int32,  2)
-
-instance TextureData Int where
-  format _ = case sizeOf (undefined :: Int) of
-                  4 -> (CUDA.Int32, 1)
-                  8 -> (CUDA.Int32, 2)
-                  _ -> error "we can never get here"
-
-instance TextureData Word where
-  format _ = case sizeOf (undefined :: Word) of
-                  4 -> (CUDA.Word32, 1)
-                  8 -> (CUDA.Word32, 2)
-                  _ -> error "we can never get here"
-
+instance TextureData Bool   where format _ = (CUDA.Word8,  1)
+#if   SIZEOF_HSINT == 4
+instance TextureData Int    where format _ = (CUDA.Int32,  1)
+#elif SIZEOF_HSINT == 8
+instance TextureData Int    where format _ = (CUDA.Int32,  2)
+#endif
+#if   SIZEOF_HSINT == 4
+instance TextureData Word   where format _ = (CUDA.Word32, 1)
+#elif SIZEOF_HSINT == 8
+instance TextureData Word   where format _ = (CUDA.Word32, 2)
+#endif
+#if SIZEOF_HSCHAR == 4
+instance TextureData Char   where format _ = (CUDA.Word32, 1)
+#endif
 
 
 -- Primitive array operations
@@ -150,20 +151,21 @@
 --
 mallocArray
     :: forall e a. (ArrayElt e, DevicePtrs e ~ CUDA.DevicePtr a, Typeable e, Typeable a, Storable a)
-    => MemoryTable
+    => Context
+    -> MemoryTable
     -> ArrayData e
     -> Int
     -> IO ()
-mallocArray !mt !ad !n0 = do
+mallocArray !ctx !mt !ad !n0 = do
   let !n = 1 `max` n0
-  exists <- isJust <$> (lookup mt ad :: IO (Maybe (CUDA.DevicePtr a)))
+  exists <- isJust <$> (lookup ctx mt ad :: IO (Maybe (CUDA.DevicePtr a)))
   unless exists $ do
     message $ "mallocArray: " ++ showBytes (n * sizeOf (undefined::a))
     ptr <- CUDA.mallocArray n `catch` \(e :: CUDAException) ->
       case e of
         ExitCode OutOfMemory -> reclaim mt >> CUDA.mallocArray n
         _                    -> throwIO e
-    insert mt ad (ptr :: CUDA.DevicePtr a)
+    insert ctx mt ad (ptr :: CUDA.DevicePtr a)
 
 
 -- A combination of 'mallocArray' and 'pokeArray' to allocate space on the
@@ -172,15 +174,16 @@
 --
 useArray
     :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr a, Typeable e, Typeable a, Storable a)
-    => MemoryTable
+    => Context
+    -> MemoryTable
     -> ArrayData e
     -> Int
     -> IO ()
-useArray !mt !ad !n0 =
+useArray !ctx !mt !ad !n0 =
   let src = ptrsOfArrayData ad
       !n  = 1 `max` n0
   in do
-    exists <- isJust <$> (lookup mt ad :: IO (Maybe (CUDA.DevicePtr a)))
+    exists <- isJust <$> (lookup ctx mt ad :: IO (Maybe (CUDA.DevicePtr a)))
     unless exists $ do
       message $ "useArray/malloc: " ++ showBytes (n * sizeOf (undefined::a))
       dst <- CUDA.mallocArray n `catch` \(e :: CUDAException) ->
@@ -188,21 +191,22 @@
           ExitCode OutOfMemory -> reclaim mt >> CUDA.mallocArray n
           _                    -> throwIO e
       CUDA.pokeArray n src dst
-      insert mt ad dst
+      insert ctx mt ad dst
 
 
 useArrayAsync
     :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr a, Typeable e, Typeable a, Storable a)
-    => MemoryTable
+    => Context
+    -> MemoryTable
     -> ArrayData e
     -> Int
     -> Maybe CUDA.Stream
     -> IO ()
-useArrayAsync !mt !ad !n0 !ms =
+useArrayAsync !ctx !mt !ad !n0 !ms =
   let src = CUDA.HostPtr (ptrsOfArrayData ad)
       !n  = 1 `max` n0
   in do
-    exists <- isJust <$> (lookup mt ad :: IO (Maybe (CUDA.DevicePtr a)))
+    exists <- isJust <$> (lookup ctx mt ad :: IO (Maybe (CUDA.DevicePtr a)))
     unless exists $ do
       message $ "useArrayAsync/malloc: " ++ showBytes (n * sizeOf (undefined::a))
       dst <- CUDA.mallocArray n `catch` \(e :: CUDAException) ->
@@ -210,20 +214,22 @@
           ExitCode OutOfMemory -> reclaim mt >> CUDA.mallocArray n
           _                    -> throwIO e
       CUDA.pokeArrayAsync n src dst ms
-      insert mt ad dst
+      insert ctx mt ad dst
 
 
 -- Read a single element from an array at the given row-major index
 --
 indexArray
-    :: (ArrayElt e, DevicePtrs e ~ CUDA.DevicePtr b, Typeable e, Typeable b, Storable b)
-    => MemoryTable
+    :: forall e a. (ArrayElt e, DevicePtrs e ~ CUDA.DevicePtr a, Typeable e, Typeable a, Storable a)
+    => Context
+    -> MemoryTable
     -> ArrayData e
     -> Int
-    -> IO b
-indexArray !mt !ad !i =
-  alloca                        $ \dst ->
-  devicePtrsOfArrayData mt ad >>= \src -> do
+    -> IO a
+indexArray !ctx !mt !ad !i =
+  alloca                            $ \dst ->
+  devicePtrsOfArrayData ctx mt ad >>= \src -> do
+    message $ "indexArray: " ++ showBytes (sizeOf (undefined::a))
     CUDA.peekArray 1 (src `CUDA.advanceDevPtr` i) dst
     peek dst
 
@@ -232,63 +238,73 @@
 -- respect to the host, but will never overlap kernel execution.
 --
 copyArray
-    :: (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr b, Typeable a, Typeable b, Typeable e, Storable b)
-    => MemoryTable
+    :: forall e a b. (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr b, Typeable a, Typeable b, Typeable e, Storable b)
+    => Context
+    -> MemoryTable
     -> ArrayData e              -- source array
     -> ArrayData e              -- destination array
     -> Int                      -- number of array elements
     -> IO ()
-copyArray !mt !from !to !n = do
-  src <- devicePtrsOfArrayData mt from
-  dst <- devicePtrsOfArrayData mt to
+copyArray !ctx !mt !from !to !n = do
+  message $ "copyArrayAsync: " ++ showBytes (n * sizeOf (undefined :: b))
+  src <- devicePtrsOfArrayData ctx mt from
+  dst <- devicePtrsOfArrayData ctx mt to
   CUDA.copyArrayAsync n src dst
 
 
 -- Copy data from the device into the associated Accelerate host-side array
 --
 peekArray
-    :: (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr a, Typeable a, Typeable e, Storable a)
-    => MemoryTable
+    :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr a, Typeable a, Typeable e, Storable a)
+    => Context
+    -> MemoryTable
     -> ArrayData e
     -> Int
     -> IO ()
-peekArray !mt !ad !n =
-  devicePtrsOfArrayData mt ad >>= \src ->
+peekArray !ctx !mt !ad !n =
+  devicePtrsOfArrayData ctx mt ad >>= \src -> do
+    message $ "peekArray: " ++ showBytes (n * sizeOf (undefined :: a))
     CUDA.peekArray n src (ptrsOfArrayData ad)
 
 peekArrayAsync
-    :: (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr a, Typeable a, Typeable e, Storable a)
-    => MemoryTable
+    :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr a, Typeable a, Typeable e, Storable a)
+    => Context
+    -> MemoryTable
     -> ArrayData e
     -> Int
     -> Maybe CUDA.Stream
     -> IO ()
-peekArrayAsync !mt !ad !n !st =
-  devicePtrsOfArrayData mt ad >>= \src ->
+peekArrayAsync !ctx !mt !ad !n !st =
+  devicePtrsOfArrayData ctx mt ad >>= \src -> do
+    message $ "peekArrayAsync: " ++ showBytes (n * sizeOf (undefined :: a))
     CUDA.peekArrayAsync n src (CUDA.HostPtr $ ptrsOfArrayData ad) st
 
 
 -- Copy data from an Accelerate array into the associated device array
 --
 pokeArray
-    :: (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr a, Typeable a, Typeable e, Storable a)
-    => MemoryTable
+    :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr a, Typeable a, Typeable e, Storable a)
+    => Context
+    -> MemoryTable
     -> ArrayData e
     -> Int
     -> IO ()
-pokeArray !mt !ad !n =
-  devicePtrsOfArrayData mt ad >>= \dst ->
+pokeArray !ctx !mt !ad !n =
+  devicePtrsOfArrayData ctx mt ad >>= \dst -> do
+    message $ "pokeArrayAsync: " ++ showBytes (n * sizeOf (undefined :: a))
     CUDA.pokeArray n (ptrsOfArrayData ad) dst
 
 pokeArrayAsync
-    :: (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr a, Typeable a, Typeable e, Storable a)
-    => MemoryTable
+    :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr a, Typeable a, Typeable e, Storable a)
+    => Context
+    -> MemoryTable
     -> ArrayData e
     -> Int
     -> Maybe CUDA.Stream
     -> IO ()
-pokeArrayAsync !mt !ad !n !st =
-  devicePtrsOfArrayData mt ad >>= \dst ->
+pokeArrayAsync !ctx !mt !ad !n !st =
+  devicePtrsOfArrayData ctx mt ad >>= \dst -> do
+    message $ "pokeArrayAsync: " ++ showBytes (n * sizeOf (undefined :: a))
     CUDA.pokeArrayAsync n (CUDA.HostPtr $ ptrsOfArrayData ad) dst st
 
 
@@ -307,24 +323,26 @@
 --
 marshalArrayData
     :: (ArrayElt e, DevicePtrs e ~ CUDA.DevicePtr b, Typeable b, Typeable e)
-    => MemoryTable
+    => Context
+    -> MemoryTable
     -> ArrayData e
     -> IO CUDA.FunParam
-marshalArrayData !mt !ad = marshalDevicePtrs ad <$> devicePtrsOfArrayData mt ad
+marshalArrayData !ctx !mt !ad = marshalDevicePtrs ad <$> devicePtrsOfArrayData ctx mt ad
 
 
 -- Bind device memory to the given texture reference, setting appropriate type
 --
 marshalTextureData
     :: forall a e. (ArrayElt e, DevicePtrs e ~ CUDA.DevicePtr a, Typeable a, Typeable e, Storable a, TextureData a)
-    => MemoryTable
+    => Context
+    -> MemoryTable
     -> ArrayData e              -- host array
     -> Int                      -- number of elements
     -> CUDA.Texture             -- texture reference to bind array to
     -> IO ()
-marshalTextureData !mt !ad !n !tex =
+marshalTextureData !ctx !mt !ad !n !tex =
   let (fmt, c) = format (undefined :: a)
-  in  devicePtrsOfArrayData mt ad >>= \ptr -> do
+  in  devicePtrsOfArrayData ctx mt ad >>= \ptr -> do
         CUDA.setFormat tex fmt c
         CUDA.bind tex ptr (fromIntegral $ n * sizeOf (undefined :: a))
 
@@ -333,11 +351,12 @@
 --
 devicePtrsOfArrayData
     :: (ArrayElt e, DevicePtrs e ~ CUDA.DevicePtr b, Typeable e, Typeable b)
-    => MemoryTable
+    => Context
+    -> MemoryTable
     -> ArrayData e
     -> IO (DevicePtrs e)
-devicePtrsOfArrayData !mt !ad = do
-  mv <- lookup mt ad
+devicePtrsOfArrayData !ctx !mt !ad = do
+  mv <- lookup ctx mt ad
   case mv of
     Just v  -> return v
     Nothing -> do
diff --git a/Data/Array/Accelerate/CUDA/Array/Table.hs b/Data/Array/Accelerate/CUDA/Array/Table.hs
--- a/Data/Array/Accelerate/CUDA/Array/Table.hs
+++ b/Data/Array/Accelerate/CUDA/Array/Table.hs
@@ -1,4 +1,7 @@
-{-# LANGUAGE BangPatterns, CPP, GADTs, PatternGuards #-}
+{-# LANGUAGE BangPatterns  #-}
+{-# LANGUAGE CPP           #-}
+{-# LANGUAGE GADTs         #-}
+{-# LANGUAGE PatternGuards #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.Array.Table
 -- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
@@ -13,7 +16,7 @@
 module Data.Array.Accelerate.CUDA.Array.Table (
 
   -- Tables for host/device memory associations
-  MemoryTable, new, lookup, insert, reclaim
+  MemoryTable, Context(..), new, lookup, insert, reclaim
 
 ) where
 
@@ -24,7 +27,7 @@
 import Data.Typeable                                    ( Typeable, gcast )
 import Control.Monad                                    ( unless )
 import Control.Exception                                ( bracket_ )
-import Control.Applicative                              ( (<$>), (<*>) )
+import Control.Applicative                              ( (<$>) )
 import System.Mem                                       ( performGC )
 import System.Mem.Weak                                  ( Weak, mkWeak, deRefWeak, finalize )
 import System.Mem.StableName                            ( StableName, makeStableName, hashStableName )
@@ -35,7 +38,7 @@
 import qualified Data.HashTable.IO                      as HT
 
 import Data.Array.Accelerate.Array.Data                 ( ArrayData )
-import qualified Data.Array.Accelerate.CUDA.Debug       as D ( message, dump_gc )
+import qualified Data.Array.Accelerate.CUDA.Debug       as D
 
 #include "accelerate.h"
 
@@ -53,15 +56,22 @@
 -- return Nothing. References from 'val' to the key are ignored (see the
 -- semantics of weak pointers in the documentation).
 --
-type HashTable key val = HT.BasicHashTable key val
-type MT                = IORef ( HashTable HostArray DeviceArray )
-data MemoryTable       = MemoryTable {-# UNPACK #-} !MT
-                                     {-# UNPACK #-} !(Weak MT)
+type HashTable key val  = HT.BasicHashTable key val
+type MT                 = IORef ( HashTable HostArray DeviceArray )
+data MemoryTable        = MemoryTable {-# UNPACK #-} !MT
+                                      {-# UNPACK #-} !(Weak MT)
 
+-- The currently active context. Finaliser threads need to check if the context
+-- is still active before attempting to release their associated memory.
+--
+data Context = Context {-# UNPACK #-} !CUDA.Context
+                       {-# UNPACK #-} !(Weak CUDA.Context)
 
+-- Arrays on the host and device
+--
 data HostArray where
   HostArray :: Typeable e
-            => {-# UNPACK #-} !CUDA.Context
+            => {-# UNPACK #-} !Int      -- unique ID relating to the parent context
             -> {-# UNPACK #-} !(StableName (ArrayData e))
             -> HostArray
 
@@ -75,8 +85,7 @@
     = maybe False (== a2) (gcast a1)
 
 instance Hashable HostArray where
-  hash (HostArray (CUDA.Context p) sn) =
-    fromIntegral (ptrToIntPtr p) `hashWithSalt` sn
+  hash (HostArray cid sn) = hashWithSalt cid sn
 
 instance Show HostArray where
   show (HostArray _ sn) = "Array #" ++ show (hashStableName sn)
@@ -98,28 +107,28 @@
 
 -- Look for the device memory corresponding to a given host-side array.
 --
-lookup :: (Typeable a, Typeable b) => MemoryTable -> ArrayData a -> IO (Maybe (DevicePtr b))
-lookup (MemoryTable ref _) !arr = do
-  sa <- makeStableArray arr
+lookup :: (Typeable a, Typeable b) => Context -> MemoryTable -> ArrayData a -> IO (Maybe (DevicePtr b))
+lookup ctx (MemoryTable ref _) !arr = do
+  sa <- makeStableArray ctx arr
   mw <- withIORef ref (`HT.lookup` sa)
   case mw of
     Nothing              -> trace ("lookup/not found: " ++ show sa) $ return Nothing
     Just (DeviceArray w) -> do
       mv <- deRefWeak w
       case mv of
-        Just v | Just p <- gcast v   -> trace ("lookup/found: " ++ show sa) $ return (Just p)
-               | otherwise           -> INTERNAL_ERROR(error) "lookup" $ "type mismatch"
-        Nothing                      ->
-          makeStableArray arr >>= \x -> INTERNAL_ERROR(error) "lookup" $ "dead weak pair: " ++ show x
+        Just v | Just p <- gcast v -> trace ("lookup/found: " ++ show sa) $ return (Just p)
+               | otherwise         -> INTERNAL_ERROR(error) "lookup" $ "type mismatch"
+        Nothing                    ->
+          makeStableArray ctx arr >>= \x -> INTERNAL_ERROR(error) "lookup" $ "dead weak pair: " ++ show x
 
 
 -- Record an association between a host-side array and a new device memory area.
 -- The device memory will be freed when the host array is garbage collected.
 --
-insert :: (Typeable a, Typeable b) => MemoryTable -> ArrayData a -> DevicePtr b -> IO ()
-insert (MemoryTable ref weak_ref) !arr !ptr = do
-  key  <- makeStableArray arr
-  dev  <- DeviceArray `fmap` mkWeak arr ptr (Just $ finalizer weak_ref key ptr)
+insert :: (Typeable a, Typeable b) => Context -> MemoryTable -> ArrayData a -> DevicePtr b -> IO ()
+insert ctx@(Context _ weak_ctx) (MemoryTable ref weak_ref) !arr !ptr = do
+  key  <- makeStableArray ctx arr
+  dev  <- DeviceArray `fmap` mkWeak arr ptr (Just $ finalizer weak_ctx weak_ref key ptr)
   tbl  <- readIORef ref
   message $ "insert: " ++ show key
   HT.insert tbl key dev
@@ -133,7 +142,8 @@
 --
 reclaim :: MemoryTable -> IO ()
 reclaim (MemoryTable _ weak_ref) = do
-  trace "reclaim" performGC
+  (free, total) <- CUDA.getMemInfo
+  performGC
   mr <- deRefWeak weak_ref
   case mr of
     Nothing  -> return ()
@@ -141,7 +151,12 @@
       flip HT.mapM_ tbl $ \(_,DeviceArray w) -> do
         alive <- isJust `fmap` deRefWeak w
         unless alive $ finalize w
-
+  --
+  D.when D.dump_gc $ do
+    (free', _)  <- CUDA.getMemInfo
+    message $ "reclaim: freed "   ++ showBytes (fromIntegral (free - free'))
+                        ++ ", "   ++ showBytes (fromIntegral free')
+                        ++ " of " ++ showBytes (fromIntegral total) ++ " remaining"
 
 -- Because a finaliser might run at any time, we must reinstate the context in
 -- which the array was allocated before attempting to release it.
@@ -152,17 +167,17 @@
 -- the hash tables --- but we must do this first before failing to use a dead
 -- context.
 --
-finalizer :: Weak MT -> HostArray -> DevicePtr b -> IO ()
-finalizer !weak_ref !key@(HostArray ctx _) !ptr = do
+finalizer :: Weak CUDA.Context -> Weak MT -> HostArray -> DevicePtr b -> IO ()
+finalizer !weak_ctx !weak_ref !key !ptr = do
   mr <- deRefWeak weak_ref
   case mr of
-    Nothing  -> trace ("finalise/dead table: " ++ show key) $ return ()
-    Just ref -> trace ("finalise: "            ++ show key) $ withIORef ref (`HT.delete` key)
+    Nothing  -> message ("finalise/dead table: " ++ show key)
+    Just ref -> trace   ("finalise: "            ++ show key) $ withIORef ref (`HT.delete` key)
   --
-  bracket_
-    (CUDA.push ctx)
-    (CUDA.pop)
-    (CUDA.free ptr)
+  mc <- deRefWeak weak_ctx
+  case mc of
+    Nothing  -> message ("finalise/dead context: " ++ show key)
+    Just ctx -> bracket_ (CUDA.push ctx) CUDA.pop (CUDA.free ptr)
 
 
 table_finalizer :: HashTable HostArray DeviceArray -> IO ()
@@ -175,8 +190,10 @@
 -- -------------
 
 {-# INLINE makeStableArray #-}
-makeStableArray :: Typeable a => ArrayData a -> IO HostArray
-makeStableArray !arr = HostArray <$> CUDA.get <*> makeStableName arr
+makeStableArray :: Typeable a => Context -> ArrayData a -> IO HostArray
+makeStableArray (Context (CUDA.Context !p) !_) !arr =
+  let cid = fromIntegral (ptrToIntPtr p)
+  in  HostArray cid <$> makeStableName arr
 
 {-# INLINE withIORef #-}
 withIORef :: IORef a -> (a -> IO b) -> IO b
@@ -185,6 +202,10 @@
 
 -- Debug
 -- -----
+
+{-# INLINE showBytes #-}
+showBytes :: Int -> String
+showBytes x = D.showFFloatSIBase (Just 0) 1024 (fromIntegral x :: Double) "B"
 
 {-# INLINE trace #-}
 trace :: String -> IO a -> IO a
diff --git a/Data/Array/Accelerate/CUDA/CodeGen.hs b/Data/Array/Accelerate/CUDA/CodeGen.hs
--- a/Data/Array/Accelerate/CUDA/CodeGen.hs
+++ b/Data/Array/Accelerate/CUDA/CodeGen.hs
@@ -184,7 +184,7 @@
           idx'    = show $ idxToInt idx
           sh      = cshape ("sh" ++ idx') (accDim avar)
           ty      = accTypeTex avar
-          arr n   = "arr" ++ idx' ++ "_a" ++ show (n::Int)
+          arr n   = "avar" ++ idx' ++ "_a" ++ show (n::Int)
       in
       sh : zipWith (\t n -> cglobal t (arr n)) (reverse ty) [0..]
 
@@ -322,7 +322,9 @@
           [x]   -> bind [cty| typename bool |] x
           _     -> INTERNAL_ERROR(error) "codegenOpenExp" "expected conditional predicate"
       --
-      return $ zipWith (\a b -> [cexp| $exp:p' ? $exp:a : $exp:b|]) t' e'
+      let cond ty a b   = addVar ty a >> addVar ty b >>
+                          return [cexp| $exp:p' ? $exp:a : $exp:b|]
+      sequence $ zipWith3 cond (expType t) t' e'
 
     -- Array indices and shapes
     --
@@ -361,8 +363,8 @@
     IndexScalar arr ix
       | OpenAcc (Avar a) <- arr ->
         let avar        = show (idxToInt a)
-            sh          = cvar ("sh"  ++ avar)
-            array x     = cvar ("arr" ++ avar ++ "_a" ++ show x)
+            sh          = cvar ("sh"   ++ avar)
+            array x     = cvar ("avar" ++ avar ++ "_a" ++ show x)
             elt         = accTypeTex arr
             n           = length elt
         in do
@@ -408,6 +410,8 @@
 addVar ty exp = case show exp of
   ('x':v:'_':'a':n) | [(v',[])] <- reads [v], [(n',[])] <- reads n
         -> use v' n' ty exp >> return True
+  ('v':n) | [(_ :: Int,[])] <- reads n
+        ->                     return True
   _     ->                     return False
 
 
diff --git a/Data/Array/Accelerate/CUDA/CodeGen/Base.hs b/Data/Array/Accelerate/CUDA/CodeGen/Base.hs
--- a/Data/Array/Accelerate/CUDA/CodeGen/Base.hs
+++ b/Data/Array/Accelerate/CUDA/CodeGen/Base.hs
@@ -116,8 +116,8 @@
        , String -> [InitGroup] )        -- const declarations and initialisation from index
 getters base arrElt expElt =
   let n                 = length arrElt
-      arr x             = "d_in" ++ shows base "_a" ++ show x
       arrParams         = zipWith (\t x -> [cparam| const $ty:(cptr t) $id:(arr x) |]) arrElt [n-1, n-2 .. 0]
+      arr x             = "arrIn" ++ shows base "_a" ++ show x
       expVars           = map (\(_,_,v) -> v) expElt
       expDecls          = map (\(_,t,v) -> [cdecl| $ty:t $id:(show v) ; |]) expElt
   in
@@ -139,7 +139,7 @@
        , String -> [Exp] -> [Stm])      -- store a value to the given index
 setters arrElt =
   let n                 = length arrElt
-      arrVars           = map (\x -> "d_out_a" ++ show x) [n-1, n-2 .. 0]
+      arrVars           = map (\x -> "arrOut_a" ++ show x) [n-1, n-2 .. 0]
       arrParams         = zipWith (\t x -> [cparam| $ty:(cptr t) $id:x |]) arrElt arrVars
       set ix a x        = [cstm| $id:a [$id:ix] = $exp:x; |]
   in
diff --git a/Data/Array/Accelerate/CUDA/CodeGen/IndexSpace.hs b/Data/Array/Accelerate/CUDA/CodeGen/IndexSpace.hs
--- a/Data/Array/Accelerate/CUDA/CodeGen/IndexSpace.hs
+++ b/Data/Array/Accelerate/CUDA/CodeGen/IndexSpace.hs
@@ -379,6 +379,7 @@
 --
 project :: Int -> String -> [Exp] -> [Stm]
 project n sh idx
+  | n   == 0    = [[cstm| $id:sh = 0; |]]
   | [e] <- idx  = [[cstm| $id:sh = $exp:e; |]]
   | otherwise   = zipWith (\i c -> [cstm| $id:sh . $id:('a':show c) = $exp:i; |]) idx [n-1,n-2..0]
 
diff --git a/Data/Array/Accelerate/CUDA/CodeGen/Monad.hs b/Data/Array/Accelerate/CUDA/CodeGen/Monad.hs
--- a/Data/Array/Accelerate/CUDA/CodeGen/Monad.hs
+++ b/Data/Array/Accelerate/CUDA/CodeGen/Monad.hs
@@ -79,10 +79,15 @@
 use :: Int -> Int -> Type -> Exp -> CGM ()
 use base prj ty var = modify variables (S.adjust (IM.insert prj (ty,var)) base)
 
--- Return the tuple components of a given variable that are actually used
+-- Return the tuple components of a given variable that are actually used. These
+-- in snoc-list ordering, i.e. with variable zero on the right.
 --
 subscripts :: Int -> CGM [(Int, Type, Exp)]
-subscripts base = map swizzle . IM.toList . flip S.index base <$> gets variables
+subscripts base
+  = reverse
+  . map swizzle
+  . IM.toList
+  . flip S.index base <$> gets variables
   where
     swizzle (i, (t,e)) = (i,t,e)
 
diff --git a/Data/Array/Accelerate/CUDA/CodeGen/PrefixSum.hs b/Data/Array/Accelerate/CUDA/CodeGen/PrefixSum.hs
--- a/Data/Array/Accelerate/CUDA/CodeGen/PrefixSum.hs
+++ b/Data/Array/Accelerate/CUDA/CodeGen/PrefixSum.hs
@@ -134,7 +134,7 @@
         $decls:decl2
 
         /*
-         * Read in previous result partial sum. We store the carry value in x0
+         * Read in previous result partial sum. We store the carry value in x2
          * and read new values from the input array into x1, since 'scanBlock'
          * will store its results into x1 on completion.
          */
@@ -154,19 +154,26 @@
             $stms:(x1 .=. getIn0 "j")
 
             if ( $exp:carry_in ) {
-                $decls:env
                 $stms:(x0 .=. x2)
+                $decls:env
                 $stms:(x1 .=. combine)
             }
 
             /*
-             * Store our input into shared memory and cooperatively scan
+             * Store our input into shared memory and perform a cooperative
+             * inclusive left scan.
              */
             $stms:(sdata "threadIdx.x" .=. x1)
             __syncthreads();
 
             $stms:(scanBlock dev elt Nothing (cvar "blockDim.x") sdata env combine)
 
+            /*
+             * Exclusive scans write the result of the previous thread to global
+             * memory. The first thread must reinstate the carry-in value which
+             * is the result of the last thread from the previous interval, or
+             * the carry-in/seed value for multi-block scans.
+             */
             if ( $exp:(cbool exclusive) ) {
                 if ( threadIdx.x == 0 ) {
                     $stms:(x1 .=. x2)
@@ -177,9 +184,9 @@
             $stms:(setOut "j" x1)
 
             /*
-             * Carry the final result from this block through x0. If this is the
-             * last section of the interval, this is the value to write out as
-             * the final (reduction) result.
+             * Carry the final result of this block through the set x2. If this
+             * is the final interval, this is the value to write out as the
+             * reduction result
              */
             if ( threadIdx.x == 0 ) {
                 const int last = min(interval_size - i, blockDim.x) - 1;
@@ -189,8 +196,7 @@
         }
 
         /*
-         * for exclusive scans, set the overall scan result and reapply the
-         * initial element at the boundaries of each interval
+         * for exclusive scans, set the overall scan result (reduction value)
          */
         if ( $exp:(cbool exclusive) && threadIdx.x == 0 && blockIdx.x == $id:lastBlock ) {
             $stms:(setSum .=. x2)
diff --git a/Data/Array/Accelerate/CUDA/CodeGen/Reduction.hs b/Data/Array/Accelerate/CUDA/CodeGen/Reduction.hs
--- a/Data/Array/Accelerate/CUDA/CodeGen/Reduction.hs
+++ b/Data/Array/Accelerate/CUDA/CodeGen/Reduction.hs
@@ -84,8 +84,8 @@
 
             for (i += gridSize; i < num_elements; i += gridSize)
             {
-                $decls:env
                 $stms:(x0 .=. getIn0 "i")
+                $decls:env
                 $stms:(x1 .=. combine)
             }
         }
@@ -124,6 +124,7 @@
           if (gridDim.x == 1) {
               $decls:env'
               $stms:(x0 .=. seed)
+              $decls:env
               $stms:(x1 .=. combine)
           }
           $stms:(setOut "blockIdx.x" x1)
@@ -171,7 +172,6 @@
         $decls:smem
         $decls:decl1
         $decls:decl0
-        $decls:env
 
         /*
          * If the intervals of an exclusive fold are empty, use all threads to
@@ -228,8 +228,8 @@
                  */
                 for (i += 2 * blockDim.x; i < end; i += blockDim.x)
                 {
-                    $decls:env
                     $stms:(x0 .=. getIn0 "i")
+                    $decls:env
                     $stms:(x1 .=. combine)
                 }
             }
@@ -252,11 +252,7 @@
              * exclusive reductions, we also combine with the seed element here.
              */
             if (threadIdx.x == 0)
-            {
-                $decls:final_decls
-                $stms:final_stms
-                $stms:(setOut "seg" x1)
-            }
+               $stm:(maybe inclusive_finish exclusive_finish mseed)
         }
     }
   |]
@@ -268,10 +264,16 @@
     (x1,   decl1)                       = locals "x1" elt
     (smem, sdata)                       = shared 0 Nothing [cexp| blockDim.x |] elt
     --
-    (final_decls, final_stms) =
-      case mseed of
-        Nothing                -> ([], [])
-        Just (CUExp env' seed) -> (env', concat [x0 .=. seed, x1 .=. combine])
+    inclusive_finish                    = [cstm| {
+        $stms:(setOut "seg" x1)
+    } |]
+    exclusive_finish (CUExp env' seed)  = [cstm| {
+        $decls:env'
+        $stms:(x0 .=. seed)
+        $decls:env
+        $stms:(x1 .=. combine)
+        $stms:(setOut "seg" x1)
+    } |]
     --
     mapseed (CUExp env' seed)           = [cstm|
       if (interval_size == 0)
@@ -411,8 +413,8 @@
                  */
                 for (i += 2 * warpSize; i < end; i += warpSize)
                 {
-                    $decls:env
                     $stms:(x0 .=. getIn0 "i")
+                    $decls:env
                     $stms:(x1 .=. combine)
                 }
             }
@@ -451,6 +453,7 @@
       if (num_elements > 0) {
           $decls:env'
           $stms:(x0 .=. seed)
+          $decls:env
           $stms:(x1 .=. combine)
       } else {
           $decls:env'
@@ -486,8 +489,8 @@
     reduce i
       | i > 1
       = [cstm| if ( $id:tid + $int:i < $id:n ) {
-                   $decls:env
                    $stms:(x0 .=. sdata ("threadIdx.x + " ++ show i))
+                   $decls:env
                    $stms:(x1 .=. combine)
                    $stms:(sdata "threadIdx.x" .=. x1)
                }
@@ -495,8 +498,8 @@
       --
       | otherwise
       = [cstm| if ( $id:tid + $int:i < $id:n ) {
-                   $decls:env
                    $stms:(x0 .=. sdata "threadIdx.x + 1")
+                   $decls:env
                    $stms:(x1 .=. combine)
                }
              |]
@@ -525,8 +528,8 @@
       | i > warpSize dev
       = [cstm| if ( $id:n > $int:i ) {
                    if ( threadIdx.x + $int:i < $id:n ) {
-                       $decls:env
                        $stms:(x0 .=. sdata ("threadIdx.x + " ++ show i))
+                       $decls:env
                        $stms:(x1 .=. combine)
                        $stms:(sdata "threadIdx.x" .=. x1)
                    }
diff --git a/Data/Array/Accelerate/CUDA/CodeGen/Stencil.hs b/Data/Array/Accelerate/CUDA/CodeGen/Stencil.hs
--- a/Data/Array/Accelerate/CUDA/CodeGen/Stencil.hs
+++ b/Data/Array/Accelerate/CUDA/CodeGen/Stencil.hs
@@ -184,8 +184,8 @@
   , \ix -> concatMap (get ix) subs )
   where
     n           = length stencil
-    sh          = "shIn"    ++ show  base
-    arr x       = "stencil" ++ shows base "_a" ++ show (x `mod` n)
+    sh          = "shIn" ++ show  base
+    arr x       = "arrIn" ++ shows base "_a" ++ show (x `mod` n)
     textures    = zipWith cglobal stencil (map arr [n-1, n-2 .. 0])
     --
     offsets     :: IArray.Array Int [Int]
diff --git a/Data/Array/Accelerate/CUDA/CodeGen/Type.hs b/Data/Array/Accelerate/CUDA/CodeGen/Type.hs
--- a/Data/Array/Accelerate/CUDA/CodeGen/Type.hs
+++ b/Data/Array/Accelerate/CUDA/CodeGen/Type.hs
@@ -35,7 +35,6 @@
 -- libraries
 import Language.C.Quote.CUDA
 import qualified Language.C                             as C
-import qualified Foreign.Storable                       as F
 
 
 #include "accelerate.h"
@@ -56,14 +55,6 @@
   | otherwise           = INTERNAL_ERROR(error) "accType" "non-scalar segment type"
 
 
---sizeOfAccTypes :: OpenAcc aenv (Sugar.Array dim e) -> [Int]
---sizeOfAccTypes = sizeOf' . Sugar.accType
---  where
---    sizeOf' :: TupleType a -> [Int]
---    sizeOf' UnitTuple           = []
---    sizeOf' x@(SingleTuple _)   = [Sugar.sizeOf x]
---    sizeOf' (PairTuple a b)     = sizeOf' a ++ sizeOf' b
-
 eltType :: Sugar.Elt a => a {- dummy -} -> [C.Type]
 eltType =  codegenTupleType . Sugar.eltType
 
@@ -95,14 +86,14 @@
 codegenNumType (FloatingNumType ty) = codegenFloatingType ty
 
 codegenIntegralType :: IntegralType a -> C.Type
-codegenIntegralType (TypeInt8    _) = typename "int8_t"
-codegenIntegralType (TypeInt16   _) = typename "int16_t"
-codegenIntegralType (TypeInt32   _) = typename "int32_t"
-codegenIntegralType (TypeInt64   _) = typename "int64_t"
-codegenIntegralType (TypeWord8   _) = typename "uint8_t"
-codegenIntegralType (TypeWord16  _) = typename "uint16_t"
-codegenIntegralType (TypeWord32  _) = typename "uint32_t"
-codegenIntegralType (TypeWord64  _) = typename "uint64_t"
+codegenIntegralType (TypeInt8    _) = typename "Int8"
+codegenIntegralType (TypeInt16   _) = typename "Int16"
+codegenIntegralType (TypeInt32   _) = typename "Int32"
+codegenIntegralType (TypeInt64   _) = typename "Int64"
+codegenIntegralType (TypeWord8   _) = typename "Word8"
+codegenIntegralType (TypeWord16  _) = typename "Word16"
+codegenIntegralType (TypeWord32  _) = typename "Word32"
+codegenIntegralType (TypeWord64  _) = typename "Word64"
 codegenIntegralType (TypeCShort  _) = [cty|short|]
 codegenIntegralType (TypeCUShort _) = [cty|unsigned short|]
 codegenIntegralType (TypeCInt    _) = [cty|int|]
@@ -111,19 +102,16 @@
 codegenIntegralType (TypeCULong  _) = [cty|unsigned long int|]
 codegenIntegralType (TypeCLLong  _) = [cty|long long int|]
 codegenIntegralType (TypeCULLong _) = [cty|unsigned long long int|]
-
-codegenIntegralType (TypeInt     _) =
-  case F.sizeOf (undefined::Int) of
-       4 -> typename "int32_t"
-       8 -> typename "int64_t"
-       _ -> error "we can never get here"
-
-codegenIntegralType (TypeWord    _) =
-  case F.sizeOf (undefined::Int) of
-       4 -> typename "uint32_t"
-       8 -> typename "uint64_t"
-       _ -> error "we can never get here"
-
+#if   SIZEOF_HSINT == 4
+codegenIntegralType (TypeInt     _) = typename "Int32"
+#elif SIZEOF_HSINT == 8
+codegenIntegralType (TypeInt     _) = typename "Int64"
+#endif
+#if   SIZEOF_HSINT == 4
+codegenIntegralType (TypeWord    _) = typename "Word32"
+#elif SIZEOF_HSINT == 8
+codegenIntegralType (TypeWord    _) = typename "Word64"
+#endif
 
 codegenFloatingType :: FloatingType a -> C.Type
 codegenFloatingType (TypeFloat   _) = [cty|float|]
@@ -132,8 +120,10 @@
 codegenFloatingType (TypeCDouble _) = [cty|double|]
 
 codegenNonNumType :: NonNumType a -> C.Type
-codegenNonNumType (TypeBool   _) = error "codegenNonNum :: Bool"
-codegenNonNumType (TypeChar   _) = error "codegenNonNum :: Char"
+codegenNonNumType (TypeBool   _) = typename "Word8"
+#if   SIZEOF_HSCHAR == 4
+codegenNonNumType (TypeChar   _) = typename "Word32"
+#endif
 codegenNonNumType (TypeCChar  _) = [cty|char|]
 codegenNonNumType (TypeCSChar _) = [cty|signed char|]
 codegenNonNumType (TypeCUChar _) = [cty|unsigned char|]
@@ -178,18 +168,16 @@
 codegenIntegralTex (TypeCULong  _) = typename "TexCULong"
 codegenIntegralTex (TypeCLLong  _) = typename "TexCLLong"
 codegenIntegralTex (TypeCULLong _) = typename "TexCULLong"
-
-codegenIntegralTex (TypeInt     _) =
-  case F.sizeOf (undefined::Int) of
-       4 -> typename "TexInt32"
-       8 -> typename "TexInt64"
-       _ -> error "we can never get here"
-
-codegenIntegralTex (TypeWord    _) =
-  case F.sizeOf (undefined::Word) of
-       4 -> typename "TexWord32"
-       8 -> typename "TexWord64"
-       _ -> error "we can never get here"
+#if   SIZEOF_HSINT == 4
+codegenIntegralTex (TypeInt     _) = typename "TexInt32"
+#elif SIZEOF_HSINT == 8
+codegenIntegralTex (TypeInt     _) = typename "TexInt64"
+#endif
+#if   SIZEOF_HSINT == 4
+codegenIntegralTex (TypeWord    _) = typename "TexWord32"
+#elif SIZEOF_HSINT == 8
+codegenIntegralTex (TypeWord    _) = typename "TexWord64"
+#endif
 
 
 codegenFloatingTex :: FloatingType a -> C.Type
@@ -199,13 +187,11 @@
 codegenFloatingTex (TypeCDouble _) = typename "TexCDouble"
 
 
--- TLM 2010-06-29:
---   Bool and Char can be implemented once the array types in
---   Data.Array.Accelerate.[CUDA.]Array.Data are made concrete.
---
 codegenNonNumTex :: NonNumType a -> C.Type
-codegenNonNumTex (TypeBool   _) = error "codegenNonNumTex :: Bool"
-codegenNonNumTex (TypeChar   _) = error "codegenNonNumTex :: Char"
+codegenNonNumTex (TypeBool   _) = typename "TexWord8"
+#if   SIZEOF_HSCHAR == 4
+codegenNonNumTex (TypeChar   _) = typename "TexWord32"
+#endif
 codegenNonNumTex (TypeCChar  _) = typename "TexCChar"
 codegenNonNumTex (TypeCSChar _) = typename "TexCSChar"
 codegenNonNumTex (TypeCUChar _) = typename "TexCUChar"
diff --git a/Data/Array/Accelerate/CUDA/Compile.hs b/Data/Array/Accelerate/CUDA/Compile.hs
--- a/Data/Array/Accelerate/CUDA/Compile.hs
+++ b/Data/Array/Accelerate/CUDA/Compile.hs
@@ -24,11 +24,12 @@
 import Data.Array.Accelerate.Tuple
 
 import Data.Array.Accelerate.CUDA.AST
-import Data.Array.Accelerate.CUDA.FullList              as FL
 import Data.Array.Accelerate.CUDA.State
 import Data.Array.Accelerate.CUDA.CodeGen
 import Data.Array.Accelerate.CUDA.Array.Sugar
 import Data.Array.Accelerate.CUDA.Analysis.Launch
+import Data.Array.Accelerate.CUDA.FullList              as FL
+import Data.Array.Accelerate.CUDA.Persistent            as KT
 import qualified Data.Array.Accelerate.CUDA.Debug       as D
 
 -- libraries
@@ -45,7 +46,6 @@
 import Data.List
 import Data.Maybe
 import Data.Monoid
-import Foreign.Storable
 import System.Directory
 import System.Exit                                      ( ExitCode(..) )
 import System.FilePath
@@ -55,7 +55,6 @@
 import Text.PrettyPrint.Mainland                        ( RDoc(..), ppr, renderCompact )
 import Data.ByteString.Internal                         ( w2c )
 import qualified Data.HashSet                           as Set
-import qualified Data.HashTable.IO                      as HT
 import qualified Data.ByteString                        as B
 import qualified Data.ByteString.Lazy                   as L
 import qualified Foreign.CUDA.Driver                    as CUDA
@@ -300,53 +299,57 @@
       --
       -- make sure kernel/stats are printed together
       --
-      message   $ intercalate "\n" [msg1, "cc: " ++ msg2]
+      message   $ intercalate "\n" [msg1, "     ... " ++ msg2]
 
 
 -- Link a compiled binary and update the associated kernel entry in the hash
 -- table. This may entail waiting for the external compilation process to
--- complete. If successfully, the temporary files are removed.
+-- complete. If successful, the temporary files are removed.
 --
 link :: KernelTable -> KernelKey -> IO CUDA.Module
 link table key =
   let intErr = INTERNAL_ERROR(error) "link" "missing kernel entry"
   in do
-    ctx                         <- CUDA.get
-    (KernelEntry cufile stat)   <- fromMaybe intErr `fmap` HT.lookup table key
-    case stat of
-      Right (KernelObject bin active)
-        | Just mdl <- FL.lookup ctx active      -> return mdl
-        | otherwise                             -> do
-            message "re-linking module for current context"
-            mdl         <- CUDA.loadData bin
-            let obj     =  KernelObject bin (FL.cons ctx mdl active)
-            HT.insert table key (KernelEntry cufile (Right obj))
-            return mdl
-      --
-      Left  pid         -> do
-        -- wait for compiler to finish and load binary object
+    ctx         <- CUDA.get
+    entry       <- fromMaybe intErr `fmap` KT.lookup table key
+    case entry of
+      CompileProcess cufile pid -> do
+        -- Wait for the compiler to finish and load the binary object into the
+        -- current context
         --
         message "waiting for nvcc..."
         waitFor pid
-        bin     <- B.readFile (replaceExtension cufile ".cubin")
-        mdl     <- CUDA.loadData bin
-        let obj =  KernelObject bin (FL.singleton ctx mdl)
+        let cubin       =  replaceExtension cufile ".cubin"
+        bin             <- B.readFile cubin
+        mdl             <- CUDA.loadData bin
 
-#ifndef ACCELERATE_CUDA_PERSISTENT_CACHE
-        -- remove build products
+        -- Update hash tables and stash the binary object into the persistent
+        -- cache
         --
+        KT.insert table key $! KernelObject bin (FL.singleton ctx mdl)
+        KT.persist cubin key
+
+        -- Remove temporary build products
+        --
         removeFile      cufile
-        removeFile      (replaceExtension cufile ".cubin")
         removeDirectory (dropFileName cufile)
           `catch` \(_ :: IOError) -> return ()          -- directory not empty
-#endif
 
-        -- update hash table
-        --
-        HT.insert table key (KernelEntry cufile (Right obj))
         return mdl
 
+      -- If we get a real object back, then this will already be in the
+      -- persistent cache, since either it was just read in from there, or we
+      -- had to generate new code and the link step above has added it.
+      --
+      KernelObject bin active
+        | Just mdl <- FL.lookup ctx active      -> return mdl
+        | otherwise                             -> do
+            message "re-linking module for current context"
+            mdl                 <- CUDA.loadData bin
+            KT.insert table key $! KernelObject bin (FL.cons ctx mdl active)
+            return mdl
 
+
 -- Generate and compile code for a single open array expression
 --
 compile :: KernelTable
@@ -355,17 +358,17 @@
         -> AccBindings aenv
         -> CIO (String, KernelKey)
 compile table dev acc fvar = do
-  exists        <- isJust `fmap` liftIO (HT.lookup table key)
+  exists        <- isJust `fmap` liftIO (KT.lookup table key)
   unless exists $ do
     message     $  unlines [ show key, map w2c (L.unpack code) ]
     nvcc        <- fromMaybe (error "nvcc: command not found") <$> liftIO (findExecutable "nvcc")
-    (file,hdl)  <- openOutputFile "dragon.cu"   -- rawr!
+    (file,hdl)  <- openTemporaryFile "dragon.cu"   -- rawr!
     flags       <- compileFlags file
     (_,_,_,pid) <- liftIO $ do
       L.hPut hdl code                 `finally`     hClose hdl
       createProcess (proc nvcc flags) `onException` removeFile file
     --
-    liftIO $ HT.insert table key (KernelEntry file (Left pid))
+    liftIO $ KT.insert table key (CompileProcess file pid)
   --
   return (entry, key)
   where
@@ -405,28 +408,25 @@
     , "-arch=sm_" ++ show (round (arch * 10) :: Int)
     , "-cubin"
     , "-o", cufile `replaceExtension` "cubin"
-    , if D.mode D.verbose then ""   else "--disable-warnings"
-    , if D.mode D.debug   then "-G" else "-O3"
+    , if D.mode D.dump_cc  then ""   else "--disable-warnings"
+    , if D.mode D.debug_cc then "-G" else "-O3"
     , machine
     , cufile ]
   where
-    wordSize                    = sizeOf (undefined::Int)
-    machine | wordSize == 4     = "-m32"
-            | wordSize == 8     = "-m64"
-            | otherwise         = error "recreational scolding?"
+#if SIZEOF_HSINT == 4
+    machine     = "-m32"
+#elif SIZEOF_HSINT == 8
+    machine     = "-m64"
+#endif
 
 
 -- Open a unique file in the temporary directory used for compilation
 -- by-products. The directory will be created if it does not exist.
 --
-openOutputFile :: String -> CIO (FilePath, Handle)
-openOutputFile template = liftIO $ do
-#ifdef ACCELERATE_CUDA_PERSISTENT_CACHE
-  dir <- (</>) <$> getDataDir            <*> pure "cache"
-#else
+openTemporaryFile :: String -> CIO (FilePath, Handle)
+openTemporaryFile template = liftIO $ do
   pid <- getProcessID
   dir <- (</>) <$> getTemporaryDirectory <*> pure ("accelerate-cuda-" ++ show pid)
-#endif
   createDirectoryIfMissing True dir
   openTempFile dir template
 
@@ -440,9 +440,9 @@
 
 {-# INLINE message #-}
 message :: MonadIO m => String -> m ()
-message msg = trace ("cc: " ++ msg) $ return ()
+message msg = trace msg $ return ()
 
 {-# INLINE trace #-}
 trace :: MonadIO m => String -> m a -> m a
-trace msg next = D.message D.dump_cc msg >> next
+trace msg next = D.message D.dump_cc ("cc: " ++ msg) >> next
 
diff --git a/Data/Array/Accelerate/CUDA/Debug.hs b/Data/Array/Accelerate/CUDA/Debug.hs
--- a/Data/Array/Accelerate/CUDA/Debug.hs
+++ b/Data/Array/Accelerate/CUDA/Debug.hs
@@ -20,8 +20,8 @@
   showFFloatSIBase,
 
   message, event, when, mode,
-  verbose, debug,
-  dump_gc, dump_cc, dump_exec,
+  verbose, flush_cache,
+  dump_gc, dump_cc, debug_cc, dump_exec,
 
 ) where
 
@@ -30,6 +30,7 @@
 import Data.Label
 import Data.IORef
 import Control.Monad.IO.Class
+import System.CPUTime
 import System.IO.Unsafe
 import System.Environment
 import System.Console.GetOpt
@@ -67,13 +68,14 @@
 data Flags = Flags
   {
     -- phase control
-    _dump_gc    :: !Bool        -- garbage collection & memory management
-  , _dump_cc    :: !Bool        -- compilation & linking
-  , _dump_exec  :: !Bool        -- kernel execution
+    _dump_gc            :: !Bool        -- garbage collection & memory management
+  , _dump_cc            :: !Bool        -- compilation & linking
+  , _debug_cc           :: !Bool        -- compile device code with debug symbols
+  , _dump_exec          :: !Bool        -- kernel execution
 
     -- general options
-  , _verbose    :: !Bool        -- additional status messages
-  , _debug      :: !Bool        -- generate device code suitable for debugging
+  , _verbose            :: !Bool        -- additional status messages
+  , _flush_cache        :: !Bool        -- delete the persistent cache directory
   }
 
 $(mkLabels [''Flags])
@@ -82,15 +84,16 @@
 flags =
   [ Option [] ["ddump-gc"]      (NoArg (set dump_gc True))      "print device memory management trace"
   , Option [] ["ddump-cc"]      (NoArg (set dump_cc True))      "print generated code and compilation information"
+  , Option [] ["ddebug-cc"]     (NoArg (set debug_cc True))     "generate debug information for device code"
   , Option [] ["ddump-exec"]    (NoArg (set dump_exec True))    "print kernel execution trace"
   , Option [] ["dverbose"]      (NoArg (set verbose True))      "print additional information"
-  , Option [] ["ddebug"]        (NoArg (set debug True))        "generate debug information for device code"
+  , Option [] ["fflush-cache"]  (NoArg (set flush_cache True))  "delete the persistent cache directory"
   ]
 
 initialise :: IO Flags
 initialise = parse `fmap` getArgs
   where
-    defaults      = Flags False False False False False
+    defaults      = Flags False False False False False False
     parse         = foldl parse1 defaults
     parse1 opts x = case filter (\(Option _ [f] _ _) -> x `isPrefixOf` ('-':f)) flags of
                       [Option _ _ (NoArg go) _] -> go opts
@@ -113,7 +116,11 @@
 {-# INLINE message #-}
 message :: MonadIO m => (Flags :-> Bool) -> String -> m ()
 #ifdef ACCELERATE_DEBUG
-message f str = when f (liftIO $ traceIO str)
+message f str
+  = when f . liftIO
+  $ do psec     <- getCPUTime
+       let sec   = fromIntegral psec * 1E-12 :: Double
+       traceIO   $ showFFloat (Just 2) sec (':':str)
 #else
 message _ _   = return ()
 #endif
diff --git a/Data/Array/Accelerate/CUDA/Execute.hs b/Data/Array/Accelerate/CUDA/Execute.hs
--- a/Data/Array/Accelerate/CUDA/Execute.hs
+++ b/Data/Array/Accelerate/CUDA/Execute.hs
@@ -641,9 +641,9 @@
     -> Val aenv
     -> Array dim a
     -> CIO (Array dim b)
-stencilOp kernel@(Kernel _ mdl _ _ _) bindings aenv in0@(Array sh0 _) = do
+stencilOp kernel bindings aenv in0@(Array sh0 _) = do
   res@(Array _ out)     <- allocateArray (toElt sh0)
-  bindStencil 0 mdl in0
+  bindAcc 0 kernel in0
   execute kernel bindings aenv (size sh0)
     (((), out)
         , toElt sh0 :: dim)
@@ -657,10 +657,10 @@
     -> Array dim a
     -> Array dim b
     -> CIO (Array dim c)
-stencil2Op kernel@(Kernel _ mdl _ _ _) bindings aenv in1@(Array sh1 _) in0@(Array sh0 _) = do
+stencil2Op kernel bindings aenv in1@(Array sh1 _) in0@(Array sh0 _) = do
   res@(Array sh out)    <- allocateArray $ toElt (sh1 `intersect` sh0)
-  bindStencil 1 mdl in1
-  bindStencil 0 mdl in0
+  bindAcc 1 kernel in1
+  bindAcc 0 kernel in0
   execute kernel bindings aenv (size sh)
     (((((), out)
           , toElt sh  :: dim)
@@ -756,40 +756,42 @@
 -- Array references in scalar code
 -- -------------------------------
 
-bindLifted :: CUDA.Module -> Val aenv -> AccBindings aenv -> CIO ()
-bindLifted mdl aenv (AccBindings vars) = mapM_ (bindAcc mdl aenv) (Set.toList vars)
-
+-- All CUDA devices have between 6-8KB of read-only texture memory per
+-- multiprocessor. Since all arrays in Accelerate are immutable, we can always
+-- access input arrays through the texture cache to reduce global memory demand
+-- when accesses do not follow the regular patterns required for coalescing.
+--
+-- This is great for older 1.x series devices, but compute 2.x devices have a
+-- dedicated 768KB L2 cache, as well as a configurable L1 cache of 16/48KB
+-- (combined with shared memory). What we really want is for the code generator
+-- to pass all inputs either as textures or global arrays, depending on what
+-- device we are currently targeting.
+--
 
-bindAcc
-    :: CUDA.Module
-    -> Val aenv
-    -> ArrayVar aenv
-    -> CIO ()
-bindAcc mdl aenv (ArrayVar idx) =
-  let idx'        = show $ idxToInt idx
-      Array sh ad = prj idx aenv
-      --
-      bindDim = liftIO $
-        CUDA.getPtr mdl ("sh" ++ idx') >>=
-        CUDA.pokeListArray (convertSh sh) . fst
-      --
-      arr n   = "arr" ++ idx' ++ "_a" ++ show (n::Int)
-      tex     = CUDA.getTex mdl . arr
-      bindTex = marshalTextureData ad (size sh) =<< liftIO (sequence' $ map tex [0..])
+bindAcc :: Int -> AccKernel a -> Array dim a' -> CIO ()
+bindAcc base (Kernel _ mdl _ _ _) (Array sh ad) =
+  let arr n     = "arrIn" ++ show base ++ "_a" ++ show (n::Int)
+      tex       = CUDA.getTex mdl . arr
   in
-  bindDim >> bindTex
+  marshalTextureData ad (size sh) =<< liftIO (sequence' $ map tex [0..])
 
 
-bindStencil
-    :: Int
-    -> CUDA.Module
-    -> Array dim e
-    -> CIO ()
-bindStencil s mdl (Array sh ad) =
-  let sten n = "stencil" ++ show s ++ "_a" ++ show (n::Int)
-      tex    = CUDA.getTex mdl . sten
-  in
-  marshalTextureData ad (size sh) =<< liftIO (sequence' $ map tex [0..])
+bindAccEnv :: AccKernel a -> Val aenv -> AccBindings aenv -> CIO ()
+bindAccEnv (Kernel _ mdl _ _ _) aenv (AccBindings vars) = mapM_ bindAvar (Set.toList vars)
+  where
+    bindAvar (ArrayVar idx) =
+      let idx'          = show $ idxToInt idx
+          Array sh ad   = prj idx aenv
+          --
+          bindDim       = liftIO $
+            CUDA.getPtr mdl ("sh" ++ idx') >>=
+            CUDA.pokeListArray (convertSh sh) . fst
+          --
+          arr n         = "avar" ++ idx' ++ "_a" ++ show (n::Int)
+          tex           = CUDA.getTex mdl . arr
+          bindTex       = marshalTextureData ad (size sh) =<< liftIO (sequence' $ map tex [0..])
+      in
+      bindDim >> bindTex
 
 
 -- Kernel execution
@@ -864,8 +866,8 @@
         -> Int
         -> args
         -> CIO ()
-execute kernel@(Kernel _ !mdl !_ !_ !_) !bindings !aenv !n !args = do
-  bindLifted mdl aenv bindings
+execute kernel@(Kernel _ !_ !_ !_ !_) !bindings !aenv !n !args = do
+  bindAccEnv kernel aenv bindings
   launch kernel (configure kernel n) args
 
 
diff --git a/Data/Array/Accelerate/CUDA/Persistent.hs b/Data/Array/Accelerate/CUDA/Persistent.hs
new file mode 100644
--- /dev/null
+++ b/Data/Array/Accelerate/CUDA/Persistent.hs
@@ -0,0 +1,274 @@
+{-# LANGUAGE BangPatterns        #-}
+{-# LANGUAGE ScopedTypeVariables #-}
+-- |
+-- Module      : Data.Array.Accelerate.CUDA.Persistent
+-- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
+--               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- License     : BSD3
+--
+-- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
+-- Stability   : experimental
+-- Portability : non-partable (GHC extensions)
+--
+
+module Data.Array.Accelerate.CUDA.Persistent (
+
+  KernelTable, KernelKey, KernelEntry(..),
+  new, lookup, insert, persist
+
+) where
+
+-- friends
+import Data.Array.Accelerate.CUDA.FullList              ( FullList )
+import qualified Data.Array.Accelerate.CUDA.Debug       as D
+import qualified Data.Array.Accelerate.CUDA.FullList    as FL
+
+-- libraries
+import Prelude                                          hiding ( lookup, catch )
+import Data.Char
+import System.IO
+import System.FilePath
+import System.Directory
+import System.Process                                   ( ProcessHandle )
+import Control.Exception
+import Control.Applicative
+import Control.Monad.Trans
+import Data.Binary
+import Data.Binary.Get
+import Data.ByteString                                  ( ByteString )
+import Data.ByteString.Internal                         ( w2c )
+import qualified Data.ByteString                        as B
+import qualified Data.ByteString.Lazy                   as L
+import qualified Data.HashTable.IO                      as HT
+
+import qualified Foreign.CUDA.Driver                    as CUDA
+import qualified Foreign.CUDA.Analysis                  as CUDA
+
+import Paths_accelerate_cuda
+
+
+-- Interface -------------------------------------------------------------------
+-- ---------                                                                  --
+
+data KernelTable = KT {-# UNPACK #-} !ProgramCache      -- first level cache
+                      {-# UNPACK #-} !PersistentCache   -- second level cache
+
+new :: IO KernelTable
+new = do
+  cacheDir <- cacheDirectory
+  createDirectoryIfMissing True cacheDir
+  --
+  local         <- HT.new
+  persistent    <- restore (cacheDir </> "persistent.db")
+  --
+  return        $! KT local persistent
+
+
+-- Lookup a kernel through the two-level cache system. If the kernel is found in
+-- the persistent cache, it is loaded and linked into the current context.
+--
+lookup :: KernelTable -> KernelKey -> IO (Maybe KernelEntry)
+lookup (KT kt pt) key = do
+  -- First check the local cache. If we get a hit, this could be:
+  --   a) currently compiling
+  --   b) compiled, but not linked into the current context
+  --   c) compiled & linked
+  --
+  v1    <- HT.lookup kt key
+  case v1 of
+    Just _      -> return v1
+    Nothing     -> do
+
+    -- Check the persistent cache. If found, read in the associated object file
+    -- and link it into the current context. Also add to the first-level cache.
+    --
+    -- TLM: maybe we should change KernelObject to hold a possibly empty list,
+    --      so we don't have to mess with the CUDA context here.
+    --
+    v2  <- HT.lookup pt key
+    case v2 of
+      Nothing   -> return Nothing
+      Just ()   -> do
+        message "found/persistent"
+        cubin   <- (</>) <$> cacheDirectory <*> pure (cacheFilePath key)
+        ctx     <- CUDA.get
+        bin     <- B.readFile cubin
+        mdl     <- CUDA.loadData bin
+        let obj  = KernelObject bin (FL.singleton ctx mdl)
+        HT.insert kt key obj
+        return  $! Just obj
+
+
+-- Insert a key/value pair into the first-level cache. This does not add the
+-- entry to the persistent database.
+--
+-- TLM: Also add to the persistent cache, or return a boolean as to whether it
+--      exists there already? Would require updating that hash table as new
+--      entries are added, which the functions currently do not do.
+--
+insert :: KernelTable -> KernelKey -> KernelEntry -> IO ()
+insert (KT kt _) key val = HT.insert kt key val
+
+
+-- Local cache -----------------------------------------------------------------
+-- -----------                                                                --
+--
+-- Kernel code that has been generated and linked into the currently running
+-- program.
+
+-- An exact association between an accelerate computation and its
+-- implementation, which is either a reference to the external compiler (nvcc)
+-- or the resulting binary module.
+--
+-- Note that since we now support running in multiple contexts, we also need to
+-- keep track of
+--   a) the compute architecture the code was compiled for
+--   b) which contexts have linked the code
+--
+-- We aren't concerned with true (typed) equality of an OpenAcc expression,
+-- since we largely want to disregard the array environment; we really only want
+-- to assert the type and index of those variables that are accessed by the
+-- computation and no more, but we can not do that. Instead, this is keyed to
+-- the generated kernel code.
+--
+type ProgramCache = HT.BasicHashTable KernelKey KernelEntry
+
+type KernelKey    = (CUDA.Compute, ByteString)
+data KernelEntry
+  -- A currently compiling external process. We record the process ID and the
+  -- path of the .cu file being compiled
+  --
+  = CompileProcess !FilePath !ProcessHandle
+
+  -- The raw compiled data, and the list of contexts that the object has already
+  -- been linked into. If we locate this entry in the ProgramCache, it may have
+  -- been inserted by an alternate but compatible device context, so just
+  -- re-link into the current context.
+  --
+  | KernelObject {-# UNPACK #-} !ByteString
+                 {-# UNPACK #-} !(FullList CUDA.Context CUDA.Module)
+
+
+-- Persistent cache ------------------------------------------------------------
+-- ----------------                                                           --
+--
+-- Stash compiled code into the user's home directory so that they are available
+-- across separate runs of the program.
+--
+-- TLM: we don't have any migration or versioning policy here, so cache files
+--      will be kept around indefinitely. This can easily clutter the cache by
+--      generating many similar kernels that differ only by, for example, an
+--      embedded constant value.
+
+type PersistentCache = HT.BasicHashTable KernelKey ()
+
+
+-- The root directory of where the various persistent cache files live; the
+-- database and each individual binary object.
+--
+-- TLM: Is this writeable, even at a 'cabal instal --global'? Maybe we should
+--      specifically choose something in the user's home directory.
+--
+cacheDirectory :: IO FilePath
+cacheDirectory = do
+  dir   <- canonicalizePath =<< getDataDir
+  return $ dir </> "cache"
+
+-- A relative path to be appended to (presumably) 'cacheDirectory'.
+--
+cacheFilePath :: KernelKey -> FilePath
+cacheFilePath (cap, key) = show cap </> foldl (flip (mangle . w2c)) ".cubin" (B.unpack key)
+  where
+    -- TODO: complete z-encoding? see: compiler/utils/Encoding.hs
+    --
+    mangle '\\'   = ("zr" ++)
+    mangle '/'    = ("zs" ++)
+    mangle c      = showLitChar c
+
+
+-- The default Binary instance for lists is (necessarily) spine and value
+-- strict for efficiency. For us it is better if we just lazily consume elements
+-- and add them directly to the hash table so they can be collected as we go.
+--
+{-# INLINE getMany #-}
+getMany :: Binary a => Int -> Get [a]
+getMany n = go n []
+  where
+    go 0 xs = return xs
+    go i xs = do
+      x <- get
+      go (i-1) (x:xs)
+
+
+-- Load the entire persistent cache index file. If it does not exist, an empty
+-- file is created, so that 'persist' can always append elements.
+--
+restore :: FilePath -> IO PersistentCache
+restore db = do
+  D.when D.flush_cache $ do
+    message $ "deleting persistent cache"
+    cacheDir <- cacheDirectory
+    removeDirectoryRecursive cacheDir
+    createDirectoryIfMissing True cacheDir
+  --
+  exists <- doesFileExist db
+  case exists of
+    False       -> encodeFile db (0::Int) >> HT.new
+    True        -> do
+      store         <- L.readFile db
+      let (n,rest,_) = runGetState get store 0
+      pt            <- HT.newSized n
+      --
+      let go []      = return ()
+          go (!k:xs) = HT.insert pt k () >> go xs
+      --
+      message $ "persist/restore: " ++ shows n " entries"
+      go (runGet (getMany n) rest)
+      pt `seq` return pt
+
+
+-- Append a single value to the persistent cache.
+--
+-- This moves the compiled object file (first argument) to the appropriate
+-- location, and updates the database on disk.
+--
+persist :: FilePath -> KernelKey -> IO ()
+persist cubin key = do
+  cacheDir <- cacheDirectory
+  let db        = cacheDir </> "persistent.db"
+      cacheFile = cacheDir </> cacheFilePath key
+  --
+  message $ "persist/save: " ++ cacheFile
+  createDirectoryIfMissing True (dropFileName cacheFile)
+  renameFile cubin cacheFile
+    -- If the temporary and cache directories are on different disks, we must
+    -- copy the file instead. Unsupported operation: (Cross-device link)
+    --
+    `catch` \(_ :: IOError) -> do
+      copyFile cubin cacheFile
+      removeFile cubin
+  --
+  withBinaryFile db ReadWriteMode $ \h -> do
+    -- The file opens with the cursor at the beginning of the file
+    --
+    n <- runGet (get :: Get Int) `fmap` L.hGet h 8
+    hSeek h AbsoluteSeek 0
+    L.hPut h $ encode (n+1)
+
+    -- Append the new entry to the end of file
+    --
+    hSeek h SeekFromEnd 0
+    L.hPut h (encode key)
+
+
+-- Debug
+-- -----
+
+{-# INLINE message #-}
+message :: MonadIO m => String -> m ()
+message msg = trace msg $ return ()
+
+{-# INLINE trace #-}
+trace :: MonadIO m => String -> m a -> m a
+trace msg next = D.message D.dump_cc ("cc: " ++ msg) >> next
+
diff --git a/Data/Array/Accelerate/CUDA/State.hs b/Data/Array/Accelerate/CUDA/State.hs
--- a/Data/Array/Accelerate/CUDA/State.hs
+++ b/Data/Array/Accelerate/CUDA/State.hs
@@ -2,7 +2,7 @@
 {-# LANGUAGE TemplateHaskell #-}
 {-# LANGUAGE TupleSections   #-}
 {-# LANGUAGE TypeOperators   #-}
-{-# OPTIONS_GHC -fno-warn-orphans #-}   -- CUDA.Context
+{-# OPTIONS_GHC -fno-warn-orphans #-}   -- Eq CUDA.Context
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.State
 -- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
@@ -20,80 +20,40 @@
 
 module Data.Array.Accelerate.CUDA.State (
 
-  -- Types
-  CIO, KernelTable, KernelKey, KernelEntry(KernelEntry), KernelObject(KernelObject),
-
   -- Evaluating computations
-  evalCUDA, defaultContext, deviceProps,
-  memoryTable, kernelTable, kernelName, kernelStatus
+  CIO, evalCUDA,
 
+  -- Querying execution state
+  defaultContext, deviceProps, activeContext, kernelTable, memoryTable
+
 ) where
 
 -- friends
-import Data.Array.Accelerate.CUDA.FullList              ( FullList )
 import Data.Array.Accelerate.CUDA.Debug                 ( message, verbose, dump_gc, showFFloatSIBase )
+import Data.Array.Accelerate.CUDA.Persistent            as KT
 import Data.Array.Accelerate.CUDA.Array.Table           as MT
 import Data.Array.Accelerate.CUDA.Analysis.Device
 
 -- library
 import Data.Label
 import Control.Exception
-import Data.ByteString                                  ( ByteString )
-import Control.Concurrent.MVar                          ( MVar, newMVar )
+import Control.Concurrent                               ( forkIO, threadDelay )
 import Control.Monad.State.Strict                       ( StateT(..), evalStateT )
-import System.Process                                   ( ProcessHandle )
 import System.Mem                                       ( performGC )
-import System.Mem.Weak                                  ( addFinalizer )
-import System.IO.Unsafe
+import System.Mem.Weak                                  ( mkWeakPtr, addFinalizer )
+import System.IO.Unsafe                                 ( unsafePerformIO )
 import Text.PrettyPrint
 import qualified Foreign.CUDA.Driver                    as CUDA hiding ( device )
 import qualified Foreign.CUDA.Driver.Context            as CUDA
-import qualified Foreign.CUDA.Analysis                  as CUDA
-import qualified Data.HashTable.IO                      as HT
 
-#ifdef ACCELERATE_CUDA_PERSISTENT_CACHE
-import Data.Binary                                      ( encodeFile, decodeFile )
-import Control.Arrow                                    ( second )
-import Paths_accelerate                                 ( getDataDir )
-#endif
 
-
--- An exact association between an accelerate computation and its
--- implementation, which is either a reference to the external compiler (nvcc)
--- or the resulting binary module.
---
--- Note that since we now support running in multiple contexts, we also need to
--- keep track of
---   a) the compute architecture the code was compiled for
---   b) which contexts have linked the code
---
--- We aren't concerned with true (typed) equality of an OpenAcc expression,
--- since we largely want to disregard the array environment; we really only want
--- to assert the type and index of those variables that are accessed by the
--- computation and no more, but we can not do that. Instead, this is keyed to
--- the generated kernel code.
---
-type KernelTable = HT.BasicHashTable KernelKey KernelEntry
-
-type KernelKey   = (CUDA.Compute, ByteString)
-data KernelEntry = KernelEntry
-  {
-    _kernelName         :: !FilePath,
-    _kernelStatus       :: !(Either ProcessHandle KernelObject)
-  }
-
-data KernelObject = KernelObject
-  {
-    _binaryData         :: !ByteString,
-    _activeContexts     :: {-# UNPACK #-} !(FullList CUDA.Context CUDA.Module)
-  }
-
 -- The state token for CUDA accelerated array operations
 --
 type CIO        = StateT CUDAState IO
 data CUDAState  = CUDAState
   {
     _deviceProps        :: !CUDA.DeviceProperties,
+    _activeContext      :: {-# UNPACK #-} !Context,
     _kernelTable        :: {-# UNPACK #-} !KernelTable,
     _memoryTable        :: {-# UNPACK #-} !MemoryTable
   }
@@ -101,7 +61,7 @@
 instance Eq CUDA.Context where
   CUDA.Context p1 == CUDA.Context p2    = p1 == p2
 
-$(mkLabels [''CUDAState, ''KernelEntry])
+$(mkLabels [''CUDAState])
 
 
 -- Execution State
@@ -117,37 +77,66 @@
       CUDA.push ctx
       dev       <- CUDA.device
       prp       <- CUDA.props dev
-      return $! CUDAState prp knl mem
+      weak_ctx  <- mkWeakPtr ctx Nothing
+      return $! CUDAState prp (Context ctx weak_ctx) theKernelTable theMemoryTable
 
-    -- one-shot top-level mutable state
-    {-# NOINLINE mem #-}
-    {-# NOINLINE knl #-}
-    mem = unsafePerformIO MT.new
-    knl = unsafePerformIO HT.new
 
+-- Top-level mutable state
+-- -----------------------
+--
+-- It is important to keep some information alive for the entire run of the
+-- program, not just a single execution. These tokens use unsafePerformIO to
+-- ensure they are executed only once, and reused for subsequent invocations.
+--
 
+{-# NOINLINE theMemoryTable #-}
+theMemoryTable :: MemoryTable
+theMemoryTable = unsafePerformIO $ do
+  message dump_gc "gc: initialise memory table"
+  keepAlive =<< MT.new
+
+
+{-# NOINLINE theKernelTable #-}
+theKernelTable :: KernelTable
+theKernelTable = unsafePerformIO $ do
+  message dump_gc "gc: initialise kernel table"
+  keepAlive =<< KT.new
+
+
 -- Select and initialise a default CUDA device, and create a new execution
 -- context. The device is selected based on compute capability and estimated
 -- maximum throughput.
 --
 {-# NOINLINE defaultContext #-}
-defaultContext :: MVar CUDA.Context
+defaultContext :: CUDA.Context
 defaultContext = unsafePerformIO $ do
   CUDA.initialise []
   (dev,prp)     <- selectBestDevice
   ctx           <- CUDA.create dev [CUDA.SchedAuto] >> CUDA.pop
-  ref           <- newMVar ctx
   --
   message dump_gc $ "gc: initialise context"
   message verbose $ deviceInfo dev prp
   --
   addFinalizer ctx $ do
-    message dump_gc $ "gc: finalise context"
+    message dump_gc $ "gc: finalise context"    -- should never happen!
     CUDA.destroy ctx
   --
-  return ref
+  keepAlive ctx
 
 
+-- Make sure the GC knows that we want to keep this thing alive past the end of
+-- 'evalCUDA'.
+--
+-- We may want to introduce some way to actually shut this down if, for example,
+-- the object has not been accessed in a while, and so let it be collected.
+--
+keepAlive :: a -> IO a
+keepAlive x = forkIO (caffeine x) >> return x
+  where
+    caffeine hit = do threadDelay 5000000 -- microseconds = 5 seconds
+                      caffeine hit
+
+
 -- Debugging
 -- ---------
 
@@ -172,36 +161,4 @@
     clock       = showFFloatSIBase (Just 2) 1000 (fromIntegral $ CUDA.clockRate prp * 1000 :: Double) "Hz"
     mem         = showFFloatSIBase (Just 0) 1024 (fromIntegral $ CUDA.totalGlobalMem prp   :: Double) "B"
     at          = char '@'
-
-
--- Persistent caching (deprecated)
--- -------------------------------
-
-#ifdef ACCELERATE_CUDA_PERSISTENT_CACHE
--- Load and save the persistent kernel index file
---
-indexFileName :: IO FilePath
-indexFileName = do
-  tmp <- (</> "cache") `fmap` getDataDir
-  dir <- createDirectoryIfMissing True tmp >> canonicalizePath tmp
-  return (dir </> "_index")
-
-saveIndexFile :: CUDAState -> IO ()
-saveIndexFile s = do
-  ind <- indexFileName
-  encodeFile ind . map (second _kernelName) =<< HT.toList (_kernelTable s)
-
--- Read the kernel index map file (if it exists), loading modules into the
--- current context
---
-loadIndexFile :: IO (KernelTable, Int)
-loadIndexFile = do
-  f <- indexFileName
-  x <- doesFileExist f
-  e <- if x then mapM reload =<< decodeFile f
-            else return []
-  (,length e) <$> HT.fromList hashAccKey e
-  where
-    reload (k,n) = (k,) . KernelEntry n . Right <$> CUDA.loadFile (n `replaceExtension` ".cubin")
-#endif
 
diff --git a/accelerate-cuda.buildinfo.in b/accelerate-cuda.buildinfo.in
new file mode 100644
--- /dev/null
+++ b/accelerate-cuda.buildinfo.in
@@ -0,0 +1,3 @@
+ghc-options: @ghc_flags@
+cc-options: @cpp_flags@
+
diff --git a/accelerate-cuda.cabal b/accelerate-cuda.cabal
--- a/accelerate-cuda.cabal
+++ b/accelerate-cuda.cabal
@@ -1,5 +1,5 @@
 Name:                   accelerate-cuda
-Version:                0.12.0.0
+Version:                0.12.1.0
 Cabal-version:          >= 1.6
 Tested-with:            GHC >= 7.4
 Build-type:             Configure
@@ -9,9 +9,11 @@
   This library implements a backend for the Accelerate language instrumented for
   parallel execution on CUDA-capable NVIDIA GPUs.
   .
-  To use this backend you need CUDA version 3.x or later installed. Note that
-  currently there is no support for 'Char' and 'Bool' arrays (this is a
-  limitation of the front-end language).
+  To use this backend you need CUDA version 3.x or later installed, which you
+  can find at the NVIDIA Developer Zone.
+  .
+  <http://developer.nvidia.com/cuda-downloads>
+  .
 
 License:                BSD3
 License-file:           LICENSE
@@ -35,22 +37,38 @@
                         cubits/accelerate_cuda_shape.h
                         cubits/accelerate_cuda_stencil.h
                         cubits/accelerate_cuda_texture.h
+                        cubits/accelerate_cuda_type.h
                         cubits/accelerate_cuda_util.h
 
 Extra-tmp-files:        config.status
                         config.log
+                        accelerate-cuda.buildinfo               -- generated by configure
                         cubits/accelerate_cuda_shape.h          -- generated by configure
 
 Extra-source-files:     configure
+                        accelerate-cuda.buildinfo.in
                         cubits/accelerate_cuda_shape.h.in
                         include/accelerate.h
 
--- Flag pcache
---   Description:          Enable the persistent caching of the compiled CUDA modules (experimental)
---   Default:              False
-
 Flag debug
-  Description:          Enable tracing message flags
+  Description:
+    Enable tracing message flags. These are read from the command-line
+    arguments, which is convenient but may cause problems interacting with the
+    user program, so is not enabled by default. The available options:
+    .
+    * -ddump-cc: print the generated code, kernel table management information,
+         nvcc compiler warnings, and thread & resource statistics
+    .
+    * -ddebug-cc: compile code with debugging symbols, suitable for 'cuda-gdb'
+    .
+    * -ddump-exec: print each kernel name as it is invoked
+    .
+    * -ddump-gc: print memory management information
+    .
+    * -dverbose: other, uncategorised messages
+    .
+    * -fflush-cache: delete the persistent kernel cache
+    .
   Default:              False
 
 Flag bounds-checks
@@ -68,7 +86,7 @@
 Library
   Include-Dirs:         include
 
-  Build-depends:        accelerate              == 0.12.*,
+  Build-depends:        accelerate              >= 0.12.1 && < 0.13,
                         array                   >= 0.3,
                         base                    == 4.*,
                         binary                  >= 0.5,
@@ -88,7 +106,6 @@
                         pretty                  >= 1.0,
                         process                 >= 1.0,
                         srcloc                  >= 0.1,
-                        symbol                  >= 0.1,
                         transformers            >= 0.2,
                         unordered-containers    >= 0.1.4
 
@@ -119,12 +136,10 @@
                         Data.Array.Accelerate.CUDA.Debug
                         Data.Array.Accelerate.CUDA.Execute
                         Data.Array.Accelerate.CUDA.FullList
+                        Data.Array.Accelerate.CUDA.Persistent
                         Data.Array.Accelerate.CUDA.State
                         Data.Array.Accelerate.Internal.Check
                         Paths_accelerate_cuda
-
---  if flag(pcache)
---    CPP-options:        -DACCELERATE_CUDA_PERSISTENT_CACHE
 
   if flag(debug)
     cpp-options:        -DACCELERATE_DEBUG
diff --git a/configure b/configure
--- a/configure
+++ b/configure
@@ -1,13 +1,13 @@
 #! /bin/sh
 # Guess values for system-dependent variables and create Makefiles.
-# Generated by GNU Autoconf 2.65 for accelerate-cuda 0.9.0.0.
+# Generated by GNU Autoconf 2.68 for accelerate-cuda 0.13.0.0.
 #
-# Report bugs to <accelerate@projects.haskell.org>.
+# Report bugs to <accelerate-haskell@googlegroups.com>.
 #
 #
 # Copyright (C) 1992, 1993, 1994, 1995, 1996, 1998, 1999, 2000, 2001,
-# 2002, 2003, 2004, 2005, 2006, 2007, 2008, 2009 Free Software Foundation,
-# Inc.
+# 2002, 2003, 2004, 2005, 2006, 2007, 2008, 2009, 2010 Free Software
+# Foundation, Inc.
 #
 #
 # This configure script is free software; the Free Software Foundation
@@ -91,6 +91,7 @@
 IFS=" ""	$as_nl"
 
 # Find who we are.  Look in the path if we contain no directory separator.
+as_myself=
 case $0 in #((
   *[\\/]* ) as_myself=$0 ;;
   *) as_save_IFS=$IFS; IFS=$PATH_SEPARATOR
@@ -215,11 +216,18 @@
   # We cannot yet assume a decent shell, so we have to provide a
 	# neutralization value for shells without unset; and this also
 	# works around shells that cannot unset nonexistent variables.
+	# Preserve -v and -x to the replacement shell.
 	BASH_ENV=/dev/null
 	ENV=/dev/null
 	(unset BASH_ENV) >/dev/null 2>&1 && unset BASH_ENV ENV
 	export CONFIG_SHELL
-	exec "$CONFIG_SHELL" "$as_myself" ${1+"$@"}
+	case $- in # ((((
+	  *v*x* | *x*v* ) as_opts=-vx ;;
+	  *v* ) as_opts=-v ;;
+	  *x* ) as_opts=-x ;;
+	  * ) as_opts= ;;
+	esac
+	exec "$CONFIG_SHELL" $as_opts "$as_myself" ${1+"$@"}
 fi
 
     if test x$as_have_required = xno; then :
@@ -230,7 +238,7 @@
     $as_echo "$0: be upgraded to zsh 4.3.4 or later."
   else
     $as_echo "$0: Please tell bug-autoconf@gnu.org and
-$0: accelerate@projects.haskell.org about your system,
+$0: accelerate-haskell@googlegroups.com about your system,
 $0: including any error possibly output before this
 $0: message. Then install a modern shell, or manually run
 $0: the script under such a shell if you do have one."
@@ -318,7 +326,7 @@
       test -d "$as_dir" && break
     done
     test -z "$as_dirs" || eval "mkdir $as_dirs"
-  } || test -d "$as_dir" || as_fn_error "cannot create directory $as_dir"
+  } || test -d "$as_dir" || as_fn_error $? "cannot create directory $as_dir"
 
 
 } # as_fn_mkdir_p
@@ -358,19 +366,19 @@
 fi # as_fn_arith
 
 
-# as_fn_error ERROR [LINENO LOG_FD]
-# ---------------------------------
+# as_fn_error STATUS ERROR [LINENO LOG_FD]
+# ----------------------------------------
 # Output "`basename $0`: error: ERROR" to stderr. If LINENO and LOG_FD are
 # provided, also output the error to LOG_FD, referencing LINENO. Then exit the
-# script with status $?, using 1 if that was 0.
+# script with STATUS, using 1 if that was 0.
 as_fn_error ()
 {
-  as_status=$?; test $as_status -eq 0 && as_status=1
-  if test "$3"; then
-    as_lineno=${as_lineno-"$2"} as_lineno_stack=as_lineno_stack=$as_lineno_stack
-    $as_echo "$as_me:${as_lineno-$LINENO}: error: $1" >&$3
+  as_status=$1; test $as_status -eq 0 && as_status=1
+  if test "$4"; then
+    as_lineno=${as_lineno-"$3"} as_lineno_stack=as_lineno_stack=$as_lineno_stack
+    $as_echo "$as_me:${as_lineno-$LINENO}: error: $2" >&$4
   fi
-  $as_echo "$as_me: error: $1" >&2
+  $as_echo "$as_me: error: $2" >&2
   as_fn_exit $as_status
 } # as_fn_error
 
@@ -532,7 +540,7 @@
 exec 6>&1
 
 # Name of the host.
-# hostname on some systems (SVR3.2, Linux) returns a bogus exit status,
+# hostname on some systems (SVR3.2, old GNU/Linux) returns a bogus exit status,
 # so uname gets run too.
 ac_hostname=`(hostname || uname -n) 2>/dev/null | sed 1q`
 
@@ -551,15 +559,18 @@
 # Identity of this package.
 PACKAGE_NAME='accelerate-cuda'
 PACKAGE_TARNAME='accelerate-cuda'
-PACKAGE_VERSION='0.9.0.0'
-PACKAGE_STRING='accelerate-cuda 0.9.0.0'
-PACKAGE_BUGREPORT='accelerate@projects.haskell.org'
+PACKAGE_VERSION='0.13.0.0'
+PACKAGE_STRING='accelerate-cuda 0.13.0.0'
+PACKAGE_BUGREPORT='accelerate-haskell@googlegroups.com'
 PACKAGE_URL=''
 
 ac_unique_file="Data/Array/Accelerate/CUDA.hs"
 ac_subst_vars='LTLIBOBJS
 LIBOBJS
-hs_int_type
+cpp_flags
+ghc_flags
+type_hs_char
+type_hs_int
 GHC
 OBJEXT
 EXEEXT
@@ -681,8 +692,9 @@
   fi
 
   case $ac_option in
-  *=*)	ac_optarg=`expr "X$ac_option" : '[^=]*=\(.*\)'` ;;
-  *)	ac_optarg=yes ;;
+  *=?*) ac_optarg=`expr "X$ac_option" : '[^=]*=\(.*\)'` ;;
+  *=)   ac_optarg= ;;
+  *)    ac_optarg=yes ;;
   esac
 
   # Accept the important Cygnus configure options, so we can diagnose typos.
@@ -727,7 +739,7 @@
     ac_useropt=`expr "x$ac_option" : 'x-*disable-\(.*\)'`
     # Reject names that are not valid shell variable names.
     expr "x$ac_useropt" : ".*[^-+._$as_cr_alnum]" >/dev/null &&
-      as_fn_error "invalid feature name: $ac_useropt"
+      as_fn_error $? "invalid feature name: $ac_useropt"
     ac_useropt_orig=$ac_useropt
     ac_useropt=`$as_echo "$ac_useropt" | sed 's/[-+.]/_/g'`
     case $ac_user_opts in
@@ -753,7 +765,7 @@
     ac_useropt=`expr "x$ac_option" : 'x-*enable-\([^=]*\)'`
     # Reject names that are not valid shell variable names.
     expr "x$ac_useropt" : ".*[^-+._$as_cr_alnum]" >/dev/null &&
-      as_fn_error "invalid feature name: $ac_useropt"
+      as_fn_error $? "invalid feature name: $ac_useropt"
     ac_useropt_orig=$ac_useropt
     ac_useropt=`$as_echo "$ac_useropt" | sed 's/[-+.]/_/g'`
     case $ac_user_opts in
@@ -957,7 +969,7 @@
     ac_useropt=`expr "x$ac_option" : 'x-*with-\([^=]*\)'`
     # Reject names that are not valid shell variable names.
     expr "x$ac_useropt" : ".*[^-+._$as_cr_alnum]" >/dev/null &&
-      as_fn_error "invalid package name: $ac_useropt"
+      as_fn_error $? "invalid package name: $ac_useropt"
     ac_useropt_orig=$ac_useropt
     ac_useropt=`$as_echo "$ac_useropt" | sed 's/[-+.]/_/g'`
     case $ac_user_opts in
@@ -973,7 +985,7 @@
     ac_useropt=`expr "x$ac_option" : 'x-*without-\(.*\)'`
     # Reject names that are not valid shell variable names.
     expr "x$ac_useropt" : ".*[^-+._$as_cr_alnum]" >/dev/null &&
-      as_fn_error "invalid package name: $ac_useropt"
+      as_fn_error $? "invalid package name: $ac_useropt"
     ac_useropt_orig=$ac_useropt
     ac_useropt=`$as_echo "$ac_useropt" | sed 's/[-+.]/_/g'`
     case $ac_user_opts in
@@ -1003,8 +1015,8 @@
   | --x-librar=* | --x-libra=* | --x-libr=* | --x-lib=* | --x-li=* | --x-l=*)
     x_libraries=$ac_optarg ;;
 
-  -*) as_fn_error "unrecognized option: \`$ac_option'
-Try \`$0 --help' for more information."
+  -*) as_fn_error $? "unrecognized option: \`$ac_option'
+Try \`$0 --help' for more information"
     ;;
 
   *=*)
@@ -1012,7 +1024,7 @@
     # Reject names that are not valid shell variable names.
     case $ac_envvar in #(
       '' | [0-9]* | *[!_$as_cr_alnum]* )
-      as_fn_error "invalid variable name: \`$ac_envvar'" ;;
+      as_fn_error $? "invalid variable name: \`$ac_envvar'" ;;
     esac
     eval $ac_envvar=\$ac_optarg
     export $ac_envvar ;;
@@ -1022,7 +1034,7 @@
     $as_echo "$as_me: WARNING: you should use --build, --host, --target" >&2
     expr "x$ac_option" : ".*[^-._$as_cr_alnum]" >/dev/null &&
       $as_echo "$as_me: WARNING: invalid host type: $ac_option" >&2
-    : ${build_alias=$ac_option} ${host_alias=$ac_option} ${target_alias=$ac_option}
+    : "${build_alias=$ac_option} ${host_alias=$ac_option} ${target_alias=$ac_option}"
     ;;
 
   esac
@@ -1030,13 +1042,13 @@
 
 if test -n "$ac_prev"; then
   ac_option=--`echo $ac_prev | sed 's/_/-/g'`
-  as_fn_error "missing argument to $ac_option"
+  as_fn_error $? "missing argument to $ac_option"
 fi
 
 if test -n "$ac_unrecognized_opts"; then
   case $enable_option_checking in
     no) ;;
-    fatal) as_fn_error "unrecognized options: $ac_unrecognized_opts" ;;
+    fatal) as_fn_error $? "unrecognized options: $ac_unrecognized_opts" ;;
     *)     $as_echo "$as_me: WARNING: unrecognized options: $ac_unrecognized_opts" >&2 ;;
   esac
 fi
@@ -1059,7 +1071,7 @@
     [\\/$]* | ?:[\\/]* )  continue;;
     NONE | '' ) case $ac_var in *prefix ) continue;; esac;;
   esac
-  as_fn_error "expected an absolute directory name for --$ac_var: $ac_val"
+  as_fn_error $? "expected an absolute directory name for --$ac_var: $ac_val"
 done
 
 # There might be people who depend on the old broken behavior: `$host'
@@ -1073,8 +1085,8 @@
 if test "x$host_alias" != x; then
   if test "x$build_alias" = x; then
     cross_compiling=maybe
-    $as_echo "$as_me: WARNING: If you wanted to set the --build type, don't use --host.
-    If a cross compiler is detected then cross compile mode will be used." >&2
+    $as_echo "$as_me: WARNING: if you wanted to set the --build type, don't use --host.
+    If a cross compiler is detected then cross compile mode will be used" >&2
   elif test "x$build_alias" != "x$host_alias"; then
     cross_compiling=yes
   fi
@@ -1089,9 +1101,9 @@
 ac_pwd=`pwd` && test -n "$ac_pwd" &&
 ac_ls_di=`ls -di .` &&
 ac_pwd_ls_di=`cd "$ac_pwd" && ls -di .` ||
-  as_fn_error "working directory cannot be determined"
+  as_fn_error $? "working directory cannot be determined"
 test "X$ac_ls_di" = "X$ac_pwd_ls_di" ||
-  as_fn_error "pwd does not report name of working directory"
+  as_fn_error $? "pwd does not report name of working directory"
 
 
 # Find the source files, if location was not specified.
@@ -1130,11 +1142,11 @@
 fi
 if test ! -r "$srcdir/$ac_unique_file"; then
   test "$ac_srcdir_defaulted" = yes && srcdir="$ac_confdir or .."
-  as_fn_error "cannot find sources ($ac_unique_file) in $srcdir"
+  as_fn_error $? "cannot find sources ($ac_unique_file) in $srcdir"
 fi
 ac_msg="sources are in $srcdir, but \`cd $srcdir' does not work"
 ac_abs_confdir=`(
-	cd "$srcdir" && test -r "./$ac_unique_file" || as_fn_error "$ac_msg"
+	cd "$srcdir" && test -r "./$ac_unique_file" || as_fn_error $? "$ac_msg"
 	pwd)`
 # When building in place, set srcdir=.
 if test "$ac_abs_confdir" = "$ac_pwd"; then
@@ -1160,7 +1172,7 @@
   # Omit some internal or obsolete options to make the list less imposing.
   # This message is too long to be a string in the A/UX 3.1 sh.
   cat <<_ACEOF
-\`configure' configures accelerate-cuda 0.9.0.0 to adapt to many kinds of systems.
+\`configure' configures accelerate-cuda 0.13.0.0 to adapt to many kinds of systems.
 
 Usage: $0 [OPTION]... [VAR=VALUE]...
 
@@ -1174,7 +1186,7 @@
       --help=short        display options specific to this package
       --help=recursive    display the short help of all the included packages
   -V, --version           display version information and exit
-  -q, --quiet, --silent   do not print \`checking...' messages
+  -q, --quiet, --silent   do not print \`checking ...' messages
       --cache-file=FILE   cache test results in FILE [disabled]
   -C, --config-cache      alias for \`--cache-file=config.cache'
   -n, --no-create         do not create output files
@@ -1221,7 +1233,7 @@
 
 if test -n "$ac_init_help"; then
   case $ac_init_help in
-     short | recursive ) echo "Configuration of accelerate-cuda 0.9.0.0:";;
+     short | recursive ) echo "Configuration of accelerate-cuda 0.13.0.0:";;
    esac
   cat <<\_ACEOF
 
@@ -1237,7 +1249,7 @@
 Use these variables to override the choices made by `configure' or to help
 it to find libraries and programs with nonstandard names/locations.
 
-Report bugs to <accelerate@projects.haskell.org>.
+Report bugs to <accelerate-haskell@googlegroups.com>.
 _ACEOF
 ac_status=$?
 fi
@@ -1300,10 +1312,10 @@
 test -n "$ac_init_help" && exit $ac_status
 if $ac_init_version; then
   cat <<\_ACEOF
-accelerate-cuda configure 0.9.0.0
-generated by GNU Autoconf 2.65
+accelerate-cuda configure 0.13.0.0
+generated by GNU Autoconf 2.68
 
-Copyright (C) 2009 Free Software Foundation, Inc.
+Copyright (C) 2010 Free Software Foundation, Inc.
 This configure script is free software; the Free Software Foundation
 gives unlimited permission to copy, distribute and modify it.
 _ACEOF
@@ -1347,7 +1359,7 @@
 
 	ac_retval=1
 fi
-  eval $as_lineno_stack; test "x$as_lineno_stack" = x && { as_lineno=; unset as_lineno;}
+  eval $as_lineno_stack; ${as_lineno_stack:+:} unset as_lineno
   as_fn_set_status $ac_retval
 
 } # ac_fn_cxx_try_compile
@@ -1355,8 +1367,8 @@
 This file contains any messages produced by compilers while
 running configure, to aid debugging if configure makes a mistake.
 
-It was created by accelerate-cuda $as_me 0.9.0.0, which was
-generated by GNU Autoconf 2.65.  Invocation command line was
+It was created by accelerate-cuda $as_me 0.13.0.0, which was
+generated by GNU Autoconf 2.68.  Invocation command line was
 
   $ $0 $@
 
@@ -1466,11 +1478,9 @@
   {
     echo
 
-    cat <<\_ASBOX
-## ---------------- ##
+    $as_echo "## ---------------- ##
 ## Cache variables. ##
-## ---------------- ##
-_ASBOX
+## ---------------- ##"
     echo
     # The following way of writing the cache mishandles newlines in values,
 (
@@ -1504,11 +1514,9 @@
 )
     echo
 
-    cat <<\_ASBOX
-## ----------------- ##
+    $as_echo "## ----------------- ##
 ## Output variables. ##
-## ----------------- ##
-_ASBOX
+## ----------------- ##"
     echo
     for ac_var in $ac_subst_vars
     do
@@ -1521,11 +1529,9 @@
     echo
 
     if test -n "$ac_subst_files"; then
-      cat <<\_ASBOX
-## ------------------- ##
+      $as_echo "## ------------------- ##
 ## File substitutions. ##
-## ------------------- ##
-_ASBOX
+## ------------------- ##"
       echo
       for ac_var in $ac_subst_files
       do
@@ -1539,11 +1545,9 @@
     fi
 
     if test -s confdefs.h; then
-      cat <<\_ASBOX
-## ----------- ##
+      $as_echo "## ----------- ##
 ## confdefs.h. ##
-## ----------- ##
-_ASBOX
+## ----------- ##"
       echo
       cat confdefs.h
       echo
@@ -1598,7 +1602,12 @@
 ac_site_file1=NONE
 ac_site_file2=NONE
 if test -n "$CONFIG_SITE"; then
-  ac_site_file1=$CONFIG_SITE
+  # We do not want a PATH search for config.site.
+  case $CONFIG_SITE in #((
+    -*)  ac_site_file1=./$CONFIG_SITE;;
+    */*) ac_site_file1=$CONFIG_SITE;;
+    *)   ac_site_file1=./$CONFIG_SITE;;
+  esac
 elif test "x$prefix" != xNONE; then
   ac_site_file1=$prefix/share/config.site
   ac_site_file2=$prefix/etc/config.site
@@ -1613,7 +1622,11 @@
     { $as_echo "$as_me:${as_lineno-$LINENO}: loading site script $ac_site_file" >&5
 $as_echo "$as_me: loading site script $ac_site_file" >&6;}
     sed 's/^/| /' "$ac_site_file" >&5
-    . "$ac_site_file"
+    . "$ac_site_file" \
+      || { { $as_echo "$as_me:${as_lineno-$LINENO}: error: in \`$ac_pwd':" >&5
+$as_echo "$as_me: error: in \`$ac_pwd':" >&2;}
+as_fn_error $? "failed to load site script $ac_site_file
+See \`config.log' for more details" "$LINENO" 5; }
   fi
 done
 
@@ -1689,7 +1702,7 @@
 $as_echo "$as_me: error: in \`$ac_pwd':" >&2;}
   { $as_echo "$as_me:${as_lineno-$LINENO}: error: changes in the environment can compromise the build" >&5
 $as_echo "$as_me: error: changes in the environment can compromise the build" >&2;}
-  as_fn_error "run \`make distclean' and/or \`rm $cache_file' and start over" "$LINENO" 5
+  as_fn_error $? "run \`make distclean' and/or \`rm $cache_file' and start over" "$LINENO" 5
 fi
 ## -------------------- ##
 ## Main body of script. ##
@@ -1703,7 +1716,7 @@
 
 
 
-ac_config_files="$ac_config_files cubits/accelerate_cuda_shape.h"
+ac_config_files="$ac_config_files accelerate-cuda.buildinfo cubits/accelerate_cuda_shape.h"
 
 
 # This doesn't pick up a user specified '--with-compiler' flag
@@ -1724,7 +1737,7 @@
 set dummy $ac_tool_prefix$ac_prog; ac_word=$2
 { $as_echo "$as_me:${as_lineno-$LINENO}: checking for $ac_word" >&5
 $as_echo_n "checking for $ac_word... " >&6; }
-if test "${ac_cv_prog_CXX+set}" = set; then :
+if ${ac_cv_prog_CXX+:} false; then :
   $as_echo_n "(cached) " >&6
 else
   if test -n "$CXX"; then
@@ -1768,7 +1781,7 @@
 set dummy $ac_prog; ac_word=$2
 { $as_echo "$as_me:${as_lineno-$LINENO}: checking for $ac_word" >&5
 $as_echo_n "checking for $ac_word... " >&6; }
-if test "${ac_cv_prog_ac_ct_CXX+set}" = set; then :
+if ${ac_cv_prog_ac_ct_CXX+:} false; then :
   $as_echo_n "(cached) " >&6
 else
   if test -n "$ac_ct_CXX"; then
@@ -1931,9 +1944,8 @@
 
 { { $as_echo "$as_me:${as_lineno-$LINENO}: error: in \`$ac_pwd':" >&5
 $as_echo "$as_me: error: in \`$ac_pwd':" >&2;}
-{ as_fn_set_status 77
-as_fn_error "C++ compiler cannot create executables
-See \`config.log' for more details." "$LINENO" 5; }; }
+as_fn_error 77 "C++ compiler cannot create executables
+See \`config.log' for more details" "$LINENO" 5; }
 else
   { $as_echo "$as_me:${as_lineno-$LINENO}: result: yes" >&5
 $as_echo "yes" >&6; }
@@ -1975,8 +1987,8 @@
 else
   { { $as_echo "$as_me:${as_lineno-$LINENO}: error: in \`$ac_pwd':" >&5
 $as_echo "$as_me: error: in \`$ac_pwd':" >&2;}
-as_fn_error "cannot compute suffix of executables: cannot compile and link
-See \`config.log' for more details." "$LINENO" 5; }
+as_fn_error $? "cannot compute suffix of executables: cannot compile and link
+See \`config.log' for more details" "$LINENO" 5; }
 fi
 rm -f conftest conftest$ac_cv_exeext
 { $as_echo "$as_me:${as_lineno-$LINENO}: result: $ac_cv_exeext" >&5
@@ -2033,9 +2045,9 @@
     else
 	{ { $as_echo "$as_me:${as_lineno-$LINENO}: error: in \`$ac_pwd':" >&5
 $as_echo "$as_me: error: in \`$ac_pwd':" >&2;}
-as_fn_error "cannot run C++ compiled programs.
+as_fn_error $? "cannot run C++ compiled programs.
 If you meant to cross compile, use \`--host'.
-See \`config.log' for more details." "$LINENO" 5; }
+See \`config.log' for more details" "$LINENO" 5; }
     fi
   fi
 fi
@@ -2046,7 +2058,7 @@
 ac_clean_files=$ac_clean_files_save
 { $as_echo "$as_me:${as_lineno-$LINENO}: checking for suffix of object files" >&5
 $as_echo_n "checking for suffix of object files... " >&6; }
-if test "${ac_cv_objext+set}" = set; then :
+if ${ac_cv_objext+:} false; then :
   $as_echo_n "(cached) " >&6
 else
   cat confdefs.h - <<_ACEOF >conftest.$ac_ext
@@ -2086,8 +2098,8 @@
 
 { { $as_echo "$as_me:${as_lineno-$LINENO}: error: in \`$ac_pwd':" >&5
 $as_echo "$as_me: error: in \`$ac_pwd':" >&2;}
-as_fn_error "cannot compute suffix of object files: cannot compile
-See \`config.log' for more details." "$LINENO" 5; }
+as_fn_error $? "cannot compute suffix of object files: cannot compile
+See \`config.log' for more details" "$LINENO" 5; }
 fi
 rm -f conftest.$ac_cv_objext conftest.$ac_ext
 fi
@@ -2097,7 +2109,7 @@
 ac_objext=$OBJEXT
 { $as_echo "$as_me:${as_lineno-$LINENO}: checking whether we are using the GNU C++ compiler" >&5
 $as_echo_n "checking whether we are using the GNU C++ compiler... " >&6; }
-if test "${ac_cv_cxx_compiler_gnu+set}" = set; then :
+if ${ac_cv_cxx_compiler_gnu+:} false; then :
   $as_echo_n "(cached) " >&6
 else
   cat confdefs.h - <<_ACEOF >conftest.$ac_ext
@@ -2134,7 +2146,7 @@
 ac_save_CXXFLAGS=$CXXFLAGS
 { $as_echo "$as_me:${as_lineno-$LINENO}: checking whether $CXX accepts -g" >&5
 $as_echo_n "checking whether $CXX accepts -g... " >&6; }
-if test "${ac_cv_prog_cxx_g+set}" = set; then :
+if ${ac_cv_prog_cxx_g+:} false; then :
   $as_echo_n "(cached) " >&6
 else
   ac_save_cxx_werror_flag=$ac_cxx_werror_flag
@@ -2220,7 +2232,7 @@
 set dummy ghc; ac_word=$2
 { $as_echo "$as_me:${as_lineno-$LINENO}: checking for $ac_word" >&5
 $as_echo_n "checking for $ac_word... " >&6; }
-if test "${ac_cv_path_GHC+set}" = set; then :
+if ${ac_cv_path_GHC+:} false; then :
   $as_echo_n "(cached) " >&6
 else
   case $GHC in
@@ -2256,20 +2268,48 @@
 fi
 
 
-{ $as_echo "$as_me:${as_lineno-$LINENO}: checking size of Int" >&5
+
+
+
+
+    { $as_echo "$as_me:${as_lineno-$LINENO}: checking size of Int" >&5
 $as_echo_n "checking size of Int... " >&6; }
 
-hs_int_size=`$GHC -w -ignore-dot-ghci -e "Foreign.sizeOf (undefined::Int)"`
-{ $as_echo "$as_me:${as_lineno-$LINENO}: result: $hs_int_size" >&5
-$as_echo "$hs_int_size" >&6; }
+    sizeof_hs_Int=`$GHC -w -ignore-dot-ghci -e "Foreign.sizeOf (undefined::Int)"`
+    { $as_echo "$as_me:${as_lineno-$LINENO}: result: $sizeof_hs_Int" >&5
+$as_echo "$sizeof_hs_Int" >&6; }
 
-case $hs_int_size in
-    4) hs_int_type=int32_t ;;
-    8) hs_int_type=int64_t ;;
-    *) as_fn_error "could not determine size of a Haskell Int" "$LINENO" 5
-esac
+    case $sizeof_hs_Int in
+        4) type_hs_int=Int32 ;;
+        8) type_hs_int=Int64 ;;
+    esac
 
+    def="-DSIZEOF_HS$(echo Int | tr [:lower:] [:upper:])=$sizeof_hs_Int"
+    cpp_flags="$cpp_flags $def"
+    ghc_flags="$ghc_flags -optP$def"
 
+
+    { $as_echo "$as_me:${as_lineno-$LINENO}: checking size of Char" >&5
+$as_echo_n "checking size of Char... " >&6; }
+
+    sizeof_hs_Char=`$GHC -w -ignore-dot-ghci -e "Foreign.sizeOf (undefined::Char)"`
+    { $as_echo "$as_me:${as_lineno-$LINENO}: result: $sizeof_hs_Char" >&5
+$as_echo "$sizeof_hs_Char" >&6; }
+
+    case $sizeof_hs_Char in
+        4) type_hs_char=Word32 ;;
+        8) type_hs_char=Word64 ;;
+    esac
+
+    def="-DSIZEOF_HS$(echo Char | tr [:lower:] [:upper:])=$sizeof_hs_Char"
+    cpp_flags="$cpp_flags $def"
+    ghc_flags="$ghc_flags -optP$def"
+
+
+
+
+
+
 cat >confcache <<\_ACEOF
 # This file is a shell script that caches the results of configure
 # tests run on this system so they can be shared between configure
@@ -2334,10 +2374,21 @@
      :end' >>confcache
 if diff "$cache_file" confcache >/dev/null 2>&1; then :; else
   if test -w "$cache_file"; then
-    test "x$cache_file" != "x/dev/null" &&
+    if test "x$cache_file" != "x/dev/null"; then
       { $as_echo "$as_me:${as_lineno-$LINENO}: updating cache $cache_file" >&5
 $as_echo "$as_me: updating cache $cache_file" >&6;}
-    cat confcache >$cache_file
+      if test ! -f "$cache_file" || test -h "$cache_file"; then
+	cat confcache >"$cache_file"
+      else
+        case $cache_file in #(
+        */* | ?:*)
+	  mv -f confcache "$cache_file"$$ &&
+	  mv -f "$cache_file"$$ "$cache_file" ;; #(
+        *)
+	  mv -f confcache "$cache_file" ;;
+	esac
+      fi
+    fi
   else
     { $as_echo "$as_me:${as_lineno-$LINENO}: not updating unwritable cache $cache_file" >&5
 $as_echo "$as_me: not updating unwritable cache $cache_file" >&6;}
@@ -2389,6 +2440,7 @@
 
 ac_libobjs=
 ac_ltlibobjs=
+U=
 for ac_i in : $LIBOBJS; do test "x$ac_i" = x: && continue
   # 1. Remove the extension, and $U if already installed.
   ac_script='s/\$U\././;s/\.o$//;s/\.obj$//'
@@ -2404,7 +2456,7 @@
 
 
 
-: ${CONFIG_STATUS=./config.status}
+: "${CONFIG_STATUS=./config.status}"
 ac_write_fail=0
 ac_clean_files_save=$ac_clean_files
 ac_clean_files="$ac_clean_files $CONFIG_STATUS"
@@ -2505,6 +2557,7 @@
 IFS=" ""	$as_nl"
 
 # Find who we are.  Look in the path if we contain no directory separator.
+as_myself=
 case $0 in #((
   *[\\/]* ) as_myself=$0 ;;
   *) as_save_IFS=$IFS; IFS=$PATH_SEPARATOR
@@ -2550,19 +2603,19 @@
 (unset CDPATH) >/dev/null 2>&1 && unset CDPATH
 
 
-# as_fn_error ERROR [LINENO LOG_FD]
-# ---------------------------------
+# as_fn_error STATUS ERROR [LINENO LOG_FD]
+# ----------------------------------------
 # Output "`basename $0`: error: ERROR" to stderr. If LINENO and LOG_FD are
 # provided, also output the error to LOG_FD, referencing LINENO. Then exit the
-# script with status $?, using 1 if that was 0.
+# script with STATUS, using 1 if that was 0.
 as_fn_error ()
 {
-  as_status=$?; test $as_status -eq 0 && as_status=1
-  if test "$3"; then
-    as_lineno=${as_lineno-"$2"} as_lineno_stack=as_lineno_stack=$as_lineno_stack
-    $as_echo "$as_me:${as_lineno-$LINENO}: error: $1" >&$3
+  as_status=$1; test $as_status -eq 0 && as_status=1
+  if test "$4"; then
+    as_lineno=${as_lineno-"$3"} as_lineno_stack=as_lineno_stack=$as_lineno_stack
+    $as_echo "$as_me:${as_lineno-$LINENO}: error: $2" >&$4
   fi
-  $as_echo "$as_me: error: $1" >&2
+  $as_echo "$as_me: error: $2" >&2
   as_fn_exit $as_status
 } # as_fn_error
 
@@ -2758,7 +2811,7 @@
       test -d "$as_dir" && break
     done
     test -z "$as_dirs" || eval "mkdir $as_dirs"
-  } || test -d "$as_dir" || as_fn_error "cannot create directory $as_dir"
+  } || test -d "$as_dir" || as_fn_error $? "cannot create directory $as_dir"
 
 
 } # as_fn_mkdir_p
@@ -2811,8 +2864,8 @@
 # report actual input values of CONFIG_FILES etc. instead of their
 # values after options handling.
 ac_log="
-This file was extended by accelerate-cuda $as_me 0.9.0.0, which was
-generated by GNU Autoconf 2.65.  Invocation command line was
+This file was extended by accelerate-cuda $as_me 0.13.0.0, which was
+generated by GNU Autoconf 2.68.  Invocation command line was
 
   CONFIG_FILES    = $CONFIG_FILES
   CONFIG_HEADERS  = $CONFIG_HEADERS
@@ -2858,17 +2911,17 @@
 Configuration files:
 $config_files
 
-Report bugs to <accelerate@projects.haskell.org>."
+Report bugs to <accelerate-haskell@googlegroups.com>."
 
 _ACEOF
 cat >>$CONFIG_STATUS <<_ACEOF || ac_write_fail=1
 ac_cs_config="`$as_echo "$ac_configure_args" | sed 's/^ //; s/[\\""\`\$]/\\\\&/g'`"
 ac_cs_version="\\
-accelerate-cuda config.status 0.9.0.0
-configured by $0, generated by GNU Autoconf 2.65,
+accelerate-cuda config.status 0.13.0.0
+configured by $0, generated by GNU Autoconf 2.68,
   with options \\"\$ac_cs_config\\"
 
-Copyright (C) 2009 Free Software Foundation, Inc.
+Copyright (C) 2010 Free Software Foundation, Inc.
 This config.status script is free software; the Free Software Foundation
 gives unlimited permission to copy, distribute and modify it."
 
@@ -2883,11 +2936,16 @@
 while test $# != 0
 do
   case $1 in
-  --*=*)
+  --*=?*)
     ac_option=`expr "X$1" : 'X\([^=]*\)='`
     ac_optarg=`expr "X$1" : 'X[^=]*=\(.*\)'`
     ac_shift=:
     ;;
+  --*=)
+    ac_option=`expr "X$1" : 'X\([^=]*\)='`
+    ac_optarg=
+    ac_shift=:
+    ;;
   *)
     ac_option=$1
     ac_optarg=$2
@@ -2909,6 +2967,7 @@
     $ac_shift
     case $ac_optarg in
     *\'*) ac_optarg=`$as_echo "$ac_optarg" | sed "s/'/'\\\\\\\\''/g"` ;;
+    '') as_fn_error $? "missing file argument" ;;
     esac
     as_fn_append CONFIG_FILES " '$ac_optarg'"
     ac_need_defaults=false;;
@@ -2919,7 +2978,7 @@
     ac_cs_silent=: ;;
 
   # This is an error.
-  -*) as_fn_error "unrecognized option: \`$1'
+  -*) as_fn_error $? "unrecognized option: \`$1'
 Try \`$0 --help' for more information." ;;
 
   *) as_fn_append ac_config_targets " $1"
@@ -2968,9 +3027,10 @@
 for ac_config_target in $ac_config_targets
 do
   case $ac_config_target in
+    "accelerate-cuda.buildinfo") CONFIG_FILES="$CONFIG_FILES accelerate-cuda.buildinfo" ;;
     "cubits/accelerate_cuda_shape.h") CONFIG_FILES="$CONFIG_FILES cubits/accelerate_cuda_shape.h" ;;
 
-  *) as_fn_error "invalid argument: \`$ac_config_target'" "$LINENO" 5;;
+  *) as_fn_error $? "invalid argument: \`$ac_config_target'" "$LINENO" 5;;
   esac
 done
 
@@ -2991,9 +3051,10 @@
 # after its creation but before its name has been assigned to `$tmp'.
 $debug ||
 {
-  tmp=
+  tmp= ac_tmp=
   trap 'exit_status=$?
-  { test -z "$tmp" || test ! -d "$tmp" || rm -fr "$tmp"; } && exit $exit_status
+  : "${ac_tmp:=$tmp}"
+  { test ! -d "$ac_tmp" || rm -fr "$ac_tmp"; } && exit $exit_status
 ' 0
   trap 'as_fn_exit 1' 1 2 13 15
 }
@@ -3001,12 +3062,13 @@
 
 {
   tmp=`(umask 077 && mktemp -d "./confXXXXXX") 2>/dev/null` &&
-  test -n "$tmp" && test -d "$tmp"
+  test -d "$tmp"
 }  ||
 {
   tmp=./conf$$-$RANDOM
   (umask 077 && mkdir "$tmp")
-} || as_fn_error "cannot create a temporary directory in ." "$LINENO" 5
+} || as_fn_error $? "cannot create a temporary directory in ." "$LINENO" 5
+ac_tmp=$tmp
 
 # Set up the scripts for CONFIG_FILES section.
 # No need to generate them if there are no CONFIG_FILES.
@@ -3023,12 +3085,12 @@
 fi
 ac_cs_awk_cr=`$AWK 'BEGIN { print "a\rb" }' </dev/null 2>/dev/null`
 if test "$ac_cs_awk_cr" = "a${ac_cr}b"; then
-  ac_cs_awk_cr='\r'
+  ac_cs_awk_cr='\\r'
 else
   ac_cs_awk_cr=$ac_cr
 fi
 
-echo 'BEGIN {' >"$tmp/subs1.awk" &&
+echo 'BEGIN {' >"$ac_tmp/subs1.awk" &&
 _ACEOF
 
 
@@ -3037,18 +3099,18 @@
   echo "$ac_subst_vars" | sed 's/.*/&!$&$ac_delim/' &&
   echo "_ACEOF"
 } >conf$$subs.sh ||
-  as_fn_error "could not make $CONFIG_STATUS" "$LINENO" 5
-ac_delim_num=`echo "$ac_subst_vars" | grep -c '$'`
+  as_fn_error $? "could not make $CONFIG_STATUS" "$LINENO" 5
+ac_delim_num=`echo "$ac_subst_vars" | grep -c '^'`
 ac_delim='%!_!# '
 for ac_last_try in false false false false false :; do
   . ./conf$$subs.sh ||
-    as_fn_error "could not make $CONFIG_STATUS" "$LINENO" 5
+    as_fn_error $? "could not make $CONFIG_STATUS" "$LINENO" 5
 
   ac_delim_n=`sed -n "s/.*$ac_delim\$/X/p" conf$$subs.awk | grep -c X`
   if test $ac_delim_n = $ac_delim_num; then
     break
   elif $ac_last_try; then
-    as_fn_error "could not make $CONFIG_STATUS" "$LINENO" 5
+    as_fn_error $? "could not make $CONFIG_STATUS" "$LINENO" 5
   else
     ac_delim="$ac_delim!$ac_delim _$ac_delim!! "
   fi
@@ -3056,7 +3118,7 @@
 rm -f conf$$subs.sh
 
 cat >>$CONFIG_STATUS <<_ACEOF || ac_write_fail=1
-cat >>"\$tmp/subs1.awk" <<\\_ACAWK &&
+cat >>"\$ac_tmp/subs1.awk" <<\\_ACAWK &&
 _ACEOF
 sed -n '
 h
@@ -3104,7 +3166,7 @@
 rm -f conf$$subs.awk
 cat >>$CONFIG_STATUS <<_ACEOF || ac_write_fail=1
 _ACAWK
-cat >>"\$tmp/subs1.awk" <<_ACAWK &&
+cat >>"\$ac_tmp/subs1.awk" <<_ACAWK &&
   for (key in S) S_is_set[key] = 1
   FS = ""
 
@@ -3136,21 +3198,29 @@
   sed "s/$ac_cr\$//; s/$ac_cr/$ac_cs_awk_cr/g"
 else
   cat
-fi < "$tmp/subs1.awk" > "$tmp/subs.awk" \
-  || as_fn_error "could not setup config files machinery" "$LINENO" 5
+fi < "$ac_tmp/subs1.awk" > "$ac_tmp/subs.awk" \
+  || as_fn_error $? "could not setup config files machinery" "$LINENO" 5
 _ACEOF
 
-# VPATH may cause trouble with some makes, so we remove $(srcdir),
-# ${srcdir} and @srcdir@ from VPATH if srcdir is ".", strip leading and
+# VPATH may cause trouble with some makes, so we remove sole $(srcdir),
+# ${srcdir} and @srcdir@ entries from VPATH if srcdir is ".", strip leading and
 # trailing colons and then remove the whole line if VPATH becomes empty
 # (actually we leave an empty line to preserve line numbers).
 if test "x$srcdir" = x.; then
-  ac_vpsub='/^[	 ]*VPATH[	 ]*=/{
-s/:*\$(srcdir):*/:/
-s/:*\${srcdir}:*/:/
-s/:*@srcdir@:*/:/
-s/^\([^=]*=[	 ]*\):*/\1/
+  ac_vpsub='/^[	 ]*VPATH[	 ]*=[	 ]*/{
+h
+s///
+s/^/:/
+s/[	 ]*$/:/
+s/:\$(srcdir):/:/g
+s/:\${srcdir}:/:/g
+s/:@srcdir@:/:/g
+s/^:*//
 s/:*$//
+x
+s/\(=[	 ]*\).*/\1/
+G
+s/\n//
 s/^[^=]*=[	 ]*$//
 }'
 fi
@@ -3168,7 +3238,7 @@
   esac
   case $ac_mode$ac_tag in
   :[FHL]*:*);;
-  :L* | :C*:*) as_fn_error "invalid tag \`$ac_tag'" "$LINENO" 5;;
+  :L* | :C*:*) as_fn_error $? "invalid tag \`$ac_tag'" "$LINENO" 5;;
   :[FH]-) ac_tag=-:-;;
   :[FH]*) ac_tag=$ac_tag:$ac_tag.in;;
   esac
@@ -3187,7 +3257,7 @@
     for ac_f
     do
       case $ac_f in
-      -) ac_f="$tmp/stdin";;
+      -) ac_f="$ac_tmp/stdin";;
       *) # Look for the file first in the build tree, then in the source tree
 	 # (if the path is not absolute).  The absolute path cannot be DOS-style,
 	 # because $ac_f cannot contain `:'.
@@ -3196,7 +3266,7 @@
 	   [\\/$]*) false;;
 	   *) test -f "$srcdir/$ac_f" && ac_f="$srcdir/$ac_f";;
 	   esac ||
-	   as_fn_error "cannot find input file: \`$ac_f'" "$LINENO" 5;;
+	   as_fn_error 1 "cannot find input file: \`$ac_f'" "$LINENO" 5;;
       esac
       case $ac_f in *\'*) ac_f=`$as_echo "$ac_f" | sed "s/'/'\\\\\\\\''/g"`;; esac
       as_fn_append ac_file_inputs " '$ac_f'"
@@ -3222,8 +3292,8 @@
     esac
 
     case $ac_tag in
-    *:-:* | *:-) cat >"$tmp/stdin" \
-      || as_fn_error "could not create $ac_file" "$LINENO" 5 ;;
+    *:-:* | *:-) cat >"$ac_tmp/stdin" \
+      || as_fn_error $? "could not create $ac_file" "$LINENO" 5 ;;
     esac
     ;;
   esac
@@ -3348,23 +3418,24 @@
 s&@abs_top_builddir@&$ac_abs_top_builddir&;t t
 $ac_datarootdir_hack
 "
-eval sed \"\$ac_sed_extra\" "$ac_file_inputs" | $AWK -f "$tmp/subs.awk" >$tmp/out \
-  || as_fn_error "could not create $ac_file" "$LINENO" 5
+eval sed \"\$ac_sed_extra\" "$ac_file_inputs" | $AWK -f "$ac_tmp/subs.awk" \
+  >$ac_tmp/out || as_fn_error $? "could not create $ac_file" "$LINENO" 5
 
 test -z "$ac_datarootdir_hack$ac_datarootdir_seen" &&
-  { ac_out=`sed -n '/\${datarootdir}/p' "$tmp/out"`; test -n "$ac_out"; } &&
-  { ac_out=`sed -n '/^[	 ]*datarootdir[	 ]*:*=/p' "$tmp/out"`; test -z "$ac_out"; } &&
+  { ac_out=`sed -n '/\${datarootdir}/p' "$ac_tmp/out"`; test -n "$ac_out"; } &&
+  { ac_out=`sed -n '/^[	 ]*datarootdir[	 ]*:*=/p' \
+      "$ac_tmp/out"`; test -z "$ac_out"; } &&
   { $as_echo "$as_me:${as_lineno-$LINENO}: WARNING: $ac_file contains a reference to the variable \`datarootdir'
-which seems to be undefined.  Please make sure it is defined." >&5
+which seems to be undefined.  Please make sure it is defined" >&5
 $as_echo "$as_me: WARNING: $ac_file contains a reference to the variable \`datarootdir'
-which seems to be undefined.  Please make sure it is defined." >&2;}
+which seems to be undefined.  Please make sure it is defined" >&2;}
 
-  rm -f "$tmp/stdin"
+  rm -f "$ac_tmp/stdin"
   case $ac_file in
-  -) cat "$tmp/out" && rm -f "$tmp/out";;
-  *) rm -f "$ac_file" && mv "$tmp/out" "$ac_file";;
+  -) cat "$ac_tmp/out" && rm -f "$ac_tmp/out";;
+  *) rm -f "$ac_file" && mv "$ac_tmp/out" "$ac_file";;
   esac \
-  || as_fn_error "could not create $ac_file" "$LINENO" 5
+  || as_fn_error $? "could not create $ac_file" "$LINENO" 5
  ;;
 
 
@@ -3379,7 +3450,7 @@
 ac_clean_files=$ac_clean_files_save
 
 test $ac_write_fail = 0 ||
-  as_fn_error "write failure creating $CONFIG_STATUS" "$LINENO" 5
+  as_fn_error $? "write failure creating $CONFIG_STATUS" "$LINENO" 5
 
 
 # configure is writing to config.log, and then calls config.status.
@@ -3400,7 +3471,7 @@
   exec 5>>config.log
   # Use ||, not &&, to avoid exiting from the if with $? = 1, which
   # would make configure fail if this is the last instruction.
-  $ac_cs_success || as_fn_exit $?
+  $ac_cs_success || as_fn_exit 1
 fi
 if test -n "$ac_unrecognized_opts" && test "$enable_option_checking" != no; then
   { $as_echo "$as_me:${as_lineno-$LINENO}: WARNING: unrecognized options: $ac_unrecognized_opts" >&5
diff --git a/cubits/accelerate_cuda_extras.h b/cubits/accelerate_cuda_extras.h
--- a/cubits/accelerate_cuda_extras.h
+++ b/cubits/accelerate_cuda_extras.h
@@ -17,6 +17,7 @@
 #include "accelerate_cuda_shape.h"
 #include "accelerate_cuda_stencil.h"
 #include "accelerate_cuda_texture.h"
+#include "accelerate_cuda_type.h"
 
 #endif
 
diff --git a/cubits/accelerate_cuda_function.h b/cubits/accelerate_cuda_function.h
--- a/cubits/accelerate_cuda_function.h
+++ b/cubits/accelerate_cuda_function.h
@@ -13,8 +13,8 @@
 #ifndef __ACCELERATE_CUDA_FUNCTION_H__
 #define __ACCELERATE_CUDA_FUNCTION_H__
 
-#include <stdint.h>
 #include <cuda_runtime.h>
+#include "accelerate_cuda_type.h"
 
 #ifdef __cplusplus
 
@@ -26,16 +26,16 @@
  * Left/Right bitwise rotation
  */
 template <typename T>
-static __inline__ __device__ T rotateL(const T x, const int32_t i)
+static __inline__ __device__ T rotateL(const T x, const Int32 i)
 {
-    const int32_t i8 = i & 8 * sizeof(x) - 1;
+    const Int32 i8 = i & 8 * sizeof(x) - 1;
     return i8 == 0 ? x : x << i8 | x >> 8 * sizeof(x) - i8;
 }
 
 template <typename T>
-static __inline__ __device__ T rotateR(const T x, const int32_t i)
+static __inline__ __device__ T rotateR(const T x, const Int32 i)
 {
-    const int32_t i8 = i & 8 * sizeof(x) - 1;
+    const Int32 i8 = i & 8 * sizeof(x) - 1;
     return i8 == 0 ? x : x >> i8 | x << 8 * sizeof(x) - i8;
 }
 
@@ -63,24 +63,24 @@
  * Type coercion
  */
 template <typename T>
-static __inline__ __device__ uint32_t reinterpret32(const T x)
+static __inline__ __device__ Word32 reinterpret32(const T x)
 {
-    union { T a; uint32_t b; } u;
+    union { T a; Word32 b; } u;
 
     u.a = x;
     return u.b;
 }
 
 template <>
-static __inline__ __device__ uint32_t reinterpret32(const float x)
+static __inline__ __device__ Word32 reinterpret32(const float x)
 {
     return __float_as_int(x);
 }
 
 template <typename T>
-static __inline__ __device__ uint64_t reinterpret64(const T x)
+static __inline__ __device__ Word64 reinterpret64(const T x)
 {
-    union { T a; uint64_t b; } u;
+    union { T a; Word64 b; } u;
 
     u.a = x;
     return u.b;
@@ -88,7 +88,7 @@
 
 #if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 130
 template <>
-static __inline__ __device__ uint64_t reinterpret64(const double x)
+static __inline__ __device__ Word64 reinterpret64(const double x)
 {
     return __double_as_longlong(x);
 }
@@ -101,36 +101,38 @@
 template <typename T>
 static __inline__ __device__ T atomicCAS32(T* address, T compare, T val)
 {
-    union { T a; uint32_t b; } u;
+    union { T a; Word32 b; } u;
 
-    u.b = atomicCAS((uint32_t*) address, reinterpret32<T>(compare), reinterpret32<T>(val));
+    u.b = atomicCAS((Word32*) address, reinterpret32<T>(compare), reinterpret32<T>(val));
     return u.a;
 }
 
+#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 110
 template <>
-static __inline__ __device__ int32_t atomicCAS32(int32_t* address, int32_t compare, int32_t val)
+static __inline__ __device__ Int32 atomicCAS32(Int32* address, Int32 compare, Int32 val)
 {
     return atomicCAS(address, compare, val);
 }
 
 template <>
-static __inline__ __device__ uint32_t atomicCAS32(uint32_t* address, uint32_t compare, uint32_t val)
+static __inline__ __device__ Word32 atomicCAS32(Word32* address, Word32 compare, Word32 val)
 {
     return atomicCAS(address, compare, val);
 }
+#endif
 
 #if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 120
 template <typename T>
 static __inline__ __device__ T atomicCAS64(T* address, T compare, T val)
 {
-    union { T a; unsigned long long int b; } u;
+    union { T a; Word64 b; } u;
 
-    u.b = atomicCAS((unsigned long long int*) address, reinterpret64<T>(compare), reinterpret64<T>(val));
+    u.b = atomicCAS((Word64*) address, reinterpret64<T>(compare), reinterpret64<T>(val));
     return u.a;
 }
 
 template <>
-static __inline__ __device__ unsigned long long int atomicCAS64(unsigned long long int* address, unsigned long long int compare, unsigned long long int val)
+static __inline__ __device__ Word64 atomicCAS64(Word64* address, Word64 compare, Word64 val)
 {
     return atomicCAS(address, compare, val);
 }
diff --git a/cubits/accelerate_cuda_shape.h b/cubits/accelerate_cuda_shape.h
--- a/cubits/accelerate_cuda_shape.h
+++ b/cubits/accelerate_cuda_shape.h
@@ -13,8 +13,8 @@
 #ifndef __ACCELERATE_CUDA_SHAPE_H__
 #define __ACCELERATE_CUDA_SHAPE_H__
 
-#include <stdint.h>
 #include <cuda_runtime.h>
+#include "accelerate_cuda_type.h"
 
 /*
  * For the time being, the D.A.A.CUDA.Execute converts shape components from
@@ -23,9 +23,9 @@
  * future hardware gains better 64-bit support and/or we need to access very
  * large arrays.
  *
- * typedef int64_t                             Ix;
+ * typedef Int64                             Ix;
  */
-typedef int32_t                                   Ix;
+typedef Int32                                     Ix;
 typedef void*                                     DIM0;
 typedef Ix                                        DIM1;
 typedef struct { Ix a1,a0; }                      DIM2;
diff --git a/cubits/accelerate_cuda_shape.h.in b/cubits/accelerate_cuda_shape.h.in
--- a/cubits/accelerate_cuda_shape.h.in
+++ b/cubits/accelerate_cuda_shape.h.in
@@ -13,8 +13,8 @@
 #ifndef __ACCELERATE_CUDA_SHAPE_H__
 #define __ACCELERATE_CUDA_SHAPE_H__
 
-#include <stdint.h>
 #include <cuda_runtime.h>
+#include "accelerate_cuda_type.h"
 
 /*
  * For the time being, the D.A.A.CUDA.Execute converts shape components from
@@ -23,9 +23,9 @@
  * future hardware gains better 64-bit support and/or we need to access very
  * large arrays.
  *
- * typedef @hs_int_type@                             Ix;
+ * typedef @type_hs_int@                             Ix;
  */
-typedef int32_t                                   Ix;
+typedef Int32                                     Ix;
 typedef void*                                     DIM0;
 typedef Ix                                        DIM1;
 typedef struct { Ix a1,a0; }                      DIM2;
diff --git a/cubits/accelerate_cuda_stencil.h b/cubits/accelerate_cuda_stencil.h
--- a/cubits/accelerate_cuda_stencil.h
+++ b/cubits/accelerate_cuda_stencil.h
@@ -53,7 +53,7 @@
 static __inline__ __device__ DIM1 clamp(const DIM1 sz, const DIM1 i)
 {
     // CUDA-4.0 does not include 64-bit min/max functions for Fermi (?)
-    return max(0, min((int) i, (int) sz-1));
+    return max(0, min(i, sz-1));
 }
 
 
@@ -88,6 +88,6 @@
     else              return i;
 }
 
-#endif
+#endif  // __cplusplus
+#endif  // __ACCELERATE_CUDA_STENCIL_H__
 
-#endif
diff --git a/cubits/accelerate_cuda_texture.h b/cubits/accelerate_cuda_texture.h
--- a/cubits/accelerate_cuda_texture.h
+++ b/cubits/accelerate_cuda_texture.h
@@ -9,8 +9,7 @@
  * Stability   : experimental
  *
  * CUDA texture definitions and access functions are defined in terms of
- * templates, and hence only available through the C++ interface. Expose some
- * dummy wrappers to enable parsing with language-c.
+ * templates, and hence only available through the C++ interface.
  *
  * We don't have a definition for `Int' or `Word', since the bitwidth of the
  * Haskell and C types may be different.
@@ -25,38 +24,38 @@
 #ifndef __ACCELERATE_CUDA_TEXTURE_H__
 #define __ACCELERATE_CUDA_TEXTURE_H__
 
-#include <stdint.h>
 #include <cuda_runtime.h>
+#include "accelerate_cuda_type.h"
 
 #if defined(__cplusplus) && defined(__CUDACC__)
 
-typedef texture<uint2,    1> TexWord64;
-typedef texture<uint32_t, 1> TexWord32;
-typedef texture<uint16_t, 1> TexWord16;
-typedef texture<uint8_t,  1> TexWord8;
-typedef texture<int2,     1> TexInt64;
-typedef texture<int32_t,  1> TexInt32;
-typedef texture<int16_t,  1> TexInt16;
-typedef texture<int8_t,   1> TexInt8;
-typedef texture<float,    1> TexFloat;
-typedef texture<char,     1> TexCChar;
+typedef texture<uint2,  1> TexWord64;
+typedef texture<Word32, 1> TexWord32;
+typedef texture<Word16, 1> TexWord16;
+typedef texture<Word8,  1> TexWord8;
+typedef texture<int2,   1> TexInt64;
+typedef texture<Int32,  1> TexInt32;
+typedef texture<Int16,  1> TexInt16;
+typedef texture<Int8,   1> TexInt8;
+typedef texture<float,  1> TexFloat;
+typedef texture<char,   1> TexCChar;
 
-static __inline__ __device__ uint8_t  indexArray(TexWord8  t, const int x) { return tex1Dfetch(t,x); }
-static __inline__ __device__ uint16_t indexArray(TexWord16 t, const int x) { return tex1Dfetch(t,x); }
-static __inline__ __device__ uint32_t indexArray(TexWord32 t, const int x) { return tex1Dfetch(t,x); }
-static __inline__ __device__ uint64_t indexArray(TexWord64 t, const int x)
+static __inline__ __device__  Word8 indexArray(TexWord8  t, const int x) { return tex1Dfetch(t,x); }
+static __inline__ __device__ Word16 indexArray(TexWord16 t, const int x) { return tex1Dfetch(t,x); }
+static __inline__ __device__ Word32 indexArray(TexWord32 t, const int x) { return tex1Dfetch(t,x); }
+static __inline__ __device__ Word64 indexArray(TexWord64 t, const int x)
 {
-  union { uint2 x; uint64_t y; } v;
+  union { uint2 x; Word64 y; } v;
   v.x = tex1Dfetch(t,x);
   return v.y;
 }
 
-static __inline__ __device__ int8_t  indexArray(TexInt8  t, const int x) { return tex1Dfetch(t,x); }
-static __inline__ __device__ int16_t indexArray(TexInt16 t, const int x) { return tex1Dfetch(t,x); }
-static __inline__ __device__ int32_t indexArray(TexInt32 t, const int x) { return tex1Dfetch(t,x); }
-static __inline__ __device__ int64_t indexArray(TexInt64 t, const int x)
+static __inline__ __device__  Int8 indexArray(TexInt8  t, const int x) { return tex1Dfetch(t,x); }
+static __inline__ __device__ Int16 indexArray(TexInt16 t, const int x) { return tex1Dfetch(t,x); }
+static __inline__ __device__ Int32 indexArray(TexInt32 t, const int x) { return tex1Dfetch(t,x); }
+static __inline__ __device__ Int64 indexArray(TexInt64 t, const int x)
 {
-  union { int2 x; int64_t y; } v;
+  union { int2 x; Int64 y; } v;
   v.x = tex1Dfetch(t,x);
   return v.y;
 }
@@ -98,35 +97,6 @@
   return __hiloint2double(v.y,v.x);
 }
 #endif
-
-#else
-
-typedef void* TexWord64;
-typedef void* TexWord32;
-typedef void* TexWord16;
-typedef void* TexWord8;
-typedef void* TexInt64;
-typedef void* TexInt32;
-typedef void* TexInt16;
-typedef void* TexInt8;
-typedef void* TexCShort;
-typedef void* TexCUShort;
-typedef void* TexCInt;
-typedef void* TexCUInt;
-typedef void* TexCLong;
-typedef void* TexCULong;
-typedef void* TexCLLong;
-typedef void* TexCULLong;
-typedef void* TexFloat;
-typedef void* TexDouble;
-typedef void* TexCFloat;
-typedef void* TexCDouble;
-typedef void* TexCChar;
-typedef void* TexCSChar;
-typedef void* TexCUChar;
-
-void* indexArray(const void*, const int);
-void* indexDArray(const void*, const int);
 
 #endif  // defined(__cplusplus) && defined(__CUDACC__)
 #endif  // __ACCELERATE_CUDA_TEXTURE_H__
diff --git a/cubits/accelerate_cuda_type.h b/cubits/accelerate_cuda_type.h
new file mode 100644
--- /dev/null
+++ b/cubits/accelerate_cuda_type.h
@@ -0,0 +1,32 @@
+/* -----------------------------------------------------------------------------
+ *
+ * Module      : Type
+ * Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
+ *               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+ * License     : BSD3
+ *
+ * Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
+ * Stability   : experimental
+ *
+ * ---------------------------------------------------------------------------*/
+
+#ifndef __ACCELERATE_CUDA_TYPE_H__
+#define __ACCELERATE_CUDA_TYPE_H__
+
+/*
+ * The word size on CUDA devices is always 32-bits. If the host is a 64-bit
+ * architecture, the system stdint.h implementation can define these incorrectly
+ * for device code.
+ *
+ */
+typedef signed char            Int8;
+typedef short                 Int16;
+typedef int                   Int32;
+typedef long long             Int64;
+
+typedef unsigned char         Word8;
+typedef unsigned short       Word16;
+typedef unsigned int         Word32;
+typedef unsigned long long   Word64;
+
+#endif