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
@@ -4,8 +4,9 @@
 {-# LANGUAGE ScopedTypeVariables #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2013] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2008..2009] Sean Lee
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
@@ -120,9 +121,7 @@
 --    '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.
+--    non-unique indices, the combination function requires compute-1.1.
 --
 --
 -- [/Debugging options:/]
diff --git a/Data/Array/Accelerate/CUDA/AST.hs b/Data/Array/Accelerate/CUDA/AST.hs
--- a/Data/Array/Accelerate/CUDA/AST.hs
+++ b/Data/Array/Accelerate/CUDA/AST.hs
@@ -4,8 +4,8 @@
 {-# LANGUAGE TypeSynonymInstances       #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.AST
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
diff --git a/Data/Array/Accelerate/CUDA/Analysis/Device.hs b/Data/Array/Accelerate/CUDA/Analysis/Device.hs
--- a/Data/Array/Accelerate/CUDA/Analysis/Device.hs
+++ b/Data/Array/Accelerate/CUDA/Analysis/Device.hs
@@ -1,7 +1,7 @@
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.Analysis.Device
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
diff --git a/Data/Array/Accelerate/CUDA/Analysis/Launch.hs b/Data/Array/Accelerate/CUDA/Analysis/Launch.hs
--- a/Data/Array/Accelerate/CUDA/Analysis/Launch.hs
+++ b/Data/Array/Accelerate/CUDA/Analysis/Launch.hs
@@ -1,9 +1,9 @@
-{-# LANGUAGE CPP   #-}
-{-# LANGUAGE GADTs #-}
+{-# LANGUAGE GADTs           #-}
+{-# LANGUAGE TemplateHaskell #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.Analysis.Launch
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
@@ -19,6 +19,7 @@
 
 -- friends
 import Data.Array.Accelerate.AST
+import Data.Array.Accelerate.Error
 import Data.Array.Accelerate.Trafo
 import Data.Array.Accelerate.Analysis.Type
 import Data.Array.Accelerate.Analysis.Shape
@@ -27,9 +28,7 @@
 import qualified Foreign.CUDA.Analysis                  as CUDA
 import qualified Foreign.CUDA.Driver                    as CUDA
 
-#include "accelerate.h"
 
-
 -- |
 -- Determine kernel launch parameters for the given array computation (as well
 -- as compiled function module). This consists of the thread block size, number
@@ -47,7 +46,7 @@
     -> ( Int                    -- block size
        , Int -> Int             -- number of blocks for input problem size (grid)
        , Int )                  -- shared memory (bytes)
-launchConfig Delayed{} _ _ = INTERNAL_ERROR(error) "launchConfig" "encountered delayed array"
+launchConfig Delayed{} _ _ = $internalError "launchConfig" "encountered delayed array"
 launchConfig (Manifest acc) dev occ =
   let cta       = CUDA.activeThreads occ `div` CUDA.activeThreadBlocks occ
       maxGrid   = CUDA.multiProcessorCount dev * CUDA.activeThreadBlocks occ
@@ -66,7 +65,7 @@
     -> CUDA.Fun                 -- corresponding __global__ entry function
     -> Int                      -- maximum number of threads per block
     -> IO CUDA.Occupancy
-determineOccupancy Delayed{} _ _ _ = INTERNAL_ERROR(error) "determineOccupancy" "encountered delayed array"
+determineOccupancy Delayed{} _ _ _ = $internalError "determineOccupancy" "encountered delayed array"
 determineOccupancy (Manifest acc) dev fn maxBlock = do
   registers     <- CUDA.requires fn CUDA.NumRegs
   static_smem   <- CUDA.requires fn CUDA.SharedSizeBytes        -- static memory only
@@ -137,17 +136,17 @@
 --
 sharedMem :: CUDA.DeviceProperties -> PreOpenAcc DelayedOpenAcc aenv a -> Int -> Int
 -- non-computation forms
-sharedMem _ Alet{}     _ = INTERNAL_ERROR(error) "sharedMem" "Let"
-sharedMem _ Avar{}     _ = INTERNAL_ERROR(error) "sharedMem" "Avar"
-sharedMem _ Apply{}    _ = INTERNAL_ERROR(error) "sharedMem" "Apply"
-sharedMem _ Acond{}    _ = INTERNAL_ERROR(error) "sharedMem" "Acond"
-sharedMem _ Awhile{}   _ = INTERNAL_ERROR(error) "sharedMem" "Awhile"
-sharedMem _ Atuple{}   _ = INTERNAL_ERROR(error) "sharedMem" "Atuple"
-sharedMem _ Aprj{}     _ = INTERNAL_ERROR(error) "sharedMem" "Aprj"
-sharedMem _ Use{}      _ = INTERNAL_ERROR(error) "sharedMem" "Use"
-sharedMem _ Unit{}     _ = INTERNAL_ERROR(error) "sharedMem" "Unit"
-sharedMem _ Reshape{}  _ = INTERNAL_ERROR(error) "sharedMem" "Reshape"
-sharedMem _ Aforeign{} _ = INTERNAL_ERROR(error) "sharedMem" "Aforeign"
+sharedMem _ Alet{}     _ = $internalError "sharedMem" "Let"
+sharedMem _ Avar{}     _ = $internalError "sharedMem" "Avar"
+sharedMem _ Apply{}    _ = $internalError "sharedMem" "Apply"
+sharedMem _ Acond{}    _ = $internalError "sharedMem" "Acond"
+sharedMem _ Awhile{}   _ = $internalError "sharedMem" "Awhile"
+sharedMem _ Atuple{}   _ = $internalError "sharedMem" "Atuple"
+sharedMem _ Aprj{}     _ = $internalError "sharedMem" "Aprj"
+sharedMem _ Use{}      _ = $internalError "sharedMem" "Use"
+sharedMem _ Unit{}     _ = $internalError "sharedMem" "Unit"
+sharedMem _ Reshape{}  _ = $internalError "sharedMem" "Reshape"
+sharedMem _ Aforeign{} _ = $internalError "sharedMem" "Aforeign"
 
 -- skeleton nodes
 sharedMem _ Generate{}          _        = 0
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
@@ -2,12 +2,13 @@
 {-# LANGUAGE CPP                 #-}
 {-# LANGUAGE GADTs               #-}
 {-# LANGUAGE ScopedTypeVariables #-}
+{-# LANGUAGE TemplateHaskell     #-}
 {-# LANGUAGE TypeFamilies        #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.Array.Data
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2013] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
---               [2013]       Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell, Robert Clifton-Everest
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
+--               [2013..2014] Robert Clifton-Everest
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
@@ -21,7 +22,7 @@
   mallocArray, indexArray,
   useArray,  useArrayAsync,
   useDevicePtrs,
-  copyArray, copyArrayPeer, copyArrayPeerAsync,
+  copyArray, copyArrayAsync, copyArrayPeer, copyArrayPeerAsync,
   peekArray, peekArrayAsync,
   pokeArray, pokeArrayAsync,
   marshalArrayData, marshalTextureData, marshalDevicePtrs,
@@ -44,6 +45,7 @@
 import Foreign.Ptr
 
 -- friends
+import Data.Array.Accelerate.Error
 import Data.Array.Accelerate.Array.Data
 import Data.Array.Accelerate.Array.Sugar                ( Array(..), Shape, Elt, toElt, EltRepr )
 import Data.Array.Accelerate.Array.Representation       ( size )
@@ -54,9 +56,7 @@
 import qualified Foreign.CUDA.Driver.Stream             as CUDA
 import qualified Foreign.CUDA.Driver.Texture            as CUDA
 
-#include "accelerate.h"
 
-
 -- Array Operations
 -- ----------------
 
@@ -246,7 +246,7 @@
 --
 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)
+  = $boundsCheck "copyArray" "shape mismatch" (sh1 == sh2)
   $ run doCopy
   where
     !n              = size sh1
@@ -261,13 +261,30 @@
         copyPrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> ArrayData e -> Int -> IO ()
         mkPrimDispatch(copyPrim,Prim.copyArray)
 
+copyArrayAsync :: (Shape dim, Elt e) => Array dim e -> Array dim e -> Maybe CUDA.Stream -> CIO ()
+copyArrayAsync (Array !sh1 !adata1) (Array !sh2 !adata2) ms
+  = $boundsCheck "copyArrayAsync" "shape mismatch" (sh1 == sh2)
+  $ run doCopy
+  where
+    !n              = size sh1
+    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 ctx mt ad1 ad2 n ms
+        --
+        copyPrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> ArrayData e -> Int -> Maybe CUDA.Stream -> IO ()
+        mkPrimDispatch(copyPrim,Prim.copyArrayAsync)
 
+
 -- |Copy data between two device arrays which reside in different contexts. This
 -- might entail copying between devices.
 --
 copyArrayPeer :: (Shape dim, Elt e) => Array dim e -> Context -> Array dim e -> Context -> CIO ()
 copyArrayPeer (Array !sh1 !adata1) !ctxSrc (Array !sh2 !adata2) !ctxDst
-  = BOUNDS_CHECK(check) "copyArrayPeer" "shape mismatch" (sh1 == sh2)
+  = $boundsCheck "copyArrayPeer" "shape mismatch" (sh1 == sh2)
   $ run doCopy
   where
     !n           = size sh1
@@ -284,7 +301,7 @@
 
 copyArrayPeerAsync :: (Shape dim, Elt e) => Array dim e -> Context -> Array dim e -> Context -> Maybe CUDA.Stream -> CIO ()
 copyArrayPeerAsync (Array !sh1 !adata1) !ctxSrc (Array !sh2 !adata2) !ctxDst !ms
-  = BOUNDS_CHECK(check) "copyArrayPeerAsync" "shape mismatch" (sh1 == sh2)
+  = $boundsCheck "copyArrayPeerAsync" "shape mismatch" (sh1 == sh2)
   $ run doCopy
   where
     !n           = size sh1
@@ -453,3 +470,4 @@
     --
     advancePrim :: ArrayEltR e -> ArrayData e -> Prim.DevicePtrs e -> Prim.DevicePtrs e
     mkPrimDispatch(advancePrim,Prim.advancePtrsOfArrayData n)
+
diff --git a/Data/Array/Accelerate/CUDA/Array/Nursery.hs b/Data/Array/Accelerate/CUDA/Array/Nursery.hs
--- a/Data/Array/Accelerate/CUDA/Array/Nursery.hs
+++ b/Data/Array/Accelerate/CUDA/Array/Nursery.hs
@@ -3,8 +3,8 @@
 {-# OPTIONS_GHC -fno-warn-orphans #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.Array.Nursery
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2013] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
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
@@ -2,12 +2,13 @@
 {-# LANGUAGE CPP                 #-}
 {-# LANGUAGE GADTs               #-}
 {-# LANGUAGE ScopedTypeVariables #-}
-{-# LANGUAGE TypeFamilies        #-}
+{-# LANGUAGE TemplateHaskell     #-}
 {-# LANGUAGE TupleSections       #-}
+{-# LANGUAGE TypeFamilies        #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.Array.Prim
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2013] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
@@ -20,9 +21,8 @@
   DevicePtrs, HostPtrs,
 
   mallocArray, indexArray,
-  useArray,  useArrayAsync,
-  useDevicePtrs,
-  copyArray, copyArrayPeer, copyArrayPeerAsync,
+  useArray,  useArrayAsync, useDevicePtrs,
+  copyArray, copyArrayAsync, copyArrayPeer, copyArrayPeerAsync,
   peekArray, peekArrayAsync,
   pokeArray, pokeArrayAsync,
   marshalDevicePtrs, marshalArrayData, marshalTextureData,
@@ -38,6 +38,7 @@
 import Data.Functor
 import Data.Typeable
 import Control.Monad
+import Language.Haskell.TH
 import System.Mem.StableName
 import Foreign.Ptr
 import Foreign.C.Types
@@ -48,14 +49,13 @@
 import qualified Foreign.CUDA.Driver.Texture            as CUDA
 
 -- friends
+import Data.Array.Accelerate.Error
 import Data.Array.Accelerate.Array.Data
 import Data.Array.Accelerate.CUDA.Context
 import Data.Array.Accelerate.CUDA.Array.Table
 import qualified Data.Array.Accelerate.CUDA.Debug       as D
 
-#include "accelerate.h"
 
-
 -- Device array representation
 -- ---------------------------
 
@@ -98,7 +98,7 @@
 primArrayEltAs(CDouble, Double)
 
 primArrayElt(Char)
-primArrayEltAs(CChar,  Int8)
+primArrayEltAs(CChar,  HTYPE_CCHAR)
 primArrayEltAs(CSChar, Int8)
 primArrayEltAs(CUChar, Word8)
 
@@ -140,12 +140,13 @@
 instance TextureData CSChar  where format _ = (CUDA.Int8,   1)
 instance TextureData CUChar  where format _ = (CUDA.Word8,  1)
 instance TextureData Char    where format _ = (CUDA.Word32, 1)
-instance TextureData Int     where format _ = (CUDA.Int32,  SIZEOF_HTYPE_INT           `div` 4)
-instance TextureData Word    where format _ = (CUDA.Word32, SIZEOF_HTYPE_WORD          `div` 4)
-instance TextureData CLong   where format _ = (CUDA.Int32,  SIZEOF_HTYPE_LONG          `div` 4)
-instance TextureData CULong  where format _ = (CUDA.Word32, SIZEOF_HTYPE_UNSIGNED_LONG `div` 4)
 
+$( runQ [d| instance TextureData Int    where format _ = (CUDA.Int32,  sizeOf (undefined::Int)    `div` 4) |] )
+$( runQ [d| instance TextureData Word   where format _ = (CUDA.Word32, sizeOf (undefined::Word)   `div` 4) |] )
+$( runQ [d| instance TextureData CLong  where format _ = (CUDA.Int32,  sizeOf (undefined::CLong)  `div` 4) |] )
+$( runQ [d| instance TextureData CULong where format _ = (CUDA.Word32, sizeOf (undefined::CULong) `div` 4) |] )
 
+
 -- Primitive array operations
 -- --------------------------
 
@@ -258,9 +259,23 @@
 copyArray !ctx !mt !from !to !n = do
   src <- devicePtrsOfArrayData ctx mt from
   dst <- devicePtrsOfArrayData ctx mt to
-  transfer "copyArrayAsync" (n * sizeOf (undefined :: b)) $
-    CUDA.copyArrayAsync n src dst
+  transfer "copyArray" (n * sizeOf (undefined :: b)) $
+    CUDA.copyArray n src dst
 
+copyArrayAsync
+    :: 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
+    -> Maybe CUDA.Stream
+    -> IO ()
+copyArrayAsync !ctx !mt !from !to !n !mst = do
+  src <- devicePtrsOfArrayData ctx mt from
+  dst <- devicePtrsOfArrayData ctx mt to
+  transfer "copyArrayAsync" (n * sizeOf (undefined :: b)) $
+    CUDA.copyArrayAsync n src dst mst
 
 -- Copy data between two device arrays that exist in different contexts and/or
 -- devices.
@@ -402,7 +417,7 @@
     Just v  -> return v
     Nothing -> do
       sn <- makeStableName ad
-      INTERNAL_ERROR(error) "devicePtrsOfArrayData" $ "lost device memory #" ++ show (hashStableName sn)
+      $internalError "devicePtrsOfArrayData" $ "lost device memory #" ++ show (hashStableName sn)
 
 
 -- Advance device pointers by a given number of elements
@@ -439,5 +454,5 @@
                                      ++ showBytes bytes ++ " @ " ++ showRate bytes gpuTime ++ ", "
                                      ++ D.elapsed gpuTime cpuTime
     in
-    D.timed D.dump_gc msg action
+    D.timed D.dump_gc msg Nothing action
 
diff --git a/Data/Array/Accelerate/CUDA/Array/Sugar.hs b/Data/Array/Accelerate/CUDA/Array/Sugar.hs
--- a/Data/Array/Accelerate/CUDA/Array/Sugar.hs
+++ b/Data/Array/Accelerate/CUDA/Array/Sugar.hs
@@ -1,7 +1,7 @@
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.Array.Sugar
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
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,12 +1,12 @@
 {-# LANGUAGE BangPatterns        #-}
-{-# LANGUAGE CPP                 #-}
 {-# LANGUAGE GADTs               #-}
 {-# LANGUAGE PatternGuards       #-}
 {-# LANGUAGE ScopedTypeVariables #-}
+{-# LANGUAGE TemplateHaskell     #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.Array.Table
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
@@ -41,15 +41,14 @@
 import qualified Foreign.CUDA.Driver                            as CUDA
 import qualified Data.HashTable.IO                              as HT
 
+import Data.Array.Accelerate.Error                              ( internalError )
 import Data.Array.Accelerate.Array.Data                         ( ArrayData )
 import Data.Array.Accelerate.CUDA.Context                       ( Context, weakContext, deviceContext )
 import Data.Array.Accelerate.CUDA.Array.Nursery                 ( Nursery(..), NRS )
 import qualified Data.Array.Accelerate.CUDA.Array.Nursery       as N
 import qualified Data.Array.Accelerate.CUDA.Debug               as D
 
-#include "accelerate.h"
 
-
 -- We use an MVar to the hash table, so that several threads may safely access
 -- it concurrently. This includes the finalisation threads that remove entries
 -- from the table.
@@ -125,7 +124,7 @@
       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"
+               | otherwise         -> $internalError "lookup" $ "type mismatch"
 
         -- Note: [Weak pointer weirdness]
         --
@@ -140,7 +139,7 @@
         -- above in the error message.
         --
         Nothing                    ->
-          makeStableArray ctx arr >>= \x -> INTERNAL_ERROR(error) "lookup" $ "dead weak pair: " ++ show x
+          makeStableArray ctx arr >>= \x -> $internalError "lookup" $ "dead weak pair: " ++ show x
 
 
 -- Allocate a new device array to be associated with the given host-side array.
diff --git a/Data/Array/Accelerate/CUDA/Async.hs b/Data/Array/Accelerate/CUDA/Async.hs
--- a/Data/Array/Accelerate/CUDA/Async.hs
+++ b/Data/Array/Accelerate/CUDA/Async.hs
@@ -1,7 +1,7 @@
 {-# LANGUAGE CPP #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.Async
--- Copyright   : [2009..2013] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2009..2014] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
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
@@ -1,14 +1,14 @@
-{-# LANGUAGE CPP                 #-}
 {-# LANGUAGE GADTs               #-}
 {-# LANGUAGE PatternGuards       #-}
 {-# LANGUAGE QuasiQuotes         #-}
 {-# LANGUAGE RecordWildCards     #-}
 {-# LANGUAGE ScopedTypeVariables #-}
+{-# LANGUAGE TemplateHaskell     #-}
 {-# OPTIONS -fno-warn-name-shadowing #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.CodeGen
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
@@ -35,6 +35,7 @@
 import qualified Data.HashSet                                   as Set
 
 -- friends
+import Data.Array.Accelerate.Error
 import Data.Array.Accelerate.Type
 import Data.Array.Accelerate.Tuple
 import Data.Array.Accelerate.Trafo
@@ -56,9 +57,7 @@
 import Data.Array.Accelerate.CUDA.CodeGen.Stencil
 import Data.Array.Accelerate.CUDA.Foreign.Import                ( canExecuteExp )
 
-#include "accelerate.h"
 
-
 -- Local environments
 --
 data Val env where
@@ -68,7 +67,7 @@
 prj :: Idx env t -> Val env -> [C.Exp]
 prj ZeroIdx      (Push _   v) = v
 prj (SuccIdx ix) (Push val _) = prj ix val
-prj _            _            = INTERNAL_ERROR(error) "prj" "inconsistent valuation"
+prj _            _            = $internalError "prj" "inconsistent valuation"
 
 
 -- Array expressions
@@ -87,7 +86,7 @@
 -- TODO: include a measure of how much shared memory a kernel requires.
 --
 codegenAcc :: forall aenv arrs. DeviceProperties -> DelayedOpenAcc aenv arrs -> Gamma aenv -> [ CUTranslSkel aenv arrs ]
-codegenAcc _   Delayed{}       _    = INTERNAL_ERROR(error) "codegenAcc" "expected manifest array"
+codegenAcc _   Delayed{}       _    = $internalError "codegenAcc" "expected manifest array"
 codegenAcc dev (Manifest pacc) aenv
   = codegen
   $ case pacc of
@@ -144,7 +143,7 @@
 
     -- code generation for delayed arrays
     travD :: (Shape sh, Elt e) => DelayedOpenAcc aenv (Array sh e) -> CUDA (CUDelayedAcc aenv sh e)
-    travD Manifest{}  = INTERNAL_ERROR(error) "codegenAcc" "expected delayed array"
+    travD Manifest{}  = $internalError "codegenAcc" "expected delayed array"
     travD Delayed{..} = CUDelayed <$> travE extentD
                                   <*> travF1 indexD
                                   <*> travF1 linearIndexD
@@ -169,8 +168,8 @@
     -- caffeine and misery
     prim :: String
     prim                = showPreAccOp pacc
-    unexpectedError     = INTERNAL_ERROR(error) "codegenAcc" $ "unexpected array primitive: " ++ prim
-    fusionError         = INTERNAL_ERROR(error) "codegenAcc" $ "unexpected fusible material: " ++ prim
+    unexpectedError     = $internalError "codegenAcc" $ "unexpected array primitive: " ++ prim
+    fusionError         = $internalError "codegenAcc" $ "unexpected fusible material: " ++ prim
 
 
 -- Scalar function abstraction
@@ -214,7 +213,7 @@
              $ \xs -> evalState (evalCGM (go xs)) n
   --
   | otherwise
-  = INTERNAL_ERROR(error) "codegenFun1" "expected unary function"
+  = $internalError "codegenFun1" "expected unary function"
 
 
 codegenFun2
@@ -240,7 +239,7 @@
              $ \xs ys -> evalState (evalCGM (go xs ys)) n
   --
   | otherwise
-  = INTERNAL_ERROR(error) "codegenFun2" "expected binary function"
+  = $internalError "codegenFun2" "expected binary function"
 
 
 -- It is important to filter output terms of a function that will not be used.
@@ -316,7 +315,7 @@
         Var ix                  -> return $ prj ix env
         PrimConst c             -> return $ [codegenPrimConst c]
         Const c                 -> return $ codegenConst (Sugar.eltType (undefined::t)) c
-        PrimApp f arg           -> return . codegenPrim f <$> cvtE arg env
+        PrimApp f x             -> return <$> primApp f x env
         Tuple t                 -> cvtT t env
         Prj i t                 -> prjT i t exp env
         Cond p t e              -> cond p t e env
@@ -361,6 +360,25 @@
       body'     <- cvtE body (env `Push` bnd')
       return body'
 
+    -- When evaluating primitive functions, we evaluate each argument to the
+    -- operation as a statement expression. This is necessary to ensure proper
+    -- short-circuit behaviour for logical operations.
+    --
+    primApp :: PrimFun (a -> b) -> DelayedOpenExp env aenv a -> Val env -> Gen C.Exp
+    primApp f x env
+      | Tuple (NilTup `SnocTup` a `SnocTup` b) <- x
+      = codegenPrim2 f <$> cvtE' a env <*> cvtE' b env
+
+      | otherwise
+      = codegenPrim1 f <$> cvtE' x env
+      where
+        cvtE' :: DelayedOpenExp env aenv a -> Val env -> Gen C.Exp
+        cvtE' e env = do
+          (b,r) <- clean $ single "primApp" <$> cvtE e env
+          if null b
+             then return r
+             else return [cexp| ({ $items:b; $exp:r; }) |]
+
     -- Convert an open expression into a sequence of C expressions. We retain
     -- snoc-list ordering, so the element at tuple index zero is at the end of
     -- the list. Note that nested tuple structures are flattened.
@@ -395,27 +413,45 @@
     prjToInt :: TupleIdx t e -> TupleType a -> Int
     prjToInt ZeroTupIdx     _                 = 0
     prjToInt (SuccTupIdx i) (b `PairTuple` a) = sizeTupleType a + prjToInt i b
-    prjToInt _              _                 = INTERNAL_ERROR(error) "prjToInt" "inconsistent valuation"
+    prjToInt _              _                 = $internalError "prjToInt" "inconsistent valuation"
 
     sizeTupleType :: TupleType a -> Int
     sizeTupleType UnitTuple       = 0
     sizeTupleType (SingleTuple _) = 1
     sizeTupleType (PairTuple a b) = sizeTupleType a + sizeTupleType b
 
-    -- Scalar conditionals. To keep the return type as an expression list we use
-    -- the ternery C condition operator (?:). For tuples this is not
-    -- particularly good, so the least we can do is make sure the predicate
-    -- result is evaluated only once and bound to a local variable.
+    -- Scalar conditionals insert a standard if/else statement block. We don't
+    -- use the ternary expression operator (?:) because this forces all
+    -- auxiliary bindings for both the true and false branches to always be
+    -- evaluated before the correct result is chosen.
     --
-    cond :: DelayedOpenExp env aenv Bool
+    cond :: forall env t. Elt t
+         => DelayedOpenExp env aenv Bool
          -> DelayedOpenExp env aenv t
          -> DelayedOpenExp env aenv t
          -> Val env -> Gen [C.Exp]
-    cond p t e env = do
+    cond p t f env = do
       p'        <- cvtE p env
       ok        <- single "Cond" <$> pushEnv p p'
-      zipWith (\a b -> [cexp| $exp:ok ? $exp:a : $exp:b |]) <$> cvtE t env <*> cvtE e env
+      ifTrue    <- clean $ cvtE t env
+      ifFalse   <- clean $ cvtE f env
 
+      -- Generate names for the result variables, which will be initialised
+      -- within each branch of the conditional. Twiddle the names a bit to
+      -- avoid clobbering.
+      var_r     <- lift fresh
+      let (_, r, declr) = locals ('l':var_r) (undefined :: t)
+          branch        = [citem| if ( $exp:ok ) {
+                                      $items:(r .=. ifTrue)
+                                  }
+                                  else {
+                                      $items:(r .=. ifFalse)
+                                  } |]
+                        : map C.BlockDecl declr
+
+      modify (\s -> s { localBindings = branch ++ localBindings s })
+      return r
+
     -- Value recursion
     --
     while :: forall env a. Elt a
@@ -435,9 +471,11 @@
            x'           <- cvtE x env
            var_acc      <- lift fresh
            var_ok       <- lift fresh
+           var_tmp      <- lift fresh
 
-           let (tn_acc, acc, _)  = locals ('l':var_acc) (undefined :: a)
-               (tn_ok,  ok,  _)  = locals ('l':var_ok)  (undefined :: Bool)
+           let (_, acc, decl_acc) = locals ('l':var_acc) (undefined :: a)
+               (_, ok,  decl_ok)  = locals ('l':var_ok)  (undefined :: Bool)
+               (tmp, _, _)        = locals ('l':var_tmp) (undefined :: a)
 
            -- Generate code for the predicate and body expressions, with the new
            -- names baked in directly. We can't use 'codegenFun1', because
@@ -446,25 +484,24 @@
            -- However, we do need to generate the function with a clean set of
            -- local bindings, and extract and new declarations afterwards.
            --
-           let cvtF :: forall env t. Elt t => DelayedOpenExp env aenv t -> Val env -> Gen ([C.BlockItem], [C.Exp])
-               cvtF e env = do
-                 old  <- state (\s -> ( localBindings s, s { localBindings = []  } ))
-                 e'   <- cvtE e env
-                 env' <- state (\s -> ( localBindings s, s { localBindings = old } ))
-                 return (reverse env', e')
-
-           p'   <- cvtF p (env `Push` acc)
-           f'   <- cvtF f (env `Push` acc)
+           p'   <- clean $ cvtE p (env `Push` acc)
+           f'   <- clean $ cvtE f (env `Push` acc)
 
            -- Piece it all together. Note that declarations are added to the
-           -- localBindings in reverse order.
+           -- localBindings in reverse order. Also, we have to be careful not to
+           -- assign the results of f' direction into acc. Why? If some of the
+           -- variables in acc are referenced in f', then we risk overwriting
+           -- values that are still needed to computer f'.
+           --
            let loop = [citem| while ( $exp:(single "while" ok) ) {
-                                  $items:(acc .=. f')
+                                  $items:(tmp .=. f')
+                                  $items:(acc .=. tmp)
                                   $items:(ok  .=. p')
                               } |]
-                    : (ok .=. p')
-                   ++ map     (\(t,n)   -> [citem| $ty:t $id:n ; |])      tn_ok
-                   ++ zipWith (\(t,n) v -> [citem| $ty:t $id:n = $v ; |]) tn_acc x'
+                    : reverse (ok  .=. p')
+                   ++ reverse (acc .=. x')
+                   ++ map C.BlockDecl decl_ok
+                   ++ map C.BlockDecl decl_acc
 
            modify (\s -> s { localBindings = loop ++ localBindings s })
            return acc
@@ -486,7 +523,7 @@
           restrict SliceNil              _       _       = []
           restrict (SliceAll   sliceIdx) slx     (sz:sl) = sz : restrict sliceIdx slx sl
           restrict (SliceFixed sliceIdx) (_:slx) ( _:sl) =      restrict sliceIdx slx sl
-          restrict _ _ _ = INTERNAL_ERROR(error) "IndexSlice" "unexpected shapes"
+          restrict _ _ _ = $internalError "IndexSlice" "unexpected shapes"
           --
           slice slix' sh' = reverse $ restrict sliceIndex (reverse slix') (reverse sh')
       in
@@ -505,7 +542,7 @@
           extend SliceNil              _        _       = []
           extend (SliceAll   sliceIdx) slx      (sz:sh) = sz : extend sliceIdx slx sh
           extend (SliceFixed sliceIdx) (sz:slx) sh      = sz : extend sliceIdx slx sh
-          extend _ _ _ = INTERNAL_ERROR(error) "IndexFull" "unexpected shapes"
+          extend _ _ _ = $internalError "IndexFull" "unexpected shapes"
           --
           replicate slix' sl' = reverse $ extend sliceIndex (reverse slix') (reverse sl')
       in
@@ -555,7 +592,7 @@
         return   $ zipWith (\t a -> indexArray dev t (cvar a) i) ty arr
       --
       | otherwise
-      = INTERNAL_ERROR(error) "Index" "expected array variable"
+      = $internalError "Index" "expected array variable"
 
 
     linearIndex :: (Shape sh, Elt e)
@@ -573,7 +610,7 @@
         return   $ zipWith (\t a -> indexArray dev t (cvar a) i) ty arr
       --
       | otherwise
-      = INTERNAL_ERROR(error) "LinearIndex" "expected array variable"
+      = $internalError "LinearIndex" "expected array variable"
 
     -- Array shapes created in this method refer to the shape of free array
     -- variables. As such, they are always passed as arguments to the kernel,
@@ -587,7 +624,7 @@
       = return $ cshape (delayedDim acc) (fst (namesOfAvar aenv idx))
 
       | otherwise
-      = INTERNAL_ERROR(error) "Shape" "expected array variable"
+      = $internalError "Shape" "expected array variable"
 
     -- The size of a shape, as the product of the extent in each dimension. The
     -- definition is inlined, but we could also call the C function helpers.
@@ -616,18 +653,28 @@
              -> Val env
              -> Gen [C.Exp]
     foreignE ff x env = case canExecuteExp ff of
-      Nothing      -> INTERNAL_ERROR(error) "codegenOpenExp" "Non-CUDA foreign expression encountered"
+      Nothing      -> $internalError "codegenOpenExp" "Non-CUDA foreign expression encountered"
       Just (hs, f) -> do
         lift $ modify (\st -> st { headers = foldl (flip Set.insert) (headers st) hs })
         args    <- cvtE x env
         mapM_ use args
         return  $  [ccall f (ccastTup (Sugar.eltType (undefined::a)) args)]
 
+    -- Execute a command in a new environment. The old environment is replaced
+    -- on exit, and the result and any new bindings generated are returned.
+    --
+    clean :: Gen a -> Gen ([C.BlockItem], a)
+    clean this = do
+      env  <- state (\s -> ( localBindings s, s { localBindings = []  } ))
+      r    <- this
+      env' <- state (\s -> ( localBindings s, s { localBindings = env } ))
+      return (reverse env', r)
+
     -- Some terms demand we extract only singly typed expressions
     --
     single :: String -> [C.Exp] -> C.Exp
     single _   [x] = x
-    single loc _   = INTERNAL_ERROR(error) loc "expected single expression"
+    single loc _   = $internalError loc "expected single expression"
 
 
 -- Scalar Primitives
@@ -639,65 +686,67 @@
 codegenPrimConst (PrimPi       ty) = codegenPi ty
 
 
-codegenPrim :: PrimFun p -> [C.Exp] -> C.Exp
-codegenPrim (PrimAdd              _) [a,b] = [cexp|$exp:a + $exp:b|]
-codegenPrim (PrimSub              _) [a,b] = [cexp|$exp:a - $exp:b|]
-codegenPrim (PrimMul              _) [a,b] = [cexp|$exp:a * $exp:b|]
-codegenPrim (PrimNeg              _) [a]   = [cexp| - $exp:a|]
-codegenPrim (PrimAbs             ty) [a]   = codegenAbs ty a
-codegenPrim (PrimSig             ty) [a]   = codegenSig ty a
-codegenPrim (PrimQuot             _) [a,b] = [cexp|$exp:a / $exp:b|]
-codegenPrim (PrimRem              _) [a,b] = [cexp|$exp:a % $exp:b|]
-codegenPrim (PrimIDiv             _) [a,b] = ccall "idiv" [a,b]
-codegenPrim (PrimMod              _) [a,b] = ccall "mod"  [a,b]
-codegenPrim (PrimBAnd             _) [a,b] = [cexp|$exp:a & $exp:b|]
-codegenPrim (PrimBOr              _) [a,b] = [cexp|$exp:a | $exp:b|]
-codegenPrim (PrimBXor             _) [a,b] = [cexp|$exp:a ^ $exp:b|]
-codegenPrim (PrimBNot             _) [a]   = [cexp|~ $exp:a|]
-codegenPrim (PrimBShiftL          _) [a,b] = [cexp|$exp:a << $exp:b|]
-codegenPrim (PrimBShiftR          _) [a,b] = [cexp|$exp:a >> $exp:b|]
-codegenPrim (PrimBRotateL         _) [a,b] = ccall "rotateL" [a,b]
-codegenPrim (PrimBRotateR         _) [a,b] = ccall "rotateR" [a,b]
-codegenPrim (PrimFDiv             _) [a,b] = [cexp|$exp:a / $exp:b|]
-codegenPrim (PrimRecip           ty) [a]   = codegenRecip ty a
-codegenPrim (PrimSin             ty) [a]   = ccall (FloatingNumType ty `postfix` "sin")   [a]
-codegenPrim (PrimCos             ty) [a]   = ccall (FloatingNumType ty `postfix` "cos")   [a]
-codegenPrim (PrimTan             ty) [a]   = ccall (FloatingNumType ty `postfix` "tan")   [a]
-codegenPrim (PrimAsin            ty) [a]   = ccall (FloatingNumType ty `postfix` "asin")  [a]
-codegenPrim (PrimAcos            ty) [a]   = ccall (FloatingNumType ty `postfix` "acos")  [a]
-codegenPrim (PrimAtan            ty) [a]   = ccall (FloatingNumType ty `postfix` "atan")  [a]
-codegenPrim (PrimAsinh           ty) [a]   = ccall (FloatingNumType ty `postfix` "asinh") [a]
-codegenPrim (PrimAcosh           ty) [a]   = ccall (FloatingNumType ty `postfix` "acosh") [a]
-codegenPrim (PrimAtanh           ty) [a]   = ccall (FloatingNumType ty `postfix` "atanh") [a]
-codegenPrim (PrimExpFloating     ty) [a]   = ccall (FloatingNumType ty `postfix` "exp")   [a]
-codegenPrim (PrimSqrt            ty) [a]   = ccall (FloatingNumType ty `postfix` "sqrt")  [a]
-codegenPrim (PrimLog             ty) [a]   = ccall (FloatingNumType ty `postfix` "log")   [a]
-codegenPrim (PrimFPow            ty) [a,b] = ccall (FloatingNumType ty `postfix` "pow")   [a,b]
-codegenPrim (PrimLogBase         ty) [a,b] = codegenLogBase ty a b
-codegenPrim (PrimTruncate     ta tb) [a]   = codegenTruncate ta tb a
-codegenPrim (PrimRound        ta tb) [a]   = codegenRound ta tb a
-codegenPrim (PrimFloor        ta tb) [a]   = codegenFloor ta tb a
-codegenPrim (PrimCeiling      ta tb) [a]   = codegenCeiling ta tb a
-codegenPrim (PrimAtan2           ty) [a,b] = ccall (FloatingNumType ty `postfix` "atan2") [a,b]
-codegenPrim (PrimLt               _) [a,b] = [cexp|$exp:a < $exp:b|]
-codegenPrim (PrimGt               _) [a,b] = [cexp|$exp:a > $exp:b|]
-codegenPrim (PrimLtEq             _) [a,b] = [cexp|$exp:a <= $exp:b|]
-codegenPrim (PrimGtEq             _) [a,b] = [cexp|$exp:a >= $exp:b|]
-codegenPrim (PrimEq               _) [a,b] = [cexp|$exp:a == $exp:b|]
-codegenPrim (PrimNEq              _) [a,b] = [cexp|$exp:a != $exp:b|]
-codegenPrim (PrimMax             ty) [a,b] = codegenMax ty a b
-codegenPrim (PrimMin             ty) [a,b] = codegenMin ty a b
-codegenPrim PrimLAnd                 [a,b] = [cexp|$exp:a && $exp:b|]
-codegenPrim PrimLOr                  [a,b] = [cexp|$exp:a || $exp:b|]
-codegenPrim PrimLNot                 [a]   = [cexp| ! $exp:a|]
-codegenPrim PrimOrd                  [a]   = codegenOrd a
-codegenPrim PrimChr                  [a]   = codegenChr a
-codegenPrim PrimBoolToInt            [a]   = codegenBoolToInt a
-codegenPrim (PrimFromIntegral ta tb) [a]   = codegenFromIntegral ta tb a
+codegenPrim1 :: PrimFun f -> C.Exp -> C.Exp
+codegenPrim1 (PrimNeg              _) a   = [cexp| - $exp:a|]
+codegenPrim1 (PrimAbs             ty) a   = codegenAbs ty a
+codegenPrim1 (PrimSig             ty) a   = codegenSig ty a
+codegenPrim1 (PrimBNot             _) a   = [cexp|~ $exp:a|]
+codegenPrim1 (PrimRecip           ty) a   = codegenRecip ty a
+codegenPrim1 (PrimSin             ty) a   = ccall (FloatingNumType ty `postfix` "sin")   [a]
+codegenPrim1 (PrimCos             ty) a   = ccall (FloatingNumType ty `postfix` "cos")   [a]
+codegenPrim1 (PrimTan             ty) a   = ccall (FloatingNumType ty `postfix` "tan")   [a]
+codegenPrim1 (PrimAsin            ty) a   = ccall (FloatingNumType ty `postfix` "asin")  [a]
+codegenPrim1 (PrimAcos            ty) a   = ccall (FloatingNumType ty `postfix` "acos")  [a]
+codegenPrim1 (PrimAtan            ty) a   = ccall (FloatingNumType ty `postfix` "atan")  [a]
+codegenPrim1 (PrimAsinh           ty) a   = ccall (FloatingNumType ty `postfix` "asinh") [a]
+codegenPrim1 (PrimAcosh           ty) a   = ccall (FloatingNumType ty `postfix` "acosh") [a]
+codegenPrim1 (PrimAtanh           ty) a   = ccall (FloatingNumType ty `postfix` "atanh") [a]
+codegenPrim1 (PrimExpFloating     ty) a   = ccall (FloatingNumType ty `postfix` "exp")   [a]
+codegenPrim1 (PrimSqrt            ty) a   = ccall (FloatingNumType ty `postfix` "sqrt")  [a]
+codegenPrim1 (PrimLog             ty) a   = ccall (FloatingNumType ty `postfix` "log")   [a]
+codegenPrim1 (PrimTruncate     ta tb) a   = codegenTruncate ta tb a
+codegenPrim1 (PrimRound        ta tb) a   = codegenRound ta tb a
+codegenPrim1 (PrimFloor        ta tb) a   = codegenFloor ta tb a
+codegenPrim1 (PrimCeiling      ta tb) a   = codegenCeiling ta tb a
+codegenPrim1 PrimLNot                 a   = [cexp| ! $exp:a|]
+codegenPrim1 PrimOrd                  a   = codegenOrd a
+codegenPrim1 PrimChr                  a   = codegenChr a
+codegenPrim1 PrimBoolToInt            a   = codegenBoolToInt a
+codegenPrim1 (PrimFromIntegral ta tb) a   = codegenFromIntegral ta tb a
+codegenPrim1 _ _ =
+  $internalError "codegenPrim1" "unknown primitive function"
 
--- If the argument lists are not the correct length
-codegenPrim _ _ =
-  INTERNAL_ERROR(error) "codegenPrim" "inconsistent valuation"
+codegenPrim2 :: PrimFun f -> C.Exp -> C.Exp -> C.Exp
+codegenPrim2 (PrimAdd              _) a b = [cexp|$exp:a + $exp:b|]
+codegenPrim2 (PrimSub              _) a b = [cexp|$exp:a - $exp:b|]
+codegenPrim2 (PrimMul              _) a b = [cexp|$exp:a * $exp:b|]
+codegenPrim2 (PrimQuot             _) a b = [cexp|$exp:a / $exp:b|]
+codegenPrim2 (PrimRem              _) a b = [cexp|$exp:a % $exp:b|]
+codegenPrim2 (PrimIDiv             _) a b = ccall "idiv" [a,b]
+codegenPrim2 (PrimMod              _) a b = ccall "mod"  [a,b]
+codegenPrim2 (PrimBAnd             _) a b = [cexp|$exp:a & $exp:b|]
+codegenPrim2 (PrimBOr              _) a b = [cexp|$exp:a | $exp:b|]
+codegenPrim2 (PrimBXor             _) a b = [cexp|$exp:a ^ $exp:b|]
+codegenPrim2 (PrimBShiftL          _) a b = [cexp|$exp:a << $exp:b|]
+codegenPrim2 (PrimBShiftR          _) a b = [cexp|$exp:a >> $exp:b|]
+codegenPrim2 (PrimBRotateL         _) a b = ccall "rotateL" [a,b]
+codegenPrim2 (PrimBRotateR         _) a b = ccall "rotateR" [a,b]
+codegenPrim2 (PrimFDiv             _) a b = [cexp|$exp:a / $exp:b|]
+codegenPrim2 (PrimFPow            ty) a b = ccall (FloatingNumType ty `postfix` "pow")   [a,b]
+codegenPrim2 (PrimLogBase         ty) a b = codegenLogBase ty a b
+codegenPrim2 (PrimAtan2           ty) a b = ccall (FloatingNumType ty `postfix` "atan2") [a,b]
+codegenPrim2 (PrimLt               _) a b = [cexp|$exp:a < $exp:b|]
+codegenPrim2 (PrimGt               _) a b = [cexp|$exp:a > $exp:b|]
+codegenPrim2 (PrimLtEq             _) a b = [cexp|$exp:a <= $exp:b|]
+codegenPrim2 (PrimGtEq             _) a b = [cexp|$exp:a >= $exp:b|]
+codegenPrim2 (PrimEq               _) a b = [cexp|$exp:a == $exp:b|]
+codegenPrim2 (PrimNEq              _) a b = [cexp|$exp:a != $exp:b|]
+codegenPrim2 (PrimMax             ty) a b = codegenMax ty a b
+codegenPrim2 (PrimMin             ty) a b = codegenMin ty a b
+codegenPrim2 PrimLAnd                 a b = [cexp|$exp:a && $exp:b|]
+codegenPrim2 PrimLOr                  a b = [cexp|$exp:a || $exp:b|]
+codegenPrim2 _ _ _ =
+  $internalError "codegenPrim2" "unknown primitive function"
 
 -- Implementation of scalar primitives
 --
@@ -868,7 +917,7 @@
                                          (ls, xs' ) = travTup l xs
                                          (rs, xs'') = travTup r xs'
                                        in (ls ++ rs, xs'')
-    travTup _ _                      = INTERNAL_ERROR(error) "ccastTup" "not enough expressions to match type"
+    travTup _ _                      = $internalError "ccastTup" "not enough expressions to match type"
 
 
 postfix :: NumType a -> String -> String
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
@@ -1,4 +1,3 @@
-{-# LANGUAGE CPP                   #-}
 {-# LANGUAGE FlexibleInstances     #-}
 {-# LANGUAGE GADTs                 #-}
 {-# LANGUAGE ImpredicativeTypes    #-}
@@ -7,10 +6,11 @@
 {-# LANGUAGE PatternGuards         #-}
 {-# LANGUAGE QuasiQuotes           #-}
 {-# LANGUAGE ScopedTypeVariables   #-}
+{-# LANGUAGE TemplateHaskell       #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.CodeGen.Base
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
@@ -49,13 +49,12 @@
 
 -- friends
 import Data.Array.Accelerate.Type
+import Data.Array.Accelerate.Error
 import Data.Array.Accelerate.Array.Sugar                ( Array, Shape, Elt )
 import Data.Array.Accelerate.Analysis.Shape
 import Data.Array.Accelerate.CUDA.CodeGen.Type
 import Data.Array.Accelerate.CUDA.AST
 
-#include "accelerate.h"
-
 -- Names
 -- -----
 
@@ -183,7 +182,7 @@
     toIndex []      []     = [cexp| $int:(0::Int) |]
     toIndex [_]     [i]    = i
     toIndex (sz:sh) (i:ix) = [cexp| $exp:(toIndex sh ix) * $exp:sz + $exp:i |]
-    toIndex _       _      = INTERNAL_ERROR(error) "toIndex" "argument mismatch"
+    toIndex _       _      = $internalError "toIndex" "argument mismatch"
 
 -- Generate code to calculate a multi-dimensional index from a linear index and a given array shape.
 -- This version creates temporary values that are reused in the computation.
@@ -256,7 +255,7 @@
         sh'             = cshape dim sh
         get ix          = ([], map (\a -> [cexp| $id:a [ $exp:ix ] |]) arrs)
         manifest        = CUDelayed (CUExp ([], sh'))
-                                    (INTERNAL_ERROR(error) "readArray" "linear indexing only")
+                                    ($internalError "readArray" "linear indexing only")
                                     (CUFun1 (zip (repeat True)) (\[i] -> get (rvalue i)))
     in ( args, manifest )
 
@@ -388,6 +387,27 @@
 class Lvalue a where
   lvalue :: a -> C.Exp -> C.BlockItem
 
+-- Note: [Mutable l-values]
+--
+-- Be careful when using mutable l-values that the same variable does not appear
+-- on both the left and right hand side. For example, the following will lead to
+-- problems (#114, #168):
+--
+--   $items:(x .=. f x)
+--
+--   $items:(y .=. combine x y)
+--
+-- If 'x' and 'y' represent values with tuple types, they will have multiple
+-- components. Since the LHS is updated as the new values are calculated, it is
+-- possible to get into a situation where computing the new value for some of
+-- the components will be using the updated values, not the original values.
+--
+-- Instead, store the result to some (temporary) variable that does not appear
+-- on the RHS, and then update old value using the fully computed result, e.g.:
+--
+--   $items:(x' .=. f x)
+--   $items:(x  .=. x')
+--
 instance Lvalue C.Exp where
   lvalue x y = [citem| $exp:x = $exp:y; |]
 
@@ -426,7 +446,7 @@
 instance Assign l r => Assign [l] [r] where
   assign []     []     = []
   assign (x:xs) (y:ys) = assign x y ++ assign xs ys
-  assign _      _      = INTERNAL_ERROR(error) ".=." "argument mismatch"
+  assign _      _      = $internalError ".=." "argument mismatch"
 
 instance Assign l r => Assign l ([C.BlockItem], r) where
   assign lhs (env, rhs) = env ++ assign lhs rhs
@@ -440,10 +460,10 @@
 zipWith :: (a -> b -> c) -> [a] -> [b] -> [c]
 zipWith f (x:xs) (y:ys) = f x y : zipWith f xs ys
 zipWith _ []     []     = []
-zipWith _ _      _      = INTERNAL_ERROR(error) "zipWith" "argument mismatch"
+zipWith _ _      _      = $internalError "zipWith" "argument mismatch"
 
 zipWith3 :: (a -> b -> c -> d) -> [a] -> [b] -> [c] -> [d]
 zipWith3 f (x:xs) (y:ys) (z:zs) = f x y z : zipWith3 f xs ys zs
 zipWith3 _ []     []     []     = []
-zipWith3 _ _      _      _      = INTERNAL_ERROR(error) "zipWith3" "argument mismatch"
+zipWith3 _ _      _      _      = $internalError "zipWith3" "argument mismatch"
 
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
@@ -1,12 +1,14 @@
-{-# LANGUAGE CPP                 #-}
 {-# LANGUAGE GADTs               #-}
+{-# LANGUAGE ImpredicativeTypes  #-}
 {-# LANGUAGE PatternGuards       #-}
 {-# LANGUAGE QuasiQuotes         #-}
 {-# LANGUAGE ScopedTypeVariables #-}
+{-# LANGUAGE TemplateHaskell     #-}
+{-# LANGUAGE ViewPatterns        #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.CodeGen.IndexSpace
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
@@ -24,19 +26,16 @@
 
 ) where
 
-import Data.List
 import Language.C.Quote.CUDA
 import Foreign.CUDA.Analysis.Device
 import qualified Language.C.Syntax                      as C
 
 import Data.Array.Accelerate.Array.Sugar                ( Array, Shape, Elt, ignore, shapeToList )
+import Data.Array.Accelerate.Error                      ( internalError )
 import Data.Array.Accelerate.CUDA.AST                   ( Gamma )
-import Data.Array.Accelerate.CUDA.CodeGen.Type
 import Data.Array.Accelerate.CUDA.CodeGen.Base
 
-#include "accelerate.h"
 
-
 -- Construct a new array by applying a function to each index. Each thread
 -- processes multiple elements, striding the array by the grid size.
 --
@@ -178,7 +177,8 @@
     permute
     (
         $params:argIn,
-        $params:argOut
+        $params:argOut,
+        typename Int32 * __restrict__ lock
     )
     {
         /*
@@ -199,25 +199,22 @@
 
             if ( ! $exp:(cignore dst) )
             {
-                $decls:decly
-                $decls:decly'
-
-                $items:(jx      .=. ctoIndex shOut dst)
-                $items:(dce_x x .=. get ix)
-                $items:(dce_y y .=. arrOut jx)
+                $items:(jx        .=. ctoIndex shOut dst)
+                $items:(dce_x x   .=. get ix)
 
-                $items:write
+                $items:(atomically jx
+                    [ dce_y y   .=. setOut jx
+                    , setOut jx .=. combine x y ]
+                )
             }
         }
     }
   |]
   where
-    sizeof                      = eltSizeOf (undefined::e)
     (texIn, argIn)              = environment dev aenv
-    (argOut, shOut, arrOut)     = writeArray "Out" (undefined :: Array sh' e)
+    (argOut, shOut, setOut)     = writeArray "Out" (undefined :: Array sh' e)
     (x, _, _)                   = locals "x" (undefined :: e)
-    (_, y,  decly)              = locals "y" (undefined :: e)
-    (_, y', decly')             = locals "_y" (undefined :: e)
+    (y, _, _)                   = locals "y" (undefined :: e)
     (sh, _, _)                  = locals "shIn" (undefined :: sh)
     (src, _, _)                 = locals "sh" (undefined :: sh)
     (dst, _, _)                 = locals "sh_" (undefined :: sh')
@@ -227,43 +224,79 @@
 
     -- If the destination index resolves to the magic index "ignore", the result
     -- is dropped from the output array.
+    --
     cignore :: Rvalue x => [x] -> C.Exp
-    cignore []  = INTERNAL_ERROR(error) "permute" "singleton arrays not supported"
+    cignore []  = $internalError "permute" "singleton arrays not supported"
     cignore xs  = foldl1 (\a b -> [cexp| $exp:a && $exp:b |])
                 $ zipWith (\a b -> [cexp| $exp:(rvalue a) == $int:b |]) xs
                 $ shapeToList (ignore :: sh')
 
-    -- Apply the combining function between old and new values. If multiple
-    -- threads attempt to write to the same location, the hardware
-    -- write-combining mechanism will accept one transaction and all other
-    -- updates will be lost.
-    --
-    -- If the hardware supports it, we can use atomicCAS (compare-and-swap) to
-    -- work around this. This requires at least compute 1.1 for 32-bit values,
-    -- and compute 1.2 for 64-bit values. If hardware support is not available,
-    -- write the result as normal and hope for the best.
+    -- If we can determine that the old values are not used in the combination
+    -- function (e.g. filter) then the lock and unlock fragments can be replaced
+    -- with a NOP.
     --
-    -- Each element of a tuple is necessarily written individually, so the tuple
-    -- as a whole is not stored atomically.
+    -- If locking is required but the hardware does not support it (compute 1.0)
+    -- then we issue a runtime error immediately instead of silently failing.
     --
-    write       = env ++ zipWith6 apply sizeof (arrOut jx) fun x (dce_y y) y'
-    (env, fun)  = combine x y
-
-    apply size out f x1 (used,y1) y1'
-      | used
-      , Just atomicCAS <- reinterpret size
-      = [citem| do {
-                       $exp:y1' = $exp:y1;
-                       $exp:y1  = $exp:atomicCAS ( & $exp:out, $exp:y1', $exp:f );
+    mustLock    = or . fst . unzip $ dce_y y
 
-                   } while ( $exp:y1 != $exp:y1' ); |]
+    -- The atomic section is acquired using a spin lock. This requires a
+    -- temporary array to represent the lock state for each element of the
+    -- output. We use 1 to represent the locked state, and 0 to represent
+    -- unlocked elements.
+    --
+    --   do {
+    --     old = atomicExch(&lock[i], 1);       // atomic exchange
+    --   } while (old == 1);
+    --
+    --   /* critical section */
+    --
+    --   atomicExch(&lock[i], 0);
+    --
+    -- The initial loop repeatedly attempts to take the lock by writing a 1 into
+    -- the slot. Once the 'old' state of the lock returns 0 (unlocked), we have
+    -- just acquired the lock, and the atomic section can be computed. Finally,
+    -- atomically write a 0 back into the slot to unlock the element.
+    --
+    -- However, there is a complication with CUDA devices because all threads in
+    -- the warp must execute in lockstep (with predicated execution). Once a
+    -- thread acquires a lock, then it will be disabled and stop participating
+    -- in the first loop, waiting until all other threads in the warp acquire
+    -- their locks. If two threads in a warp are attempting to acquire the same
+    -- lock, once the lock is acquired by the first thread, it sits idle while
+    -- the second thread spins attempting to grab a lock that will never be
+    -- released, because the first thread can not progress. DEADLOCK.
+    --
+    -- So, we need to invert the algorithm so that threads can always make
+    -- progress, until each thread in the warp has committed their result.
+    --
+    --   done = 0;
+    --   do {
+    --       if (atomicExch(&lock[i], 1) == 0) {
+    --
+    --           /* critical section */
+    --
+    --           done = 1;
+    --           atomicExch(&lock[i], 0);
+    --       }
+    --   } while (done == 0)
+    --
+    atomically :: (C.Type, Name) -> [[C.BlockItem]] -> [C.BlockItem]
+    atomically (_,i) (concat -> body)
+      | not mustLock            = body
+      | sm < Compute 1 1        = $internalError "permute" "Requires at least compute compatibility 1.1"
+      | otherwise               =
+        [ [citem| typename Int32 done = 0; |]
+        , [citem| do {
+                      __threadfence();
 
-      | otherwise
-      = [citem| $exp:out = $exp:(rvalue x1); |]
+                      if ( atomicExch(&lock[ $exp:(cvar i) ], 1) == 0 ) {
+                          $items:body
 
-    --
-    reinterpret :: Int -> Maybe C.Exp
-    reinterpret 4 | sm >= Compute 1 1   = Just [cexp| $id:("atomicCAS32") |]
-    reinterpret 8 | sm >= Compute 1 2   = Just [cexp| $id:("atomicCAS64") |]
-    reinterpret _                       = Nothing
+                          done = 1;
+                          atomicExch(&lock[ $exp:(cvar i) ], 0);
+                      }
+                  } while (done == 0);
+                |]
+        ]
 
diff --git a/Data/Array/Accelerate/CUDA/CodeGen/Mapping.hs b/Data/Array/Accelerate/CUDA/CodeGen/Mapping.hs
--- a/Data/Array/Accelerate/CUDA/CodeGen/Mapping.hs
+++ b/Data/Array/Accelerate/CUDA/CodeGen/Mapping.hs
@@ -1,11 +1,12 @@
 {-# LANGUAGE GADTs               #-}
+{-# LANGUAGE ImpredicativeTypes  #-}
 {-# LANGUAGE PatternGuards       #-}
 {-# LANGUAGE QuasiQuotes         #-}
 {-# LANGUAGE ScopedTypeVariables #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.CodeGen.Mapping
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
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
@@ -2,8 +2,8 @@
 {-# OPTIONS_GHC -fno-warn-orphans #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.CodeGen.Monad
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
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
@@ -1,12 +1,13 @@
 {-# LANGUAGE CPP                 #-}
 {-# LANGUAGE GADTs               #-}
+{-# LANGUAGE ImpredicativeTypes  #-}
 {-# LANGUAGE PatternGuards       #-}
 {-# LANGUAGE QuasiQuotes         #-}
 {-# LANGUAGE ScopedTypeVariables #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.CodeGen.PrefixSum
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
@@ -183,6 +184,7 @@
     )
     {
         $decls:smem
+        $decls:declt
         $decls:declx
         $decls:decly
         $decls:declz
@@ -222,7 +224,8 @@
              * Carry in the result from the privous segment
              */
             if ( $exp:carryIn ) {
-                $items:(x .=. combine z x)
+                $items:(t .=. combine z x)
+                $items:(x .=. t)
             }
 
             /*
@@ -232,7 +235,7 @@
             $items:(sdata "threadIdx.x" .=. x)
             __syncthreads();
 
-            $items:(scanBlock dev fun x y sdata Nothing)
+            $items:(scanBlock dev fun x y t sdata Nothing)
 
             /*
              * Exclusive scans write the result of the previous thread to global
@@ -273,6 +276,7 @@
     (argOut, _, setOut)         = writeArray "Out" (undefined :: Vector e)
     (argSum, _, totalSum)       = writeArray "Sum" (undefined :: Vector e)
     (argBlk, _, blkSum)         = writeArray "Blk" (undefined :: Vector e)
+    (_, t, declt)               = locals "t" (undefined :: e)
     (_, x, declx)               = locals "x" (undefined :: e)
     (_, y, decly)               = locals "y" (undefined :: e)
     (_, z, declz)               = locals "z" (undefined :: e)
@@ -351,6 +355,7 @@
     )
     {
         $decls:smem
+        $decls:declt
         $decls:declx
         $decls:decly
         $items:(sh .=. shIn)
@@ -376,7 +381,8 @@
             $items:(x .=. get ix)
 
             if ( threadIdx.x == 0 && carryIn ) {
-                $items:(x .=. combine y x)
+                $items:(t .=. combine y x)
+                $items:(x .=. t)
             }
 
             /*
@@ -385,7 +391,7 @@
             $items:(sdata "threadIdx.x" .=. x)
             __syncthreads();
 
-            $items:(scanBlock dev fun x y sdata Nothing)
+            $items:(scanBlock dev fun x y t sdata Nothing)
 
             /*
              * Store the final result of the block to be carried in
@@ -411,6 +417,7 @@
     (argOut, _, setOut)         = writeArray "Out" (undefined :: Vector e)
     (_, x, declx)               = locals "x" (undefined :: e)
     (_, y, decly)               = locals "y" (undefined :: e)
+    (_, t, declt)               = locals "t" (undefined :: e)
     (sh, _, _)                  = locals "sh" (undefined :: DIM1)
     (smem, sdata)               = shared (undefined :: e) "sdata" [cexp| blockDim.x |] Nothing
     ix                          = [cvar "ix"]
@@ -443,13 +450,13 @@
     :: forall aenv e. Elt e
     => DeviceProperties
     -> CUFun2 aenv (e -> e -> e)
-    -> [C.Exp] -> [C.Exp]
+    -> [C.Exp] -> [C.Exp] -> [C.Exp]
     -> (Name -> [C.Exp])
     -> Maybe C.Exp
     -> [C.BlockItem]
-scanBlock dev f x0 x1 sdata mlim
+scanBlock dev f x0 x1 x2 sdata mlim
   | shflOK dev (undefined :: e) = error "shfl-scan"
-  | otherwise                   = scanBlockTree dev f x0 x1 sdata mlim
+  | otherwise                   = scanBlockTree dev f x0 x1 x2 sdata mlim
 
 
 -- Use a thread block to scan values in shared memory. Each thread must have
@@ -461,11 +468,11 @@
     :: forall aenv e. Elt e
     => DeviceProperties
     -> CUFun2 aenv (e -> e -> e)
-    -> [C.Exp] -> [C.Exp]               -- temporary variables x0 and x1
+    -> [C.Exp] -> [C.Exp] -> [C.Exp]    -- input variables x0 and x1, plus a temporary to store the intermediate value
     -> (Name -> [C.Exp])                -- index elements from shared memory
     -> Maybe C.Exp                      -- partially full block bounds check?
     -> [C.BlockItem]
-scanBlockTree dev (CUFun2 _ _ f) x0 x1 sdata mlim = map (scan . pow2) [ 0 .. maxThreads ]
+scanBlockTree dev (CUFun2 _ _ f) x0 x1 x2 sdata mlim = map (scan . pow2) [ 0 .. maxThreads ]
   where
     pow2 :: Int -> Int
     pow2 x      = 2 ^ x
@@ -479,7 +486,8 @@
       if ( blockDim.x > $int:n ) {
           if ( $exp:(inrange n) ) {
               $items:(x1 .=. sdata ("threadIdx.x - " ++ show n))
-              $items:(x0 .=. f x1 x0)
+              $items:(x2 .=. f x1 x0)
+              $items:(x0 .=. x2)
           }
           __syncthreads();
           $items:(sdata "threadIdx.x" .=. x0)
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
@@ -1,12 +1,13 @@
-{-# LANGUAGE CPP                 #-}
 {-# LANGUAGE GADTs               #-}
+{-# LANGUAGE ImpredicativeTypes  #-}
 {-# LANGUAGE QuasiQuotes         #-}
 {-# LANGUAGE ScopedTypeVariables #-}
+{-# LANGUAGE TemplateHaskell     #-}
 {-# LANGUAGE TypeOperators       #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.CodeGen.Reduction
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
@@ -24,16 +25,15 @@
 import Language.C.Quote.CUDA
 import qualified Language.C.Syntax                      as C
 
-import Data.Array.Accelerate.Type                       ( IsIntegral )
-import Data.Array.Accelerate.Array.Sugar                ( Array, Shape, Elt, Z(..), (:.)(..) )
 import Data.Array.Accelerate.Analysis.Shape
+import Data.Array.Accelerate.Array.Sugar                ( Array, Shape, Elt, Z(..), (:.)(..) )
+import Data.Array.Accelerate.Error                      ( internalError )
+import Data.Array.Accelerate.Type                       ( IsIntegral )
 import Data.Array.Accelerate.CUDA.AST
 import Data.Array.Accelerate.CUDA.CodeGen.Base
 import Data.Array.Accelerate.CUDA.CodeGen.Type
 
-#include "accelerate.h"
 
-
 -- Reduce an array along the innermost dimension. The function must be
 -- associative to enable efficient parallel implementation.
 --
@@ -125,6 +125,7 @@
         $decls:smem
         $decls:declx
         $decls:decly
+        $decls:declz
 
         $items:(sh .=. shIn)
 
@@ -159,7 +160,8 @@
             for ( ix += gridSize; ix < shapeSize; ix += gridSize )
             {
                 $items:(x .=. get ix)
-                $items:(y .=. combine x y)
+                $items:(z .=. combine y x)
+                $items:(y .=. z)
             }
         }
 
@@ -170,7 +172,7 @@
         $items:(sdata "threadIdx.x" .=. y)
 
         ix = min(shapeSize - blockIdx.x * blockDim.x, blockDim.x);
-        $items:(reduceBlock dev fun x y sdata (cvar "ix"))
+        $items:(reduceBlock dev fun x y z sdata (cvar "ix"))
 
         /*
          * Write the results of this block back to global memory. If we are the last
@@ -192,6 +194,7 @@
 
     (_, x, declx)               = locals "x" (undefined :: e)
     (_, y, decly)               = locals "y" (undefined :: e)
+    (_, z, declz)               = locals "z" (undefined :: e)
     (sh, _, _)                  = locals "sh" (undefined :: sh :. Int)
     ix                          = [cvar "ix"]
     (smem, sdata)               = shared (undefined :: e) "sdata" [cexp| blockDim.x |] Nothing
@@ -201,7 +204,8 @@
       if ( shapeSize > 0 ) {
           if ( gridDim.x == 1 ) {
               $items:(x .=. seed)
-              $items:(y .=. combine x y)
+              $items:(z .=. combine y x)
+              $items:(y .=. z)
           }
           $items:(setOut "blockIdx.x" .=. y)
       }
@@ -239,6 +243,7 @@
         $decls:smem
         $decls:declx
         $decls:decly
+        $decls:declz
 
         $items:(sh .=. shIn)
 
@@ -291,7 +296,8 @@
                     $items:(x .=. get [cvar "ix + blockDim.x"])
 
                     if ( ix >= start ) {
-                        $items:(y .=. combine x y)
+                        $items:(z .=. combine y x)
+                        $items:(y .=. z)
                     }
                     else {
                         $items:(y .=. x)
@@ -304,7 +310,8 @@
                 for ( ix += 2 * blockDim.x; ix < end; ix += blockDim.x )
                 {
                     $items:(x .=. get ix)
-                    $items:(y .=. combine x y)
+                    $items:(z .=. combine y x)
+                    $items:(y .=. z)
                 }
             }
             else
@@ -317,7 +324,7 @@
              * cooperatively reduces this to a single value.
              */
             $items:(sdata "threadIdx.x" .=. y)
-            $items:(reduceBlock dev fun x y sdata (cvar "n"))
+            $items:(reduceBlock dev fun x y z sdata (cvar "n"))
 
             /*
              * Finally, the first thread writes the result for this segment. For
@@ -336,6 +343,7 @@
     (argOut, shOut, setOut)     = writeArray "Out" (undefined :: Array sh e)
     (_, x, declx)               = locals "x" (undefined :: e)
     (_, y, decly)               = locals "y" (undefined :: e)
+    (_, z, declz)               = locals "z" (undefined :: e)
     (sh, _, _)                  = locals "sh" (undefined :: sh :. Int)
     ix                          = [cvar "ix"]
     (smem, sdata)               = shared (undefined :: e) "sdata" [cexp| blockDim.x |] Nothing
@@ -355,7 +363,8 @@
     --
     exclusive_finish (CUExp seed)
       = concat [ x .=. seed
-               , y .=. combine x y ]
+               , z .=. combine y x
+               , y .=. z ]
 
 
 -- Segmented reduction along the innermost dimension of an array. Performs one
@@ -451,6 +460,7 @@
         $decls:smem
         $decls:declx
         $decls:decly
+        $decls:declz
         $items:(sh .=. shIn)
 
         /*
@@ -499,7 +509,8 @@
                     $items:(x .=. get [cvar "ix + warpSize"])
 
                     if ( ix >= start ) {
-                        $items:(y .=. combine x y)
+                        $items:(z .=. combine y x)
+                        $items:(y .=. z)
                     }
                     else {
                         $items:(y .=. x)
@@ -512,7 +523,8 @@
                 for ( ix += 2 * warpSize; ix < end; ix += warpSize )
                 {
                     $items:(x .=. get ix)
-                    $items:(y .=. combine x y)
+                    $items:(z .=. combine y x)
+                    $items:(y .=. z)
                 }
             }
             else if ( start + thread_lane < end )
@@ -525,7 +537,7 @@
              */
             ix = min(num_elements, warpSize);
             $items:(sdata "threadIdx.x" .=. y)
-            $items:(reduceWarp dev fun x y sdata (cvar "ix") (cvar "thread_lane"))
+            $items:(reduceWarp dev fun x y z sdata (cvar "ix") (cvar "thread_lane"))
 
             /*
              * Finally, the first thread writes the result for this segment
@@ -544,6 +556,7 @@
     (argOut, shOut, setOut)     = writeArray "Out" (undefined :: Array (sh :. Int) e)
     (_, x, declx)               = locals "x" (undefined :: e)
     (_, y, decly)               = locals "y" (undefined :: e)
+    (_, z, declz)               = locals "z" (undefined :: e)
     (sh, _, _)                  = locals "sh" (undefined :: sh :. Int)
     (smem, sdata)               = shared (undefined :: e) "sdata" [cexp| blockDim.x |] (Just $ [cexp| &s_ptrs[vectors_per_block][2] |])
     --
@@ -554,7 +567,8 @@
     exclusive_finish (CUExp seed)
       = [[citem| if ( num_elements > 0 ) {
                      $items:(x .=. seed)
-                     $items:(y .=. combine x y)
+                     $items:(z .=. combine y x)
+                     $items:(y .=. z)
                  } else {
                      $items:(y .=. seed)
                  } |]]
@@ -572,28 +586,28 @@
     :: forall aenv e. Elt e
     => DeviceProperties
     -> CUFun2 aenv (e -> e -> e)
-    -> [C.Exp] -> [C.Exp]               -- temporary variables x0 and x1
+    -> [C.Exp] -> [C.Exp] -> [C.Exp]    -- input variables x0 and x1, plus a temporary to store the intermediate result
     -> (Name -> [C.Exp])                -- index elements from shared memory
     -> C.Exp                            -- number of elements
     -> C.Exp                            -- thread identifier: usually lane or thread ID
     -> [C.BlockItem]
-reduceWarp dev fun x0 x1 sdata n tid
+reduceWarp dev fun x0 x1 x2 sdata n tid
   | shflOK dev (undefined :: e) = return
-                                $ reduceWarpShfl dev fun x0 x1       n tid
-  | otherwise                   = reduceWarpTree dev fun x0 x1 sdata n tid
+                                $ reduceWarpShfl dev fun x0 x1          n tid
+  | otherwise                   = reduceWarpTree dev fun x0 x1 x2 sdata n tid
 
 
 reduceBlock
     :: forall aenv e. Elt e
     => DeviceProperties
     -> CUFun2 aenv (e -> e -> e)
-    -> [C.Exp] -> [C.Exp]               -- temporary variables x0 and x1
+    -> [C.Exp] -> [C.Exp] -> [C.Exp]    -- input variables x0 and x1, plus a temporary to store the intermediate result
     -> (Name -> [C.Exp])                -- index elements from shared memory
     -> C.Exp                            -- number of elements
     -> [C.BlockItem]
-reduceBlock dev fun x0 x1 sdata n
-  | shflOK dev (undefined :: e) = reduceBlockShfl dev fun x0 x1 sdata n
-  | otherwise                   = reduceBlockTree dev fun x0 x1 sdata n
+reduceBlock dev fun x0 x1 x2 sdata n
+  | shflOK dev (undefined :: e) = reduceBlockShfl dev fun x0 x1    sdata n
+  | otherwise                   = reduceBlockTree dev fun x0 x1 x2 sdata n
 
 
 -- Tree reduction
@@ -603,12 +617,12 @@
     :: Elt e
     => DeviceProperties
     -> CUFun2 aenv (e -> e -> e)
-    -> [C.Exp] -> [C.Exp]               -- temporary variables x0 and x1
+    -> [C.Exp] -> [C.Exp] -> [C.Exp]    -- input variables x0 and x1, plus a temporary to store the intermediate result
     -> (Name -> [C.Exp])                -- index elements from shared memory
     -> C.Exp                            -- number of elements
     -> C.Exp                            -- thread identifier: usually lane or thread ID
     -> [C.BlockItem]
-reduceWarpTree dev (CUFun2 _ _ f) x0 x1 sdata n tid
+reduceWarpTree dev (CUFun2 _ _ f) x0 x1 x2 sdata n tid
   = map (reduce . pow2) [v, v-1 .. 0]
   where
     v = floor (logBase 2 (fromIntegral $ warpSize dev :: Double))
@@ -620,12 +634,14 @@
     reduce 0
       = [citem| if ( $exp:tid < $exp:n ) {
                     $items:(x0 .=. sdata "threadIdx.x + 1")
-                    $items:(x1 .=. f x1 x0)
+                    $items:(x2 .=. f x1 x0)
+                    $items:(x1 .=. x2)
                 } |]
     reduce i
       = [citem| if ( $exp:tid + $int:i < $exp:n ) {
                     $items:(x0 .=. sdata ("threadIdx.x + " ++ show i))
-                    $items:(x1 .=. f x1 x0)
+                    $items:(x2 .=. f x1 x0)
+                    $items:(x1 .=. x2)
                     $items:(sdata "threadIdx.x" .=. x1)
                 } |]
 
@@ -633,11 +649,11 @@
     :: Elt e
     => DeviceProperties
     -> CUFun2 aenv (e -> e -> e)
-    -> [C.Exp] -> [C.Exp]               -- temporary variables x0 and x1
+    -> [C.Exp] -> [C.Exp] -> [C.Exp]    -- input variables x0 and x1, plus a temporary to store the intermediate result
     -> (Name -> [C.Exp])                -- index elements from shared memory
     -> C.Exp                            -- number of elements
     -> [C.BlockItem]
-reduceBlockTree dev fun@(CUFun2 _ _ f) x0 x1 sdata n
+reduceBlockTree dev fun@(CUFun2 _ _ f) x0 x1 x2 sdata n
   = flip (foldr1 (.)) []
   $ map (reduce . pow2) [u-1, u-2 .. v]
 
@@ -664,7 +680,8 @@
       = [citem| __syncthreads(); |]
       : [citem| if ( threadIdx.x + $int:i < $exp:n ) {
                     $items:(x0 .=. sdata ("threadIdx.x + " ++ show i))
-                    $items:(x1 .=. f x1 x0)
+                    $items:(x2 .=. f x1 x0)
+                    $items:(x1 .=. x2)
                 } |]
       : [citem| __syncthreads(); |]
       : (sdata "threadIdx.x" .=. x1)
@@ -677,7 +694,7 @@
       | otherwise
       = [citem| __syncthreads(); |]
       : [citem| if ( threadIdx.x < $int:(warpSize dev) ) {
-                    $items:(reduceWarpTree dev fun x0 x1 sdata n (cvar "threadIdx.x"))
+                    $items:(reduceWarpTree dev fun x0 x1 x2 sdata n (cvar "threadIdx.x"))
                 } |]
       : rest
 
@@ -718,7 +735,7 @@
       where
         shfl 4  = "shfl_xor32"
         shfl 8  = "shfl_xor64"
-        shfl _  = INTERNAL_ERROR(error) "shfl_xor" "I only know about 32- and 64-bit types"
+        shfl _  = $internalError "shfl_xor" "I only know about 32- and 64-bit types"
 
 
 -- Reduce a block of values in butterfly fashion using __shfl_xor(). Each warp
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
@@ -1,11 +1,12 @@
 {-# LANGUAGE BangPatterns        #-}
 {-# LANGUAGE GADTs               #-}
+{-# LANGUAGE ImpredicativeTypes  #-}
 {-# LANGUAGE QuasiQuotes         #-}
 {-# LANGUAGE ScopedTypeVariables #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.CodeGen.Stencil
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
diff --git a/Data/Array/Accelerate/CUDA/CodeGen/Stencil/Extra.hs b/Data/Array/Accelerate/CUDA/CodeGen/Stencil/Extra.hs
--- a/Data/Array/Accelerate/CUDA/CodeGen/Stencil/Extra.hs
+++ b/Data/Array/Accelerate/CUDA/CodeGen/Stencil/Extra.hs
@@ -1,12 +1,12 @@
-{-# LANGUAGE CPP                 #-}
 {-# LANGUAGE ImpredicativeTypes  #-}
 {-# LANGUAGE QuasiQuotes         #-}
 {-# LANGUAGE ScopedTypeVariables #-}
+{-# LANGUAGE TemplateHaskell     #-}
 {-# LANGUAGE ViewPatterns        #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.CodeGen.Stencil.Extra
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2013] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
@@ -33,6 +33,7 @@
 import qualified Language.C.Syntax                      as C
 
 -- friends
+import Data.Array.Accelerate.Error                      ( internalError )
 import Data.Array.Accelerate.Type                       ( Boundary(..) )
 import Data.Array.Accelerate.Array.Sugar                ( Array, Shape, Elt, shapeToList )
 import Data.Array.Accelerate.Analysis.Shape
@@ -42,9 +43,7 @@
 import Data.Array.Accelerate.CUDA.CodeGen.Base
 import Data.Array.Accelerate.CUDA.CodeGen.Type
 
-#include "accelerate.h"
 
-
 -- Stencil Access
 -- --------------
 
@@ -190,7 +189,7 @@
 -- Test whether an index lies within the boundaries of a shape (first argument)
 --
 cinRange :: [C.Exp] -> [C.Exp] -> C.Exp
-cinRange []    []    = INTERNAL_ERROR(error) "inRange" "singleton index"
+cinRange []    []    = $internalError "inRange" "singleton index"
 cinRange shape index = foldl1 and (zipWith inside shape index)
   where
     inside sz i = [cexp| ({ const $ty:cint _i = $exp:i; _i >= 0 && _i < $exp:sz; }) |]
@@ -264,12 +263,12 @@
 zipWith :: (a -> b -> c) -> [a] -> [b] -> [c]
 zipWith f (x:xs) (y:ys) = f x y : zipWith f xs ys
 zipWith _ []     []     = []
-zipWith _ _      _      = INTERNAL_ERROR(error) "zipWith" "argument mismatch"
+zipWith _ _      _      = $internalError "zipWith" "argument mismatch"
 
 zipWith3 :: (a -> b -> c -> d) -> [a] -> [b] -> [c] -> [d]
 zipWith3 f (x:xs) (y:ys) (z:zs) = f x y z : zipWith3 f xs ys zs
 zipWith3 _ []     []     []     = []
-zipWith3 _ _      _      _      = INTERNAL_ERROR(error) "zipWith3" "argument mismatch"
+zipWith3 _ _      _      _      = $internalError "zipWith3" "argument mismatch"
 
 
 -- Split a list into segments of given length
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
@@ -1,11 +1,11 @@
-{-# LANGUAGE CPP           #-}
-{-# LANGUAGE GADTs         #-}
-{-# LANGUAGE PatternGuards #-}
-{-# LANGUAGE QuasiQuotes   #-}
+{-# LANGUAGE GADTs           #-}
+{-# LANGUAGE PatternGuards   #-}
+{-# LANGUAGE QuasiQuotes     #-}
+{-# LANGUAGE TemplateHaskell #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.CodeGen
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
@@ -27,18 +27,20 @@
 ) where
 
 -- friends
+import Data.Array.Accelerate.Array.Data
+import Data.Array.Accelerate.Error
 import Data.Array.Accelerate.Type
 import Data.Array.Accelerate.Trafo
 import qualified Data.Array.Accelerate.Array.Sugar      as Sugar
 import qualified Data.Array.Accelerate.Analysis.Type    as Sugar
 
 -- libraries
+import Data.Bits
+import Data.Typeable
 import Language.C.Quote.CUDA
 import qualified Language.C                             as C
 
-#include "accelerate.h"
 
-
 typename :: String -> C.Type
 typename name = [cty| typename $id:name |]
 
@@ -55,7 +57,7 @@
 segmentsType :: DelayedOpenAcc aenv (Sugar.Segments i) -> C.Type
 segmentsType seg
   | [s] <- accType seg  = s
-  | otherwise           = INTERNAL_ERROR(error) "accType" "non-scalar segment type"
+  | otherwise           = $internalError "accType" "non-scalar segment type"
 
 
 eltType :: Sugar.Elt a => a {- dummy -} -> [C.Type]
@@ -105,9 +107,8 @@
 codegenIntegralType (TypeCULong  _) = [cty|unsigned long int|]
 codegenIntegralType (TypeCLLong  _) = [cty|long long int|]
 codegenIntegralType (TypeCULLong _) = [cty|unsigned long long int|]
--- XXX: GHC's inbuilt CPP system can't handle stringification
-codegenIntegralType (TypeInt     _) = typename $ "Int"  ++ show (SIZEOF_HTYPE_INT * 8 :: Int)
-codegenIntegralType (TypeWord    _) = typename $ "Word" ++ show (SIZEOF_HTYPE_INT * 8 :: Int)
+codegenIntegralType (TypeInt     _) = typename (showsTypeRep (typeOf (undefined::HTYPE_INT))  "")
+codegenIntegralType (TypeWord    _) = typename (showsTypeRep (typeOf (undefined::HTYPE_WORD)) "")
 
 codegenFloatingType :: FloatingType a -> C.Type
 codegenFloatingType (TypeFloat   _) = [cty|float|]
@@ -162,9 +163,8 @@
 codegenIntegralTex (TypeCULong  _) = typename "TexCULong"
 codegenIntegralTex (TypeCLLong  _) = typename "TexCLLong"
 codegenIntegralTex (TypeCULLong _) = typename "TexCULLong"
--- XXX: GHC's inbuilt CPP system can't handle stringification
-codegenIntegralTex (TypeInt     _) = typename $ "TexInt"  ++ show (SIZEOF_HTYPE_INT * 8 :: Int)
-codegenIntegralTex (TypeWord    _) = typename $ "TexWord" ++ show (SIZEOF_HTYPE_INT * 8 :: Int)
+codegenIntegralTex (TypeInt     _) = typename ("TexInt"  ++ show (finiteBitSize (undefined::Int)))
+codegenIntegralTex (TypeWord    _) = typename ("TexWord" ++ show (finiteBitSize (undefined::Word)))
 
 codegenFloatingTex :: FloatingType a -> C.Type
 codegenFloatingTex (TypeFloat   _) = typename "TexFloat"
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
@@ -3,11 +3,12 @@
 {-# LANGUAGE PatternGuards       #-}
 {-# LANGUAGE RecordWildCards     #-}
 {-# LANGUAGE ScopedTypeVariables #-}
+{-# LANGUAGE TemplateHaskell     #-}
 {-# LANGUAGE TupleSections       #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.Compile
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
@@ -22,9 +23,8 @@
 
 ) where
 
-#include "accelerate.h"
-
 -- friends
+import Data.Array.Accelerate.Error
 import Data.Array.Accelerate.Tuple
 import Data.Array.Accelerate.Trafo
 import Data.Array.Accelerate.CUDA.AST
@@ -50,6 +50,7 @@
 import Control.Concurrent
 import Crypto.Hash.MD5                                          ( hashlazy )
 import Data.List                                                ( intercalate )
+import Data.Bits
 import Data.Maybe
 import Data.Monoid
 import System.Directory
@@ -58,7 +59,6 @@
 import System.IO
 import System.IO.Error
 import System.IO.Unsafe
-import System.Time
 import System.Process
 import System.Mem.Weak
 import Text.PrettyPrint.Mainland                                ( ppr, renderCompact, displayLazyText )
@@ -72,6 +72,9 @@
 
 import GHC.Conc                                                 ( getNumProcessors )
 
+#ifdef ACCELERATE_DEBUG
+import System.Time
+#endif
 #ifdef VERSION_unix
 import System.Posix.Process
 #else
@@ -114,7 +117,7 @@
     -- are merged at every step.
     --
     traverseAcc :: forall aenv arrs. DelayedOpenAcc aenv arrs -> CIO (ExecOpenAcc aenv arrs)
-    traverseAcc Delayed{} = INTERNAL_ERROR(error) "compileOpenAcc" "unexpected delayed array"
+    traverseAcc Delayed{} = $internalError "compileOpenAcc" "unexpected delayed array"
     traverseAcc topAcc@(Manifest pacc) =
       case pacc of
         -- Environment and control flow
@@ -196,10 +199,10 @@
         travF (Lam  f)  = liftA Lam  <$> travF f
 
         noKernel :: FL.FullList () (AccKernel a)
-        noKernel =  FL.FL () (INTERNAL_ERROR(error) "compile" "no kernel module for this node") FL.Nil
+        noKernel =  FL.FL () ($internalError "compile" "no kernel module for this node") FL.Nil
 
         fullOfList :: [a] -> FL.FullList () a
-        fullOfList []       = INTERNAL_ERROR(error) "fullList" "empty list"
+        fullOfList []       = $internalError "fullList" "empty list"
         fullOfList [x]      = FL.singleton () x
         fullOfList (x:xs)   = FL.cons () x (fullOfList xs)
 
@@ -215,7 +218,7 @@
           Nothing       -> liftA2 (Aforeign ff)          <$> pure <$> compileAfun afun <*> travA a
           Just _        -> liftA  (Aforeign ff err)      <$> travA a
             where
-              err = INTERNAL_ERROR(error) "compile" "Executing pure version of a CUDA foreign function"
+              err = $internalError "compile" "Executing pure version of a CUDA foreign function"
 
     -- Traverse a scalar expression
     --
@@ -303,7 +306,7 @@
 
         bind :: (Shape sh, Elt e) => ExecOpenAcc aenv (Array sh e) -> Free aenv
         bind (ExecAcc _ _ (Avar ix)) = freevar ix
-        bind _                       = INTERNAL_ERROR(error) "bind" "expected array variable"
+        bind _                       = $internalError "bind" "expected array variable"
 
 
 -- Applicative
@@ -368,7 +371,7 @@
 --
 link :: Context -> KernelTable -> KernelKey -> IO CUDA.Module
 link context table key =
-  let intErr    = INTERNAL_ERROR(error) "link" "missing kernel entry"
+  let intErr    = $internalError "link" "missing kernel entry"
       ctx       = deviceContext context
       weak_ctx  = weakContext context
   in do
@@ -394,7 +397,7 @@
         -- cache
         --
         KT.insert table key $! KernelObject bin (FL.singleton ctx mdl)
-        KT.persist cubin key
+        KT.persist table cubin key
 
         -- Remove temporary build products.
         -- If compiling kernels with debugging symbols, leave the source files
@@ -467,10 +470,10 @@
   where
     warnings    = D.mode D.dump_cc && D.mode D.verbose
     debug       = D.mode D.debug_cc
-    machine     = case (SIZEOF_HTYPE_INT :: Int) of
-                    4   -> "-m32"
-                    8   -> "-m64"
-                    _   -> INTERNAL_ERROR(error) "compileFlags" "unknown 'Int' size"
+    machine     = case finiteBitSize (undefined :: Int) of
+                    32  -> "-m32"
+                    64  -> "-m64"
+                    _   -> $internalError "compileFlags" "unknown 'Int' size"
 
 
 -- Open a unique file in the temporary directory used for compilation
diff --git a/Data/Array/Accelerate/CUDA/Context.hs b/Data/Array/Accelerate/CUDA/Context.hs
--- a/Data/Array/Accelerate/CUDA/Context.hs
+++ b/Data/Array/Accelerate/CUDA/Context.hs
@@ -3,7 +3,7 @@
 {-# LANGUAGE ViewPatterns  #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.Context
--- Copyright   : [2013] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2013..2014] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
@@ -71,6 +71,7 @@
   return actx
 
 -- |Given a device context, construct a new context around it.
+--
 fromDeviceContext :: CUDA.Device -> CUDA.Context -> IO Context
 fromDeviceContext dev ctx = do
   prp           <- CUDA.props dev
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
@@ -6,8 +6,8 @@
 {-# OPTIONS -fno-warn-unused-imports      #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.Debug
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
@@ -40,6 +40,7 @@
 import System.IO.Unsafe
 import System.Environment
 import System.Console.GetOpt
+import Foreign.CUDA.Driver.Stream                       ( Stream )
 import qualified Foreign.CUDA.Driver.Event              as Event
 
 import GHC.Float
@@ -178,18 +179,19 @@
     :: MonadIO m
     => (Flags :-> Bool)
     -> (Double -> Double -> String)
+    -> Maybe Stream
     -> m ()
     -> m ()
-timed _f _str action
+timed _f _str _stream action
 #ifdef ACCELERATE_DEBUG
   | mode _f
   = do
       gpuBegin  <- liftIO $ Event.create []
       gpuEnd    <- liftIO $ Event.create []
       cpuBegin  <- liftIO getCPUTime
-      liftIO $ Event.record gpuBegin Nothing
+      liftIO $ Event.record gpuBegin _stream
       action
-      liftIO $ Event.record gpuEnd Nothing
+      liftIO $ Event.record gpuEnd _stream
       cpuEnd    <- liftIO getCPUTime
 
       -- Wait for the GPU to finish executing then display the timing execution
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
@@ -6,13 +6,14 @@
 {-# LANGUAGE NoForeignFunctionInterface #-}
 {-# LANGUAGE PatternGuards              #-}
 {-# LANGUAGE ScopedTypeVariables        #-}
+{-# LANGUAGE TemplateHaskell            #-}
 {-# LANGUAGE TypeOperators              #-}
 {-# LANGUAGE TypeSynonymInstances       #-}
 {-# LANGUAGE UndecidableInstances       #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.Execute
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
@@ -42,6 +43,7 @@
 import qualified Data.Array.Accelerate.CUDA.Execute.Event       as Event
 import qualified Data.Array.Accelerate.CUDA.Execute.Stream      as Stream
 
+import Data.Array.Accelerate.Error
 import Data.Array.Accelerate.Tuple
 import Data.Array.Accelerate.Interpreter                        ( evalPrim, evalPrimConst, evalPrj )
 import Data.Array.Accelerate.Array.Data                         ( ArrayElt, ArrayData )
@@ -61,13 +63,11 @@
 import Data.Word
 import Data.Maybe
 
-import Foreign.CUDA.Analysis.Device                             ( DeviceProperties, computeCapability, Compute(..) )
+import Foreign.CUDA.Analysis.Device                             ( computeCapability, Compute(..) )
 import qualified Foreign.CUDA.Driver                            as CUDA
 import qualified Data.HashMap.Strict                            as Map
 
-#include "accelerate.h"
 
-
 -- Asynchronous kernel execution
 -- -----------------------------
 
@@ -88,7 +88,7 @@
 aprj :: Idx env t -> Aval env -> Async t
 aprj ZeroIdx       (Apush _   x) = x
 aprj (SuccIdx idx) (Apush val _) = aprj idx val
-aprj _             _             = INTERNAL_ERROR(error) "aprj" "inconsistent valuation"
+aprj _             _             = $internalError "aprj" "inconsistent valuation"
 
 
 -- All work submitted to the given stream will occur after the asynchronous
@@ -108,11 +108,11 @@
 
 -- Execute the given computation in a unique execution stream.
 --
-streaming :: (Stream -> CIO a) -> CIO (Async a)
-streaming action = do
+streaming :: (Stream -> CIO a) -> (Async a -> CIO b) -> CIO b
+streaming first second = do
   context   <- asks activeContext
   reservoir <- gets streamReservoir
-  uncurry Async <$> Stream.streaming context reservoir action
+  Stream.streaming context reservoir first (\e a -> second (Async e a))
 
 
 -- Array expression evaluation
@@ -132,20 +132,21 @@
 --    skeleton are invoked
 --
 executeAcc :: Arrays a => ExecAcc a -> CIO a
-executeAcc !acc = wait =<< streaming (executeOpenAcc acc Aempty)
+executeAcc !acc = streaming (executeOpenAcc acc Aempty) wait
 
 executeAfun1 :: (Arrays a, Arrays b) => ExecAfun (a -> b) -> a -> CIO b
 executeAfun1 !afun !arrs = do
-  Async event () <- streaming $ \st -> useArrays (arrays arrs) (fromArr arrs) st
-  executeOpenAfun1 afun Aempty (Async event arrs)
+  streaming (useArrays (arrays arrs) (fromArr arrs))
+            (\(Async event ()) -> executeOpenAfun1 afun Aempty (Async event arrs))
   where
     useArrays :: ArraysR arrs -> arrs -> Stream -> CIO ()
     useArrays ArraysRunit         ()       _  = return ()
     useArrays (ArraysRpair r1 r0) (a1, a0) st = useArrays r1 a1 st >> useArrays r0 a0 st
     useArrays ArraysRarray        arr      st = useArrayAsync arr (Just st)
 
+
 executeOpenAfun1 :: PreOpenAfun ExecOpenAcc aenv (a -> b) -> Aval aenv -> Async a -> CIO b
-executeOpenAfun1 (Alam (Abody f)) aenv x = wait =<< streaming (executeOpenAcc f (aenv `Apush` x))
+executeOpenAfun1 (Alam (Abody f)) aenv x = streaming (executeOpenAcc f (aenv `Apush` x)) wait
 executeOpenAfun1 _                _    _ = error "the sword comes out after you swallow it, right?"
 
 
@@ -158,7 +159,7 @@
     -> Stream
     -> CIO arrs
 executeOpenAcc EmbedAcc{} _ _
-  = INTERNAL_ERROR(error) "execute" "unexpected delayed array"
+  = $internalError "execute" "unexpected delayed array"
 executeOpenAcc (ExecAcc (FL () kernel more) !gamma !pacc) !aenv !stream
   = case pacc of
 
@@ -168,10 +169,10 @@
 
       -- Environment manipulation
       Avar ix                   -> after stream (aprj ix aenv)
-      Alet bnd body             -> streamA bnd >>= \x -> executeOpenAcc body (aenv `Apush` x) stream
+      Alet bnd body             -> streaming (executeOpenAcc bnd aenv) (\x -> executeOpenAcc body (aenv `Apush` x) stream)
+      Apply f a                 -> streaming (executeOpenAcc a aenv)   (executeOpenAfun1 f aenv)
       Atuple tup                -> toTuple <$> travT tup
       Aprj ix tup               -> evalPrj ix . fromTuple <$> travA tup
-      Apply f a                 -> executeOpenAfun1 f aenv =<< streamA a
       Acond p t e               -> travE p >>= \x -> if x then travA t else travA e
       Awhile p f a              -> awhile p f =<< travA a
 
@@ -206,15 +207,12 @@
       ZipWith _ _ _             -> fusionError
 
   where
-    fusionError = INTERNAL_ERROR(error) "executeOpenAcc" "unexpected fusible matter"
+    fusionError = $internalError "executeOpenAcc" "unexpected fusible matter"
 
     -- term traversals
     travA :: ExecOpenAcc aenv a -> CIO a
     travA !acc = executeOpenAcc acc aenv stream
 
-    streamA :: ExecOpenAcc aenv a -> CIO (Async a)
-    streamA !acc = streaming $ \st -> executeOpenAcc acc aenv st
-
     travE :: ExecExp aenv t -> CIO t
     travE !exp = executeExp exp aenv stream
 
@@ -232,7 +230,7 @@
 
     -- get the extent of an embedded array
     extent :: Shape sh => ExecOpenAcc aenv (Array sh e) -> CIO sh
-    extent ExecAcc{}     = INTERNAL_ERROR(error) "executeOpenAcc" "expected delayed array"
+    extent ExecAcc{}     = $internalError "executeOpenAcc" "expected delayed array"
     extent (EmbedAcc sh) = travE sh
 
     -- Skeleton implementation
@@ -252,7 +250,7 @@
     --
     reshapeOp :: Shape sh => sh -> Array sh' e -> Array sh e
     reshapeOp sh (Array sh' adata)
-      = BOUNDS_CHECK(check) "reshape" "shape mismatch" (size sh == R.size sh')
+      = $boundsCheck "reshape" "shape mismatch" (size sh == R.size sh')
       $ Array (fromElt sh) adata
 
     -- Executing fold operations depend on whether we are recursively collapsing
@@ -261,7 +259,7 @@
     --
     fold1Op :: (Shape sh, Elt e) => (sh :. Int) -> CIO (Array sh e)
     fold1Op !sh@(_ :. sz)
-      = BOUNDS_CHECK(check) "fold1" "empty array" (sz > 0)
+      = $boundsCheck "fold1" "empty array" (sz > 0)
       $ foldCore sh
 
     foldOp :: (Shape sh, Elt e) => (sh :. Int) -> CIO (Array sh e)
@@ -295,7 +293,7 @@
                 foldRec out
 
       | otherwise
-      = INTERNAL_ERROR(error) "foldRec" "missing phase-2 kernel module"
+      = $internalError "foldRec" "missing phase-2 kernel module"
 
     -- Segmented reduction. Subtract one from the size of the segments vector as
     -- this is the result of an exclusive scan to calculate segment offsets.
@@ -363,15 +361,22 @@
           execute kernel gamma aenv numElements (numElements, d_body, blk, d_sum) stream
 
       | otherwise
-      = INTERNAL_ERROR(error) "scanOp" "missing multi-block kernel module(s)"
+      = $internalError "scanOp" "missing multi-block kernel module(s)"
 
     -- Forward permutation
     --
-    permuteOp :: (Shape sh, Shape sh', Elt e) => sh -> Array sh' e -> CIO (Array sh' e)
+    permuteOp :: forall sh sh' e. (Shape sh, Shape sh', Elt e) => sh -> Array sh' e -> CIO (Array sh' e)
     permuteOp !sh !dfs = do
-      out <- allocateArray (shape dfs)
-      copyArray dfs out
-      execute kernel gamma aenv (size sh) out stream
+      let sh'   = shape dfs
+          n'    = size sh'
+
+      out               <- allocateArray sh'
+      Array _ locks     <- allocateArray sh'            :: CIO (Array sh' Int32)
+      ((), d_locks)     <- devicePtrsOfArrayData locks  :: CIO ((), CUDA.DevicePtr Int32)
+
+      liftIO $ CUDA.memsetAsync d_locks n' 0 (Just stream)      -- TLM: overlap these two operations?
+      copyArrayAsync dfs out (Just stream)
+      execute kernel gamma aenv (size sh) (out, d_locks) stream
       return out
 
     -- Stencil operations. NOTE: the arguments to 'namesOfArray' must be the
@@ -412,7 +417,7 @@
           return out
 
       | otherwise
-      = INTERNAL_ERROR(error) "stencil2Op" "missing stencil specialisation kernel"
+      = $internalError "stencil2Op" "missing stencil specialisation kernel"
 
 
 -- Scalar expression evaluation
@@ -659,7 +664,7 @@
 --
 launch :: AccKernel a -> (Int,Int,Int) -> [CUDA.FunParam] -> Stream -> CIO ()
 launch (AccKernel entry !fn _ _ _ _ _) !(cta, grid, smem) !args !stream
-  = D.timed D.dump_exec msg
+  = D.timed D.dump_exec msg (Just stream)
   $ liftIO $ CUDA.launchKernel fn (grid,1,1) (cta,1,1) smem (Just stream) args
   where
     msg gpuTime cpuTime
diff --git a/Data/Array/Accelerate/CUDA/Execute/Event.hs b/Data/Array/Accelerate/CUDA/Execute/Event.hs
--- a/Data/Array/Accelerate/CUDA/Execute/Event.hs
+++ b/Data/Array/Accelerate/CUDA/Execute/Event.hs
@@ -1,10 +1,8 @@
-{-# LANGUAGE BangPatterns  #-}
-{-# LANGUAGE MagicHash     #-}
-{-# LANGUAGE UnboxedTuples #-}
+{-# LANGUAGE BangPatterns #-}
+{-# LANGUAGE CPP          #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.Execute.Event
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2013] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2013..2014] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
@@ -13,30 +11,29 @@
 --
 module Data.Array.Accelerate.CUDA.Execute.Event (
 
-  Event, create, waypoint, after, block,
+  Event, create, waypoint, after, block, Event.destroy,
 
 ) where
 
 -- friends
+#ifdef ACCELERATE_DEBUG
 import qualified Data.Array.Accelerate.CUDA.Debug               as D
+#endif
 
 -- libraries
 import Foreign.CUDA.Driver.Event                                ( Event(..) )
-import Foreign.CUDA.Driver.Stream                               ( Stream )
+import Foreign.CUDA.Driver.Stream                               ( Stream(..) )
 import qualified Foreign.CUDA.Driver.Event                      as Event
 
-import GHC.Base
-import GHC.Ptr
 
-
--- Create a new event that will be automatically garbage collected. The event is
+-- Create a new event. It will not be automatically garbage collected, and is
 -- not suitable for timing purposes.
 --
 {-# INLINE create #-}
 create :: IO Event
 create = do
   event <- Event.create [Event.DisableTiming]
-  addEventFinalizer event $ trace ("destroy " ++ show event) (Event.destroy event)
+  message ("create " ++ showEvent event)
   return event
 
 -- Create a new event marker that will be filled once execution in the specified
@@ -47,7 +44,7 @@
 waypoint stream = do
   event <- create
   Event.record event (Just stream)
---  message $ "waypoint " ++ show event ++ " in " ++ show stream
+  message $ "waypoint " ++ showEvent event ++ " in " ++ showStream stream
   return event
 
 -- Make all future work submitted to the given stream wait until the event
@@ -56,7 +53,7 @@
 {-# INLINE after #-}
 after :: Event -> Stream -> IO ()
 after event stream = do
---  message $ "after " ++ show event ++ " in " ++ show stream
+  message $ "after " ++ showEvent event ++ " in " ++ showStream stream
   Event.wait event (Just stream) []
 
 -- Block the calling thread until the event is recorded
@@ -68,9 +65,9 @@
 
 -- Add a finaliser to an event token
 --
-addEventFinalizer :: Event -> IO () -> IO ()
-addEventFinalizer e@(Event (Ptr e#)) f = IO $ \s ->
-  case mkWeak# e# e f s of (# s', _w #) -> (# s', () #)
+-- addEventFinalizer :: Event -> IO () -> IO ()
+-- addEventFinalizer e@(Event (Ptr e#)) f = IO $ \s ->
+--   case mkWeak# e# e f s of (# s', _w #) -> (# s', () #)
 
 
 -- Debug
@@ -78,9 +75,21 @@
 
 {-# INLINE trace #-}
 trace :: String -> IO a -> IO a
-trace msg next = D.message D.dump_exec ("event: " ++ msg) >> next
+trace _msg next = do
+#ifdef ACCELERATE_DEBUG
+  D.when D.verbose $ D.message D.dump_exec ("event: " ++ _msg)
+#endif
+  next
 
--- {-# INLINE message #-}
--- message :: String -> IO ()
--- message s = s `trace` return ()
+{-# INLINE message #-}
+message :: String -> IO ()
+message s = s `trace` return ()
+
+{-# INLINE showEvent #-}
+showEvent :: Event -> String
+showEvent (Event e) = show e
+
+{-# INLINE showStream #-}
+showStream :: Stream -> String
+showStream (Stream s) = show s
 
diff --git a/Data/Array/Accelerate/CUDA/Execute/Stream.hs b/Data/Array/Accelerate/CUDA/Execute/Stream.hs
--- a/Data/Array/Accelerate/CUDA/Execute/Stream.hs
+++ b/Data/Array/Accelerate/CUDA/Execute/Stream.hs
@@ -1,7 +1,8 @@
 {-# LANGUAGE BangPatterns #-}
+{-# LANGUAGE CPP          #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.Execute.Stream
--- Copyright   : [2013] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2013..2014] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
@@ -22,16 +23,17 @@
 import Data.Array.Accelerate.CUDA.Execute.Event                 ( Event )
 import qualified Data.Array.Accelerate.CUDA.Execute.Event       as Event
 import qualified Data.Array.Accelerate.CUDA.FullList            as FL
+
+#ifdef ACCELERATE_DEBUG
 import qualified Data.Array.Accelerate.CUDA.Debug               as D
+#endif
 
 -- libraries
-import Control.Monad                                            ( void )
 import Control.Monad.Trans                                      ( MonadIO, liftIO )
 import Control.Exception                                        ( bracket_ )
-import Control.Concurrent                                       ( forkIO )
 import Control.Concurrent.MVar                                  ( MVar, newMVar, withMVar, mkWeakMVar )
 import System.Mem.Weak                                          ( Weak, deRefWeak )
-import Foreign.CUDA.Driver.Stream                               ( Stream )
+import Foreign.CUDA.Driver.Stream                               ( Stream(..) )
 import qualified Foreign.CUDA.Driver                            as CUDA
 import qualified Foreign.CUDA.Driver.Stream                     as Stream
 
@@ -55,16 +57,21 @@
 -- Executing operations in streams
 -- -------------------------------
 
--- Execute an operation in a unique execution stream.
+-- Execute an operation in a unique execution stream. The (asynchronous) result
+-- is passed to a second operation together with an event that will be signalled
+-- once the operation is complete. The stream and event are released after the
+-- second operation completes.
 --
 {-# INLINE streaming #-}
-streaming :: MonadIO m => Context -> Reservoir -> (Stream -> m a) -> m (Event, a)
-streaming !ctx !rsv@(Reservoir !_ !weak_rsv) !action = do
+streaming :: MonadIO m => Context -> Reservoir -> (Stream -> m a) -> (Event -> a -> m b) -> m b
+streaming !ctx !rsv@(Reservoir !_ !weak_rsv) !action !after = do
   stream <- liftIO $ create ctx rsv
-  result <- action stream
+  first  <- action stream
   end    <- liftIO $ Event.waypoint stream
-  liftIO $ merge (weakContext ctx) weak_rsv stream
-  return $ (end, result)
+  final  <- after end first
+  liftIO $! destroy (weakContext ctx) weak_rsv stream
+  liftIO $! Event.destroy end
+  return final
 
 
 -- Primitive operations
@@ -86,54 +93,60 @@
 --
 {-# INLINE create #-}
 create :: Context -> Reservoir -> IO Stream
-create !ctx (Reservoir !ref _) = withMVar ref $ \tbl -> do
+create !ctx (Reservoir !ref !_) = withMVar ref $ \tbl -> do
   let key = deviceContext ctx
   --
   ms    <- HT.lookup tbl key
   case ms of
     Nothing -> do
-        stream <- Stream.create []
-        message ("new " ++ showStream stream)
-        return stream
+      stream <- Stream.create []
+      message ("new " ++ showStream stream)
+      return stream
 
     Just (FL () stream rest) -> do
       case rest of
         FL.Nil          -> HT.delete tbl key
         FL.Cons () s ss -> HT.insert tbl key (FL () s ss)
-      --
+        --
       return stream
 
 
--- Merge a stream back into the reservoir. This is done asynchronously, once all
+-- Merge a stream back into the reservoir. This must only be done once all
 -- pending operations in the stream have completed.
 --
-{-# INLINE merge #-}
-merge :: Weak CUDA.Context -> Weak RSV -> Stream -> IO ()
-merge !weak_ctx !weak_rsv !stream
-  = void . forkIO
-  $ do  -- wait for all operations to complete
-        -- Stream.block stream
+{-# INLINE destroy #-}
+destroy :: Weak CUDA.Context -> Weak RSV -> Stream -> IO ()
+destroy !weak_ctx !weak_rsv !stream = do
+  -- Wait for all preceding operations submitted to the stream to complete. Not
+  -- necessary because of the setup of 'streaming'.
+  -- Stream.block stream
 
-        -- Now check whether the context and reservoir are still active. Return
-        -- the stream back to the reservoir for later reuse if we can, otherwise
-        -- destroy it.
-        --
-        mc <- deRefWeak weak_ctx
-        case mc of
-          Nothing       -> message ("finalise/dead context " ++ showStream stream)
-          Just ctx      -> do
-            --
-            mr <- deRefWeak weak_rsv
-            case mr of
-              Nothing   -> trace ("merge/free " ++ showStream stream) $ Stream.destroy stream
-              Just ref  -> trace ("merge/save " ++ showStream stream) $ withMVar ref $ \tbl -> do
-                --
-                ms <- HT.lookup tbl ctx
-                case ms of
-                  Nothing       -> HT.insert tbl ctx (FL.singleton () stream)
-                  Just ss       -> HT.insert tbl ctx (FL.cons () stream ss)
+  -- Now check whether the context and reservoir are still active. Return
+  -- the stream back to the reservoir for later reuse if we can, otherwise
+  -- destroy it.
+  mc <- deRefWeak weak_ctx
+  case mc of
+    Nothing       -> message ("finalise/dead context " ++ showStream stream)
+    Just ctx      -> do
+      --
+      mr <- deRefWeak weak_rsv
+      case mr of
+        Nothing   -> trace ("destroy/free " ++ showStream stream) $ Stream.destroy stream
+        Just ref  -> trace ("destroy/save " ++ showStream stream) $ withMVar ref $ \tbl -> do
+          --
+          ms <- HT.lookup tbl ctx
+          case ms of
+            Nothing       -> HT.insert tbl ctx (FL.singleton () stream)
+            Just ss       -> HT.insert tbl ctx (FL.cons () stream ss)
 
 
+-- Add a finaliser to an execution stream
+--
+-- addStreamFinalizer :: Stream -> IO () -> IO ()
+-- addStreamFinalizer st@(Stream (Ptr st#)) f = IO $ \s ->
+--   case mkWeak# st# st f s of (# s', _w #) -> (# s', () #)
+
+
 -- Destroy all streams in the reservoir.
 --
 flush :: HashTable CUDA.Context (FullList () Stream) -> IO ()
@@ -150,7 +163,11 @@
 
 {-# INLINE trace #-}
 trace :: String -> IO a -> IO a
-trace msg next = D.message D.dump_exec ("stream: " ++ msg) >> next
+trace _msg next = do
+#ifdef ACCELERATE_DEBUG
+  D.when D.verbose $ D.message D.dump_exec ("stream: " ++ _msg)
+#endif
+  next
 
 {-# INLINE message #-}
 message :: String -> IO ()
@@ -158,5 +175,5 @@
 
 {-# INLINE showStream #-}
 showStream :: Stream -> String
-showStream (Stream.Stream s) = show s
+showStream (Stream s) = show s
 
diff --git a/Data/Array/Accelerate/CUDA/Foreign.hs b/Data/Array/Accelerate/CUDA/Foreign.hs
--- a/Data/Array/Accelerate/CUDA/Foreign.hs
+++ b/Data/Array/Accelerate/CUDA/Foreign.hs
@@ -1,6 +1,6 @@
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.Foreign
--- Copyright   : [2013] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell, Robert Clifton-Everest
+-- Copyright   : [2013..2014] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell, Robert Clifton-Everest
 -- License     : BSD3
 --
 -- Maintainer  : Robert Clifton-Everest <robertce@cse.unsw.edu.au>
diff --git a/Data/Array/Accelerate/CUDA/Foreign/Export.hs b/Data/Array/Accelerate/CUDA/Foreign/Export.hs
--- a/Data/Array/Accelerate/CUDA/Foreign/Export.hs
+++ b/Data/Array/Accelerate/CUDA/Foreign/Export.hs
@@ -12,7 +12,7 @@
 {-# OPTIONS_GHC -fno-warn-name-shadowing -fno-warn-orphans #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.Foreign.Export
--- Copyright   : [2013] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell, Robert Clifton-Everest
+-- Copyright   : [2013..2014] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell, Robert Clifton-Everest
 -- License     : BSD3
 --
 -- Maintainer  : Robert Clifton-Everest <robertce@cse.unsw.edu.au>
@@ -232,7 +232,7 @@
   genCompileFun n ename ty
 
 genCompileFun :: Name -> String -> Type -> Q [Dec]
-genCompileFun fname ename ty@(AppT (AppT ArrowT (AppT _ a)) (AppT _ b))
+genCompileFun fname ename (AppT (AppT ArrowT (AppT _ _)) (AppT _ _))
   = sequence [sig, dec, expt]
   where
     initName = mkName ename
diff --git a/Data/Array/Accelerate/CUDA/Foreign/Import.hs b/Data/Array/Accelerate/CUDA/Foreign/Import.hs
--- a/Data/Array/Accelerate/CUDA/Foreign/Import.hs
+++ b/Data/Array/Accelerate/CUDA/Foreign/Import.hs
@@ -12,7 +12,7 @@
 {-# OPTIONS_GHC -fno-warn-orphans #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.Foreign.Import
--- Copyright   : [2013] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell, Robert Clifton-Everest
+-- Copyright   : [2013..2014] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell, Robert Clifton-Everest
 -- License     : BSD3
 --
 -- Maintainer  : Robert Clifton-Everest <robertce@cse.unsw.edu.au>
@@ -45,10 +45,11 @@
   -- * Manipulating arrays
   DevicePtrs,
   devicePtrsOfArray,
-  indexArray, copyArray,
+  indexArray,
   useArray,  useArrayAsync,
   peekArray, peekArrayAsync,
   pokeArray, pokeArrayAsync,
+  copyArray, copyArrayAsync,
   allocateArray, newArray,
 
   -- * Running IO actions in an Accelerate context
@@ -78,7 +79,7 @@
                  -> (as -> CIO bs)              -- operation to execute
                  -> CUDAForeignAcc as bs
 
-deriving instance Typeable2 CUDAForeignAcc
+deriving instance Typeable CUDAForeignAcc
 
 instance Foreign CUDAForeignAcc where
   strForeign (CUDAForeignAcc n _) = n
@@ -106,7 +107,7 @@
                  -> String                      -- name of the foreign function
                  -> CUDAForeignExp x y
 
-deriving instance Typeable2 CUDAForeignExp
+deriving instance Typeable CUDAForeignExp
 
 instance Foreign CUDAForeignExp where
   strForeign (CUDAForeignExp _ n) = n
diff --git a/Data/Array/Accelerate/CUDA/FullList.hs b/Data/Array/Accelerate/CUDA/FullList.hs
--- a/Data/Array/Accelerate/CUDA/FullList.hs
+++ b/Data/Array/Accelerate/CUDA/FullList.hs
@@ -2,8 +2,8 @@
 {-# LANGUAGE PatternGuards #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.FullList
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
diff --git a/Data/Array/Accelerate/CUDA/Persistent.hs b/Data/Array/Accelerate/CUDA/Persistent.hs
--- a/Data/Array/Accelerate/CUDA/Persistent.hs
+++ b/Data/Array/Accelerate/CUDA/Persistent.hs
@@ -1,10 +1,11 @@
 {-# LANGUAGE BangPatterns        #-}
 {-# LANGUAGE ScopedTypeVariables #-}
+{-# LANGUAGE TemplateHaskell     #-}
 {-# OPTIONS_GHC -fno-warn-orphans #-}
 -- |
 -- 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
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
@@ -22,6 +23,7 @@
 ) where
 
 -- friends
+import Data.Array.Accelerate.Error
 import Data.Array.Accelerate.CUDA.Context
 import Data.Array.Accelerate.CUDA.FullList              ( FullList )
 import qualified Data.Array.Accelerate.CUDA.Debug       as D
@@ -46,8 +48,9 @@
 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.ByteString                        as BS
+import qualified Data.ByteString.Lazy                   as BL
+import qualified Data.ByteString.Lazy.Internal          as BL
 import qualified Data.HashTable.IO                      as HT
 
 import qualified Foreign.CUDA.Driver                    as CUDA
@@ -67,6 +70,8 @@
 -- Interface -------------------------------------------------------------------
 -- ---------                                                                  --
 
+type HashTable key val = HT.BasicHashTable key val
+
 data KernelTable = KT {-# UNPACK #-} !ProgramCache      -- first level, in-memory cache
                       {-# UNPACK #-} !PersistentCache   -- second level, on-disk cache
 
@@ -76,7 +81,7 @@
   cacheDir <- cacheDirectory
   createDirectoryIfMissing True cacheDir
   --
-  local         <- HT.new
+  local         <- newMVar =<< HT.new
   persistent    <- restore (cacheDir </> "persistent.db")
   --
   return        $! KT local persistent
@@ -86,7 +91,7 @@
 -- the persistent cache, it is loaded and linked into the current context.
 --
 lookup :: Context -> KernelTable -> KernelKey -> IO (Maybe KernelEntry)
-lookup context (KT kt pt) !key = do
+lookup context (KT !kt_ref !pt_ref) !key = withMVar kt_ref $ \kt -> 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
@@ -95,7 +100,7 @@
   v1    <- HT.lookup kt key
   case v1 of
     Just _      -> return v1
-    Nothing     -> do
+    Nothing     -> withMVar pt_ref $ \pt -> 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.
@@ -109,7 +114,7 @@
       Just ()   -> do
         message "found/persistent"
         cubin   <- (</>) <$> cacheDirectory <*> pure (cacheFilePath key)
-        bin     <- B.readFile cubin
+        bin     <- BS.readFile cubin
         !mdl    <- CUDA.loadData bin
         let obj  = KernelObject bin (FL.singleton (deviceContext context) mdl)
         addFinalizer mdl (module_finalizer (weakContext context) key mdl)
@@ -124,8 +129,9 @@
 --      exists there already? Would require updating that hash table as new
 --      entries are added, which the functions currently do not do.
 --
+{-# INLINE insert #-}
 insert :: KernelTable -> KernelKey -> KernelEntry -> IO ()
-insert (KT kt _) !key !val = HT.insert kt key val
+insert (KT !kt_ref !_) !key !val = withMVar kt_ref $ \kt -> HT.insert kt key val
 
 
 -- Unload a kernel module from the specified context
@@ -160,7 +166,7 @@
 -- 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 ProgramCache = MVar ( HashTable KernelKey KernelEntry )
 
 type KernelKey    = (CUDA.Compute, ByteString)
 data KernelEntry
@@ -185,12 +191,16 @@
 -- 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
+-- 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.
+--
+--      One way to handle this is to put a maximum size on the cache (either as
+--      disk space consumed or number of kernels) and once the maximum size is
+--      reached keep only the most recently used files.
 
-type PersistentCache = HT.BasicHashTable KernelKey ()
+type PersistentCache = MVar ( HashTable KernelKey () )
 
 
 -- The root directory of where the various persistent cache files live; the
@@ -210,7 +220,7 @@
 --
 cacheFilePath :: KernelKey -> FilePath
 cacheFilePath (cap, key) =
-  show cap </> zEncodeString (B.foldl (flip (showLitChar . w2c)) [] key)
+  show cap </> zEncodeString (BS.foldl (flip (showLitChar . w2c)) [] key)
 
 -- stolen from compiler/utils/Encoding.hs
 --
@@ -288,7 +298,7 @@
 -- file is created, so that 'persist' can always append elements.
 --
 restore :: FilePath -> IO PersistentCache
-restore db = do
+restore !db = do
   D.when D.flush_cache $ do
     message "deleting persistent cache"
     cacheDir <- cacheDirectory
@@ -296,19 +306,36 @@
     createDirectoryIfMissing True cacheDir
   --
   exists <- doesFileExist db
-  case exists of
+  pt     <- case exists of
     False       -> encodeFile db (0::Int) >> HT.new
     True        -> do
-      store         <- L.readFile db
-      let (n,rest,_) = runGetState get store 0
+      store         <- BL.readFile db
+
+      -- Just read the start of the input to extract the number of entries
+      -- in the persistent kernel table.
+      let (n, rest) = setup (runGetIncremental get) store
+
+          setup (Done s _ r)   lbs = (r, BL.Chunk s lbs)
+          setup (Partial k)    lbs = setup (k (takeHeadChunk lbs)) (dropHeadChunk lbs)
+          setup (Fail _ p msg) _   = $internalError "restore" $ show p ++ ": " ++ msg
+
+          takeHeadChunk (BL.Chunk h _) = Just h
+          takeHeadChunk _              = Nothing
+
+          dropHeadChunk (BL.Chunk _ t) = t
+          dropHeadChunk _              = BL.empty
+
+      -- Allocate the persistent hash table and populate it with entries decoded
+      -- from the index file
       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)
       evaluate pt
+  --
+  newMVar pt
 
 
 -- Append a single value to the persistent cache.
@@ -316,8 +343,8 @@
 -- 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
+persist :: KernelTable -> FilePath -> KernelKey -> IO ()
+persist (KT !_ !pt_ref) !cubin !key = withMVar pt_ref $ \_ -> do
   cacheDir <- cacheDirectory
   let db        = cacheDir </> "persistent.db"
       cacheFile = cacheDir </> cacheFilePath key
@@ -335,14 +362,14 @@
   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
+    n <- runGet (get :: Get Int) `fmap` BL.hGet h 8
     hSeek h AbsoluteSeek 0
-    L.hPut h (encode (n+1))
+    BL.hPut h (encode (n+1))
 
     -- Append the new entry to the end of file
     --
     hSeek h SeekFromEnd 0
-    L.hPut h (encode key)
+    BL.hPut h (encode key)
 
 
 -- Debug
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
@@ -1,11 +1,11 @@
 {-# LANGUAGE BangPatterns               #-}
-{-# LANGUAGE CPP                        #-}
 {-# LANGUAGE GeneralizedNewtypeDeriving #-}
 {-# LANGUAGE ScopedTypeVariables        #-}
+{-# LANGUAGE TemplateHaskell            #-}
 -- |
 -- Module      : Data.Array.Accelerate.CUDA.State
--- Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
---               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+-- Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+--               [2009..2014] Trevor L. McDonell
 -- License     : BSD3
 --
 -- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
@@ -28,6 +28,7 @@
 ) where
 
 -- friends
+import Data.Array.Accelerate.Error
 import Data.Array.Accelerate.CUDA.Context
 import Data.Array.Accelerate.CUDA.Debug                 ( message, dump_gc )
 import Data.Array.Accelerate.CUDA.Persistent            as KT ( KernelTable, new )
@@ -47,9 +48,7 @@
 import Foreign.CUDA.Driver.Error
 import qualified Foreign.CUDA.Driver                    as CUDA
 
-#include "accelerate.h"
 
-
 -- Execution State
 -- ---------------
 
@@ -83,7 +82,7 @@
 evalCUDA !ctx !acc =
   runInBoundThread (bracket_ setup teardown action)
   `catch`
-  \e -> INTERNAL_ERROR(error) "unhandled" (show (e :: CUDAException))
+  \e -> $internalError "unhandled" (show (e :: CUDAException))
   where
     setup       = push ctx
     teardown    = pop >> performGC
diff --git a/Data/Array/Accelerate/Internal/Check.hs b/Data/Array/Accelerate/Internal/Check.hs
deleted file mode 100644
--- a/Data/Array/Accelerate/Internal/Check.hs
+++ /dev/null
@@ -1,123 +0,0 @@
-{-# LANGUAGE CPP #-}
--- |
--- Module      : Data.Array.Accelerate.Internal.Check
--- Copyright   : Roman Lechinskiy, Trevor L. McDonell
--- License     : BSD3
---
--- Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
--- Stability   : experimental
--- Portability : non-portable (GHC extensions)
---
--- Bounds checking infrastructure
---
--- Stolen from the Vector package by Roman Leshchinskiy. This code has a
--- BSD-style license. <http://hackage.haskell.org/package/vector>
---
-
-module Data.Array.Accelerate.Internal.Check (
-
-  -- * Bounds checking and assertion infrastructure
-  Checks(..), doChecks,
-  error, check, warning, assert, checkIndex, checkLength, checkSlice
-
-) where
-
-import Prelude                          hiding ( error )
-import Debug.Trace
-import qualified Prelude                as P
-
-data Checks = Bounds | Unsafe | Internal deriving( Eq )
-
-doBoundsChecks :: Bool
-#ifdef ACCELERATE_BOUNDS_CHECKS
-doBoundsChecks = True
-#else
-doBoundsChecks = False
-#endif
-
-doUnsafeChecks :: Bool
-#ifdef ACCELERATE_UNSAFE_CHECKS
-doUnsafeChecks = True
-#else
-doUnsafeChecks = False
-#endif
-
-doInternalChecks :: Bool
-#ifdef ACCELERATE_INTERNAL_CHECKS
-doInternalChecks = True
-#else
-doInternalChecks = False
-#endif
-
-
-doChecks :: Checks -> Bool
-{-# INLINE doChecks #-}
-doChecks Bounds   = doBoundsChecks
-doChecks Unsafe   = doUnsafeChecks
-doChecks Internal = doInternalChecks
-
-message :: String -> Int -> Checks -> String -> String -> String
-{-# INLINE message #-}
-message file line kind loc msg
-  = unlines
-  $ (if kind == Internal
-       then ([""
-             ,"*** Internal error in package accelerate ***"
-             ,"*** Please submit a bug report at https://github.com/AccelerateHS/accelerate/issues"]++)
-       else id)
-    [ file ++ ":" ++ show line ++ " (" ++ loc ++ "): " ++ msg ]
-
-error :: String -> Int -> Checks -> String -> String -> a
-{-# INLINE error #-}
-error file line kind loc msg
-  = P.error (message file line kind loc msg)
-
-check :: String -> Int -> Checks -> String -> String -> Bool -> a -> a
-{-# INLINE check #-}
-check file line kind loc msg cond x
-  | not (doChecks kind) || cond = x
-  | otherwise = error file line kind loc msg
-
-warning :: String -> Int -> Checks -> String -> String -> Bool -> a -> a
-{-# INLINE warning #-}
-warning file line kind loc msg cond x
-  | not (doChecks kind) || cond = x
-  | otherwise                   = trace (message file line kind loc msg) x
-
-assert_msg :: String
-assert_msg = "assertion failure"
-
-assert :: String -> Int -> Checks -> String -> Bool -> a -> a
-{-# INLINE assert #-}
-assert file line kind loc = check file line kind loc assert_msg
-
-checkIndex_msg :: Int -> Int -> String
-{-# NOINLINE checkIndex_msg #-}
-checkIndex_msg i n = "index out of bounds " ++ show (i,n)
-
-checkIndex :: String -> Int -> Checks -> String -> Int -> Int -> a -> a
-{-# INLINE checkIndex #-}
-checkIndex file line kind loc i n x
-  = check file line kind loc (checkIndex_msg i n) (i >= 0 && i<n) x
-
-
-checkLength_msg :: Int -> String
-{-# NOINLINE checkLength_msg #-}
-checkLength_msg n = "negative length " ++ show n
-
-checkLength :: String -> Int -> Checks -> String -> Int -> a -> a
-{-# INLINE checkLength #-}
-checkLength file line kind loc n x
-  = check file line kind loc (checkLength_msg n) (n >= 0) x
-
-
-checkSlice_msg :: Int -> Int -> Int -> String
-{-# NOINLINE checkSlice_msg #-}
-checkSlice_msg i m n = "invalid slice " ++ show (i,m,n)
-
-checkSlice :: String -> Int -> Checks -> String -> Int -> Int -> Int -> a -> a
-{-# INLINE checkSlice #-}
-checkSlice file line kind loc i m n x
-  = check file line kind loc (checkSlice_msg i m n)
-                             (i >= 0 && m >= 0 && i+m <= n) x
-
diff --git a/Setup.hs b/Setup.hs
--- a/Setup.hs
+++ b/Setup.hs
@@ -1,20 +1,3 @@
-#! /usr/bin/env runhaskell
-
-import Control.Monad
 import Distribution.Simple
-import Distribution.Simple.Setup
-import Distribution.Simple.Utils
-import System.Directory
-
-main :: IO ()
-main = defaultMainWithHooks autoconfUserHooks { preConf = preConfHook }
-  where
-    preConfHook args flags = do
-      let verbosity = fromFlag (configVerbosity flags)
-
-      confExists <- doesFileExist "configure"
-      unless confExists $
-        rawSystemExit verbosity "autoconf" []
-
-      preConf autoconfUserHooks args flags
+main = defaultMain
 
diff --git a/accelerate-cuda.buildinfo.in b/accelerate-cuda.buildinfo.in
deleted file mode 100644
--- a/accelerate-cuda.buildinfo.in
+++ /dev/null
@@ -1,3 +0,0 @@
-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,8 +1,8 @@
 Name:                   accelerate-cuda
-Version:                0.14.0.0
+Version:                0.15.0.0
 Cabal-version:          >= 1.6
-Tested-with:            GHC == 7.6.*
-Build-type:             Custom
+Tested-with:            GHC == 7.8.*
+Build-type:             Simple
 
 Synopsis:               Accelerate backend for NVIDIA GPUs
 Description:
@@ -14,7 +14,7 @@
   To use this backend you will need:
   .
     1. A CUDA-enabled NVIDIA GPU with, for full functionality, compute
-       capability 1.2 or greater. See the table on Wikipedia for supported GPUs:
+       capability 1.3 or greater. See the table on Wikipedia for supported GPUs:
        <http://en.wikipedia.org/wiki/CUDA#Supported_GPUs>
   .
     2. The CUDA SDK, available from the NVIDIA Developer Zone:
@@ -47,15 +47,6 @@
                         cubits/accelerate_cuda_type.h
                         include/AccFFI.h
 
-Extra-tmp-files:        config.status
-                        config.log
-                        autom4te.cache
-                        accelerate-cuda.buildinfo               -- generated by configure
-
-Extra-source-files:     configure
-                        accelerate-cuda.buildinfo.in
-                        include/accelerate.h
-
 Flag debug
   Description:
     Enable tracing message flags. These are read from the command-line
@@ -90,33 +81,30 @@
   Default:              False
 
 Library
-  Include-Dirs:         include
-
-  Build-depends:        accelerate              == 0.14.*,
+  Build-depends:        accelerate              == 0.15.*,
                         array                   >= 0.3,
-                        base                    == 4.6.*,
-                        binary                  >= 0.5          && < 0.7,
-                            -- binary 0.7 untested
-                        bytestring              >= 0.9          && < 0.11,
-                        cryptohash              >= 0.7          && < 0.12,
-                        cuda                    >= 0.5.1.1      && < 0.6,
+                        base                    == 4.7.*,
+                        binary                  >= 0.7,
+                        bytestring              >= 0.9,
+                        cryptohash              >= 0.7,
+                        cuda                    >= 0.6.0.2,
                         directory               >= 1.0,
-                        fclabels                >= 2.0          && < 2.1,
+                        fclabels                >= 2.0,
                         filepath                >= 1.0,
-                        hashable                >= 1.1          && < 1.3,
-                        hashtables              >= 1.0.1        && < 1.2,
-                        language-c-quote        >= 0.4.4        && < 0.8,
-                        mainland-pretty         >= 0.2          && < 0.3,
-                        mtl                     >= 2.0          && < 2.2,
+                        hashable                >= 1.1,
+                        hashtables              >= 1.0.1,
+                        language-c-quote        >= 0.4.4,
+                        mainland-pretty         >= 0.2,
+                        mtl                     >= 2.0,
                         old-time                >= 1.0,
                         pretty                  >= 1.0,
                         process                 >= 1.0,
-                        SafeSemaphore           >= 0.9          && < 0.10,
-                        srcloc                  >= 0.2          && < 0.5,
-                        text                    >= 0.11         && < 0.12,
+                        SafeSemaphore           >= 0.9,
+                        srcloc                  >= 0.2,
+                        text                    >= 0.11,
                         template-haskell        >= 2.2,
-                        transformers            >= 0.2          && < 0.4,
-                        unordered-containers    >= 0.1.4        && < 0.3
+                        transformers            >= 0.2,
+                        unordered-containers    >= 0.1.4
 
   if os(windows)
     cpp-options:        -DWIN32
@@ -158,7 +146,6 @@
                         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(debug)
diff --git a/configure b/configure
deleted file mode 100644
--- a/configure
+++ /dev/null
@@ -1,3568 +0,0 @@
-#! /bin/sh
-# Guess values for system-dependent variables and create Makefiles.
-# Generated by GNU Autoconf 2.69 for accelerate-cuda 0.14.0.0.
-#
-# Report bugs to <accelerate-haskell@googlegroups.com>.
-#
-#
-# Copyright (C) 1992-1996, 1998-2012 Free Software Foundation, Inc.
-#
-#
-# This configure script is free software; the Free Software Foundation
-# gives unlimited permission to copy, distribute and modify it.
-## -------------------- ##
-## M4sh Initialization. ##
-## -------------------- ##
-
-# Be more Bourne compatible
-DUALCASE=1; export DUALCASE # for MKS sh
-if test -n "${ZSH_VERSION+set}" && (emulate sh) >/dev/null 2>&1; then :
-  emulate sh
-  NULLCMD=:
-  # Pre-4.2 versions of Zsh do word splitting on ${1+"$@"}, which
-  # is contrary to our usage.  Disable this feature.
-  alias -g '${1+"$@"}'='"$@"'
-  setopt NO_GLOB_SUBST
-else
-  case `(set -o) 2>/dev/null` in #(
-  *posix*) :
-    set -o posix ;; #(
-  *) :
-     ;;
-esac
-fi
-
-
-as_nl='
-'
-export as_nl
-# Printing a long string crashes Solaris 7 /usr/bin/printf.
-as_echo='\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\'
-as_echo=$as_echo$as_echo$as_echo$as_echo$as_echo
-as_echo=$as_echo$as_echo$as_echo$as_echo$as_echo$as_echo
-# Prefer a ksh shell builtin over an external printf program on Solaris,
-# but without wasting forks for bash or zsh.
-if test -z "$BASH_VERSION$ZSH_VERSION" \
-    && (test "X`print -r -- $as_echo`" = "X$as_echo") 2>/dev/null; then
-  as_echo='print -r --'
-  as_echo_n='print -rn --'
-elif (test "X`printf %s $as_echo`" = "X$as_echo") 2>/dev/null; then
-  as_echo='printf %s\n'
-  as_echo_n='printf %s'
-else
-  if test "X`(/usr/ucb/echo -n -n $as_echo) 2>/dev/null`" = "X-n $as_echo"; then
-    as_echo_body='eval /usr/ucb/echo -n "$1$as_nl"'
-    as_echo_n='/usr/ucb/echo -n'
-  else
-    as_echo_body='eval expr "X$1" : "X\\(.*\\)"'
-    as_echo_n_body='eval
-      arg=$1;
-      case $arg in #(
-      *"$as_nl"*)
-	expr "X$arg" : "X\\(.*\\)$as_nl";
-	arg=`expr "X$arg" : ".*$as_nl\\(.*\\)"`;;
-      esac;
-      expr "X$arg" : "X\\(.*\\)" | tr -d "$as_nl"
-    '
-    export as_echo_n_body
-    as_echo_n='sh -c $as_echo_n_body as_echo'
-  fi
-  export as_echo_body
-  as_echo='sh -c $as_echo_body as_echo'
-fi
-
-# The user is always right.
-if test "${PATH_SEPARATOR+set}" != set; then
-  PATH_SEPARATOR=:
-  (PATH='/bin;/bin'; FPATH=$PATH; sh -c :) >/dev/null 2>&1 && {
-    (PATH='/bin:/bin'; FPATH=$PATH; sh -c :) >/dev/null 2>&1 ||
-      PATH_SEPARATOR=';'
-  }
-fi
-
-
-# IFS
-# We need space, tab and new line, in precisely that order.  Quoting is
-# there to prevent editors from complaining about space-tab.
-# (If _AS_PATH_WALK were called with IFS unset, it would disable word
-# splitting by setting IFS to empty value.)
-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
-for as_dir in $PATH
-do
-  IFS=$as_save_IFS
-  test -z "$as_dir" && as_dir=.
-    test -r "$as_dir/$0" && as_myself=$as_dir/$0 && break
-  done
-IFS=$as_save_IFS
-
-     ;;
-esac
-# We did not find ourselves, most probably we were run as `sh COMMAND'
-# in which case we are not to be found in the path.
-if test "x$as_myself" = x; then
-  as_myself=$0
-fi
-if test ! -f "$as_myself"; then
-  $as_echo "$as_myself: error: cannot find myself; rerun with an absolute file name" >&2
-  exit 1
-fi
-
-# Unset variables that we do not need and which cause bugs (e.g. in
-# pre-3.0 UWIN ksh).  But do not cause bugs in bash 2.01; the "|| exit 1"
-# suppresses any "Segmentation fault" message there.  '((' could
-# trigger a bug in pdksh 5.2.14.
-for as_var in BASH_ENV ENV MAIL MAILPATH
-do eval test x\${$as_var+set} = xset \
-  && ( (unset $as_var) || exit 1) >/dev/null 2>&1 && unset $as_var || :
-done
-PS1='$ '
-PS2='> '
-PS4='+ '
-
-# NLS nuisances.
-LC_ALL=C
-export LC_ALL
-LANGUAGE=C
-export LANGUAGE
-
-# CDPATH.
-(unset CDPATH) >/dev/null 2>&1 && unset CDPATH
-
-# Use a proper internal environment variable to ensure we don't fall
-  # into an infinite loop, continuously re-executing ourselves.
-  if test x"${_as_can_reexec}" != xno && test "x$CONFIG_SHELL" != x; then
-    _as_can_reexec=no; export _as_can_reexec;
-    # 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
-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+"$@"}
-# Admittedly, this is quite paranoid, since all the known shells bail
-# out after a failed `exec'.
-$as_echo "$0: could not re-execute with $CONFIG_SHELL" >&2
-as_fn_exit 255
-  fi
-  # We don't want this to propagate to other subprocesses.
-          { _as_can_reexec=; unset _as_can_reexec;}
-if test "x$CONFIG_SHELL" = x; then
-  as_bourne_compatible="if test -n \"\${ZSH_VERSION+set}\" && (emulate sh) >/dev/null 2>&1; then :
-  emulate sh
-  NULLCMD=:
-  # Pre-4.2 versions of Zsh do word splitting on \${1+\"\$@\"}, which
-  # is contrary to our usage.  Disable this feature.
-  alias -g '\${1+\"\$@\"}'='\"\$@\"'
-  setopt NO_GLOB_SUBST
-else
-  case \`(set -o) 2>/dev/null\` in #(
-  *posix*) :
-    set -o posix ;; #(
-  *) :
-     ;;
-esac
-fi
-"
-  as_required="as_fn_return () { (exit \$1); }
-as_fn_success () { as_fn_return 0; }
-as_fn_failure () { as_fn_return 1; }
-as_fn_ret_success () { return 0; }
-as_fn_ret_failure () { return 1; }
-
-exitcode=0
-as_fn_success || { exitcode=1; echo as_fn_success failed.; }
-as_fn_failure && { exitcode=1; echo as_fn_failure succeeded.; }
-as_fn_ret_success || { exitcode=1; echo as_fn_ret_success failed.; }
-as_fn_ret_failure && { exitcode=1; echo as_fn_ret_failure succeeded.; }
-if ( set x; as_fn_ret_success y && test x = \"\$1\" ); then :
-
-else
-  exitcode=1; echo positional parameters were not saved.
-fi
-test x\$exitcode = x0 || exit 1
-test -x / || exit 1"
-  as_suggested="  as_lineno_1=";as_suggested=$as_suggested$LINENO;as_suggested=$as_suggested" as_lineno_1a=\$LINENO
-  as_lineno_2=";as_suggested=$as_suggested$LINENO;as_suggested=$as_suggested" as_lineno_2a=\$LINENO
-  eval 'test \"x\$as_lineno_1'\$as_run'\" != \"x\$as_lineno_2'\$as_run'\" &&
-  test \"x\`expr \$as_lineno_1'\$as_run' + 1\`\" = \"x\$as_lineno_2'\$as_run'\"' || exit 1"
-  if (eval "$as_required") 2>/dev/null; then :
-  as_have_required=yes
-else
-  as_have_required=no
-fi
-  if test x$as_have_required = xyes && (eval "$as_suggested") 2>/dev/null; then :
-
-else
-  as_save_IFS=$IFS; IFS=$PATH_SEPARATOR
-as_found=false
-for as_dir in /bin$PATH_SEPARATOR/usr/bin$PATH_SEPARATOR$PATH
-do
-  IFS=$as_save_IFS
-  test -z "$as_dir" && as_dir=.
-  as_found=:
-  case $as_dir in #(
-	 /*)
-	   for as_base in sh bash ksh sh5; do
-	     # Try only shells that exist, to save several forks.
-	     as_shell=$as_dir/$as_base
-	     if { test -f "$as_shell" || test -f "$as_shell.exe"; } &&
-		    { $as_echo "$as_bourne_compatible""$as_required" | as_run=a "$as_shell"; } 2>/dev/null; then :
-  CONFIG_SHELL=$as_shell as_have_required=yes
-		   if { $as_echo "$as_bourne_compatible""$as_suggested" | as_run=a "$as_shell"; } 2>/dev/null; then :
-  break 2
-fi
-fi
-	   done;;
-       esac
-  as_found=false
-done
-$as_found || { if { test -f "$SHELL" || test -f "$SHELL.exe"; } &&
-	      { $as_echo "$as_bourne_compatible""$as_required" | as_run=a "$SHELL"; } 2>/dev/null; then :
-  CONFIG_SHELL=$SHELL as_have_required=yes
-fi; }
-IFS=$as_save_IFS
-
-
-      if test "x$CONFIG_SHELL" != x; then :
-  export CONFIG_SHELL
-             # 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
-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+"$@"}
-# Admittedly, this is quite paranoid, since all the known shells bail
-# out after a failed `exec'.
-$as_echo "$0: could not re-execute with $CONFIG_SHELL" >&2
-exit 255
-fi
-
-    if test x$as_have_required = xno; then :
-  $as_echo "$0: This script requires a shell more modern than all"
-  $as_echo "$0: the shells that I found on your system."
-  if test x${ZSH_VERSION+set} = xset ; then
-    $as_echo "$0: In particular, zsh $ZSH_VERSION has bugs and should"
-    $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-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."
-  fi
-  exit 1
-fi
-fi
-fi
-SHELL=${CONFIG_SHELL-/bin/sh}
-export SHELL
-# Unset more variables known to interfere with behavior of common tools.
-CLICOLOR_FORCE= GREP_OPTIONS=
-unset CLICOLOR_FORCE GREP_OPTIONS
-
-## --------------------- ##
-## M4sh Shell Functions. ##
-## --------------------- ##
-# as_fn_unset VAR
-# ---------------
-# Portably unset VAR.
-as_fn_unset ()
-{
-  { eval $1=; unset $1;}
-}
-as_unset=as_fn_unset
-
-# as_fn_set_status STATUS
-# -----------------------
-# Set $? to STATUS, without forking.
-as_fn_set_status ()
-{
-  return $1
-} # as_fn_set_status
-
-# as_fn_exit STATUS
-# -----------------
-# Exit the shell with STATUS, even in a "trap 0" or "set -e" context.
-as_fn_exit ()
-{
-  set +e
-  as_fn_set_status $1
-  exit $1
-} # as_fn_exit
-
-# as_fn_mkdir_p
-# -------------
-# Create "$as_dir" as a directory, including parents if necessary.
-as_fn_mkdir_p ()
-{
-
-  case $as_dir in #(
-  -*) as_dir=./$as_dir;;
-  esac
-  test -d "$as_dir" || eval $as_mkdir_p || {
-    as_dirs=
-    while :; do
-      case $as_dir in #(
-      *\'*) as_qdir=`$as_echo "$as_dir" | sed "s/'/'\\\\\\\\''/g"`;; #'(
-      *) as_qdir=$as_dir;;
-      esac
-      as_dirs="'$as_qdir' $as_dirs"
-      as_dir=`$as_dirname -- "$as_dir" ||
-$as_expr X"$as_dir" : 'X\(.*[^/]\)//*[^/][^/]*/*$' \| \
-	 X"$as_dir" : 'X\(//\)[^/]' \| \
-	 X"$as_dir" : 'X\(//\)$' \| \
-	 X"$as_dir" : 'X\(/\)' \| . 2>/dev/null ||
-$as_echo X"$as_dir" |
-    sed '/^X\(.*[^/]\)\/\/*[^/][^/]*\/*$/{
-	    s//\1/
-	    q
-	  }
-	  /^X\(\/\/\)[^/].*/{
-	    s//\1/
-	    q
-	  }
-	  /^X\(\/\/\)$/{
-	    s//\1/
-	    q
-	  }
-	  /^X\(\/\).*/{
-	    s//\1/
-	    q
-	  }
-	  s/.*/./; q'`
-      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"
-
-
-} # as_fn_mkdir_p
-
-# as_fn_executable_p FILE
-# -----------------------
-# Test if FILE is an executable regular file.
-as_fn_executable_p ()
-{
-  test -f "$1" && test -x "$1"
-} # as_fn_executable_p
-# as_fn_append VAR VALUE
-# ----------------------
-# Append the text in VALUE to the end of the definition contained in VAR. Take
-# advantage of any shell optimizations that allow amortized linear growth over
-# repeated appends, instead of the typical quadratic growth present in naive
-# implementations.
-if (eval "as_var=1; as_var+=2; test x\$as_var = x12") 2>/dev/null; then :
-  eval 'as_fn_append ()
-  {
-    eval $1+=\$2
-  }'
-else
-  as_fn_append ()
-  {
-    eval $1=\$$1\$2
-  }
-fi # as_fn_append
-
-# as_fn_arith ARG...
-# ------------------
-# Perform arithmetic evaluation on the ARGs, and store the result in the
-# global $as_val. Take advantage of shells that can avoid forks. The arguments
-# must be portable across $(()) and expr.
-if (eval "test \$(( 1 + 1 )) = 2") 2>/dev/null; then :
-  eval 'as_fn_arith ()
-  {
-    as_val=$(( $* ))
-  }'
-else
-  as_fn_arith ()
-  {
-    as_val=`expr "$@" || test $? -eq 1`
-  }
-fi # as_fn_arith
-
-
-# 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.
-as_fn_error ()
-{
-  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: $2" >&2
-  as_fn_exit $as_status
-} # as_fn_error
-
-if expr a : '\(a\)' >/dev/null 2>&1 &&
-   test "X`expr 00001 : '.*\(...\)'`" = X001; then
-  as_expr=expr
-else
-  as_expr=false
-fi
-
-if (basename -- /) >/dev/null 2>&1 && test "X`basename -- / 2>&1`" = "X/"; then
-  as_basename=basename
-else
-  as_basename=false
-fi
-
-if (as_dir=`dirname -- /` && test "X$as_dir" = X/) >/dev/null 2>&1; then
-  as_dirname=dirname
-else
-  as_dirname=false
-fi
-
-as_me=`$as_basename -- "$0" ||
-$as_expr X/"$0" : '.*/\([^/][^/]*\)/*$' \| \
-	 X"$0" : 'X\(//\)$' \| \
-	 X"$0" : 'X\(/\)' \| . 2>/dev/null ||
-$as_echo X/"$0" |
-    sed '/^.*\/\([^/][^/]*\)\/*$/{
-	    s//\1/
-	    q
-	  }
-	  /^X\/\(\/\/\)$/{
-	    s//\1/
-	    q
-	  }
-	  /^X\/\(\/\).*/{
-	    s//\1/
-	    q
-	  }
-	  s/.*/./; q'`
-
-# Avoid depending upon Character Ranges.
-as_cr_letters='abcdefghijklmnopqrstuvwxyz'
-as_cr_LETTERS='ABCDEFGHIJKLMNOPQRSTUVWXYZ'
-as_cr_Letters=$as_cr_letters$as_cr_LETTERS
-as_cr_digits='0123456789'
-as_cr_alnum=$as_cr_Letters$as_cr_digits
-
-
-  as_lineno_1=$LINENO as_lineno_1a=$LINENO
-  as_lineno_2=$LINENO as_lineno_2a=$LINENO
-  eval 'test "x$as_lineno_1'$as_run'" != "x$as_lineno_2'$as_run'" &&
-  test "x`expr $as_lineno_1'$as_run' + 1`" = "x$as_lineno_2'$as_run'"' || {
-  # Blame Lee E. McMahon (1931-1989) for sed's syntax.  :-)
-  sed -n '
-    p
-    /[$]LINENO/=
-  ' <$as_myself |
-    sed '
-      s/[$]LINENO.*/&-/
-      t lineno
-      b
-      :lineno
-      N
-      :loop
-      s/[$]LINENO\([^'$as_cr_alnum'_].*\n\)\(.*\)/\2\1\2/
-      t loop
-      s/-\n.*//
-    ' >$as_me.lineno &&
-  chmod +x "$as_me.lineno" ||
-    { $as_echo "$as_me: error: cannot create $as_me.lineno; rerun with a POSIX shell" >&2; as_fn_exit 1; }
-
-  # If we had to re-execute with $CONFIG_SHELL, we're ensured to have
-  # already done that, so ensure we don't try to do so again and fall
-  # in an infinite loop.  This has already happened in practice.
-  _as_can_reexec=no; export _as_can_reexec
-  # Don't try to exec as it changes $[0], causing all sort of problems
-  # (the dirname of $[0] is not the place where we might find the
-  # original and so on.  Autoconf is especially sensitive to this).
-  . "./$as_me.lineno"
-  # Exit status is that of the last command.
-  exit
-}
-
-ECHO_C= ECHO_N= ECHO_T=
-case `echo -n x` in #(((((
--n*)
-  case `echo 'xy\c'` in
-  *c*) ECHO_T='	';;	# ECHO_T is single tab character.
-  xy)  ECHO_C='\c';;
-  *)   echo `echo ksh88 bug on AIX 6.1` > /dev/null
-       ECHO_T='	';;
-  esac;;
-*)
-  ECHO_N='-n';;
-esac
-
-rm -f conf$$ conf$$.exe conf$$.file
-if test -d conf$$.dir; then
-  rm -f conf$$.dir/conf$$.file
-else
-  rm -f conf$$.dir
-  mkdir conf$$.dir 2>/dev/null
-fi
-if (echo >conf$$.file) 2>/dev/null; then
-  if ln -s conf$$.file conf$$ 2>/dev/null; then
-    as_ln_s='ln -s'
-    # ... but there are two gotchas:
-    # 1) On MSYS, both `ln -s file dir' and `ln file dir' fail.
-    # 2) DJGPP < 2.04 has no symlinks; `ln -s' creates a wrapper executable.
-    # In both cases, we have to default to `cp -pR'.
-    ln -s conf$$.file conf$$.dir 2>/dev/null && test ! -f conf$$.exe ||
-      as_ln_s='cp -pR'
-  elif ln conf$$.file conf$$ 2>/dev/null; then
-    as_ln_s=ln
-  else
-    as_ln_s='cp -pR'
-  fi
-else
-  as_ln_s='cp -pR'
-fi
-rm -f conf$$ conf$$.exe conf$$.dir/conf$$.file conf$$.file
-rmdir conf$$.dir 2>/dev/null
-
-if mkdir -p . 2>/dev/null; then
-  as_mkdir_p='mkdir -p "$as_dir"'
-else
-  test -d ./-p && rmdir ./-p
-  as_mkdir_p=false
-fi
-
-as_test_x='test -x'
-as_executable_p=as_fn_executable_p
-
-# Sed expression to map a string onto a valid CPP name.
-as_tr_cpp="eval sed 'y%*$as_cr_letters%P$as_cr_LETTERS%;s%[^_$as_cr_alnum]%_%g'"
-
-# Sed expression to map a string onto a valid variable name.
-as_tr_sh="eval sed 'y%*+%pp%;s%[^_$as_cr_alnum]%_%g'"
-
-
-test -n "$DJDIR" || exec 7<&0 </dev/null
-exec 6>&1
-
-# Name of the host.
-# 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`
-
-#
-# Initializations.
-#
-ac_default_prefix=/usr/local
-ac_clean_files=
-ac_config_libobj_dir=.
-LIBOBJS=
-cross_compiling=no
-subdirs=
-MFLAGS=
-MAKEFLAGS=
-
-# Identity of this package.
-PACKAGE_NAME='accelerate-cuda'
-PACKAGE_TARNAME='accelerate-cuda'
-PACKAGE_VERSION='0.14.0.0'
-PACKAGE_STRING='accelerate-cuda 0.14.0.0'
-PACKAGE_BUGREPORT='accelerate-haskell@googlegroups.com'
-PACKAGE_URL=''
-
-ac_unique_file="Data/Array/Accelerate/CUDA.hs"
-ac_subst_vars='LTLIBOBJS
-LIBOBJS
-cpp_flags
-ghc_flags
-OBJEXT
-EXEEXT
-ac_ct_CXX
-CPPFLAGS
-LDFLAGS
-CXXFLAGS
-CXX
-GHC
-target_alias
-host_alias
-build_alias
-LIBS
-ECHO_T
-ECHO_N
-ECHO_C
-DEFS
-mandir
-localedir
-libdir
-psdir
-pdfdir
-dvidir
-htmldir
-infodir
-docdir
-oldincludedir
-includedir
-localstatedir
-sharedstatedir
-sysconfdir
-datadir
-datarootdir
-libexecdir
-sbindir
-bindir
-program_transform_name
-prefix
-exec_prefix
-PACKAGE_URL
-PACKAGE_BUGREPORT
-PACKAGE_STRING
-PACKAGE_VERSION
-PACKAGE_TARNAME
-PACKAGE_NAME
-PATH_SEPARATOR
-SHELL'
-ac_subst_files=''
-ac_user_opts='
-enable_option_checking
-with_compiler
-with_nvcc
-with_gcc
-'
-      ac_precious_vars='build_alias
-host_alias
-target_alias
-CXX
-CXXFLAGS
-LDFLAGS
-LIBS
-CPPFLAGS
-CCC'
-
-
-# Initialize some variables set by options.
-ac_init_help=
-ac_init_version=false
-ac_unrecognized_opts=
-ac_unrecognized_sep=
-# The variables have the same names as the options, with
-# dashes changed to underlines.
-cache_file=/dev/null
-exec_prefix=NONE
-no_create=
-no_recursion=
-prefix=NONE
-program_prefix=NONE
-program_suffix=NONE
-program_transform_name=s,x,x,
-silent=
-site=
-srcdir=
-verbose=
-x_includes=NONE
-x_libraries=NONE
-
-# Installation directory options.
-# These are left unexpanded so users can "make install exec_prefix=/foo"
-# and all the variables that are supposed to be based on exec_prefix
-# by default will actually change.
-# Use braces instead of parens because sh, perl, etc. also accept them.
-# (The list follows the same order as the GNU Coding Standards.)
-bindir='${exec_prefix}/bin'
-sbindir='${exec_prefix}/sbin'
-libexecdir='${exec_prefix}/libexec'
-datarootdir='${prefix}/share'
-datadir='${datarootdir}'
-sysconfdir='${prefix}/etc'
-sharedstatedir='${prefix}/com'
-localstatedir='${prefix}/var'
-includedir='${prefix}/include'
-oldincludedir='/usr/include'
-docdir='${datarootdir}/doc/${PACKAGE_TARNAME}'
-infodir='${datarootdir}/info'
-htmldir='${docdir}'
-dvidir='${docdir}'
-pdfdir='${docdir}'
-psdir='${docdir}'
-libdir='${exec_prefix}/lib'
-localedir='${datarootdir}/locale'
-mandir='${datarootdir}/man'
-
-ac_prev=
-ac_dashdash=
-for ac_option
-do
-  # If the previous option needs an argument, assign it.
-  if test -n "$ac_prev"; then
-    eval $ac_prev=\$ac_option
-    ac_prev=
-    continue
-  fi
-
-  case $ac_option in
-  *=?*) ac_optarg=`expr "X$ac_option" : '[^=]*=\(.*\)'` ;;
-  *=)   ac_optarg= ;;
-  *)    ac_optarg=yes ;;
-  esac
-
-  # Accept the important Cygnus configure options, so we can diagnose typos.
-
-  case $ac_dashdash$ac_option in
-  --)
-    ac_dashdash=yes ;;
-
-  -bindir | --bindir | --bindi | --bind | --bin | --bi)
-    ac_prev=bindir ;;
-  -bindir=* | --bindir=* | --bindi=* | --bind=* | --bin=* | --bi=*)
-    bindir=$ac_optarg ;;
-
-  -build | --build | --buil | --bui | --bu)
-    ac_prev=build_alias ;;
-  -build=* | --build=* | --buil=* | --bui=* | --bu=*)
-    build_alias=$ac_optarg ;;
-
-  -cache-file | --cache-file | --cache-fil | --cache-fi \
-  | --cache-f | --cache- | --cache | --cach | --cac | --ca | --c)
-    ac_prev=cache_file ;;
-  -cache-file=* | --cache-file=* | --cache-fil=* | --cache-fi=* \
-  | --cache-f=* | --cache-=* | --cache=* | --cach=* | --cac=* | --ca=* | --c=*)
-    cache_file=$ac_optarg ;;
-
-  --config-cache | -C)
-    cache_file=config.cache ;;
-
-  -datadir | --datadir | --datadi | --datad)
-    ac_prev=datadir ;;
-  -datadir=* | --datadir=* | --datadi=* | --datad=*)
-    datadir=$ac_optarg ;;
-
-  -datarootdir | --datarootdir | --datarootdi | --datarootd | --dataroot \
-  | --dataroo | --dataro | --datar)
-    ac_prev=datarootdir ;;
-  -datarootdir=* | --datarootdir=* | --datarootdi=* | --datarootd=* \
-  | --dataroot=* | --dataroo=* | --dataro=* | --datar=*)
-    datarootdir=$ac_optarg ;;
-
-  -disable-* | --disable-*)
-    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"
-    ac_useropt_orig=$ac_useropt
-    ac_useropt=`$as_echo "$ac_useropt" | sed 's/[-+.]/_/g'`
-    case $ac_user_opts in
-      *"
-"enable_$ac_useropt"
-"*) ;;
-      *) ac_unrecognized_opts="$ac_unrecognized_opts$ac_unrecognized_sep--disable-$ac_useropt_orig"
-	 ac_unrecognized_sep=', ';;
-    esac
-    eval enable_$ac_useropt=no ;;
-
-  -docdir | --docdir | --docdi | --doc | --do)
-    ac_prev=docdir ;;
-  -docdir=* | --docdir=* | --docdi=* | --doc=* | --do=*)
-    docdir=$ac_optarg ;;
-
-  -dvidir | --dvidir | --dvidi | --dvid | --dvi | --dv)
-    ac_prev=dvidir ;;
-  -dvidir=* | --dvidir=* | --dvidi=* | --dvid=* | --dvi=* | --dv=*)
-    dvidir=$ac_optarg ;;
-
-  -enable-* | --enable-*)
-    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"
-    ac_useropt_orig=$ac_useropt
-    ac_useropt=`$as_echo "$ac_useropt" | sed 's/[-+.]/_/g'`
-    case $ac_user_opts in
-      *"
-"enable_$ac_useropt"
-"*) ;;
-      *) ac_unrecognized_opts="$ac_unrecognized_opts$ac_unrecognized_sep--enable-$ac_useropt_orig"
-	 ac_unrecognized_sep=', ';;
-    esac
-    eval enable_$ac_useropt=\$ac_optarg ;;
-
-  -exec-prefix | --exec_prefix | --exec-prefix | --exec-prefi \
-  | --exec-pref | --exec-pre | --exec-pr | --exec-p | --exec- \
-  | --exec | --exe | --ex)
-    ac_prev=exec_prefix ;;
-  -exec-prefix=* | --exec_prefix=* | --exec-prefix=* | --exec-prefi=* \
-  | --exec-pref=* | --exec-pre=* | --exec-pr=* | --exec-p=* | --exec-=* \
-  | --exec=* | --exe=* | --ex=*)
-    exec_prefix=$ac_optarg ;;
-
-  -gas | --gas | --ga | --g)
-    # Obsolete; use --with-gas.
-    with_gas=yes ;;
-
-  -help | --help | --hel | --he | -h)
-    ac_init_help=long ;;
-  -help=r* | --help=r* | --hel=r* | --he=r* | -hr*)
-    ac_init_help=recursive ;;
-  -help=s* | --help=s* | --hel=s* | --he=s* | -hs*)
-    ac_init_help=short ;;
-
-  -host | --host | --hos | --ho)
-    ac_prev=host_alias ;;
-  -host=* | --host=* | --hos=* | --ho=*)
-    host_alias=$ac_optarg ;;
-
-  -htmldir | --htmldir | --htmldi | --htmld | --html | --htm | --ht)
-    ac_prev=htmldir ;;
-  -htmldir=* | --htmldir=* | --htmldi=* | --htmld=* | --html=* | --htm=* \
-  | --ht=*)
-    htmldir=$ac_optarg ;;
-
-  -includedir | --includedir | --includedi | --included | --include \
-  | --includ | --inclu | --incl | --inc)
-    ac_prev=includedir ;;
-  -includedir=* | --includedir=* | --includedi=* | --included=* | --include=* \
-  | --includ=* | --inclu=* | --incl=* | --inc=*)
-    includedir=$ac_optarg ;;
-
-  -infodir | --infodir | --infodi | --infod | --info | --inf)
-    ac_prev=infodir ;;
-  -infodir=* | --infodir=* | --infodi=* | --infod=* | --info=* | --inf=*)
-    infodir=$ac_optarg ;;
-
-  -libdir | --libdir | --libdi | --libd)
-    ac_prev=libdir ;;
-  -libdir=* | --libdir=* | --libdi=* | --libd=*)
-    libdir=$ac_optarg ;;
-
-  -libexecdir | --libexecdir | --libexecdi | --libexecd | --libexec \
-  | --libexe | --libex | --libe)
-    ac_prev=libexecdir ;;
-  -libexecdir=* | --libexecdir=* | --libexecdi=* | --libexecd=* | --libexec=* \
-  | --libexe=* | --libex=* | --libe=*)
-    libexecdir=$ac_optarg ;;
-
-  -localedir | --localedir | --localedi | --localed | --locale)
-    ac_prev=localedir ;;
-  -localedir=* | --localedir=* | --localedi=* | --localed=* | --locale=*)
-    localedir=$ac_optarg ;;
-
-  -localstatedir | --localstatedir | --localstatedi | --localstated \
-  | --localstate | --localstat | --localsta | --localst | --locals)
-    ac_prev=localstatedir ;;
-  -localstatedir=* | --localstatedir=* | --localstatedi=* | --localstated=* \
-  | --localstate=* | --localstat=* | --localsta=* | --localst=* | --locals=*)
-    localstatedir=$ac_optarg ;;
-
-  -mandir | --mandir | --mandi | --mand | --man | --ma | --m)
-    ac_prev=mandir ;;
-  -mandir=* | --mandir=* | --mandi=* | --mand=* | --man=* | --ma=* | --m=*)
-    mandir=$ac_optarg ;;
-
-  -nfp | --nfp | --nf)
-    # Obsolete; use --without-fp.
-    with_fp=no ;;
-
-  -no-create | --no-create | --no-creat | --no-crea | --no-cre \
-  | --no-cr | --no-c | -n)
-    no_create=yes ;;
-
-  -no-recursion | --no-recursion | --no-recursio | --no-recursi \
-  | --no-recurs | --no-recur | --no-recu | --no-rec | --no-re | --no-r)
-    no_recursion=yes ;;
-
-  -oldincludedir | --oldincludedir | --oldincludedi | --oldincluded \
-  | --oldinclude | --oldinclud | --oldinclu | --oldincl | --oldinc \
-  | --oldin | --oldi | --old | --ol | --o)
-    ac_prev=oldincludedir ;;
-  -oldincludedir=* | --oldincludedir=* | --oldincludedi=* | --oldincluded=* \
-  | --oldinclude=* | --oldinclud=* | --oldinclu=* | --oldincl=* | --oldinc=* \
-  | --oldin=* | --oldi=* | --old=* | --ol=* | --o=*)
-    oldincludedir=$ac_optarg ;;
-
-  -prefix | --prefix | --prefi | --pref | --pre | --pr | --p)
-    ac_prev=prefix ;;
-  -prefix=* | --prefix=* | --prefi=* | --pref=* | --pre=* | --pr=* | --p=*)
-    prefix=$ac_optarg ;;
-
-  -program-prefix | --program-prefix | --program-prefi | --program-pref \
-  | --program-pre | --program-pr | --program-p)
-    ac_prev=program_prefix ;;
-  -program-prefix=* | --program-prefix=* | --program-prefi=* \
-  | --program-pref=* | --program-pre=* | --program-pr=* | --program-p=*)
-    program_prefix=$ac_optarg ;;
-
-  -program-suffix | --program-suffix | --program-suffi | --program-suff \
-  | --program-suf | --program-su | --program-s)
-    ac_prev=program_suffix ;;
-  -program-suffix=* | --program-suffix=* | --program-suffi=* \
-  | --program-suff=* | --program-suf=* | --program-su=* | --program-s=*)
-    program_suffix=$ac_optarg ;;
-
-  -program-transform-name | --program-transform-name \
-  | --program-transform-nam | --program-transform-na \
-  | --program-transform-n | --program-transform- \
-  | --program-transform | --program-transfor \
-  | --program-transfo | --program-transf \
-  | --program-trans | --program-tran \
-  | --progr-tra | --program-tr | --program-t)
-    ac_prev=program_transform_name ;;
-  -program-transform-name=* | --program-transform-name=* \
-  | --program-transform-nam=* | --program-transform-na=* \
-  | --program-transform-n=* | --program-transform-=* \
-  | --program-transform=* | --program-transfor=* \
-  | --program-transfo=* | --program-transf=* \
-  | --program-trans=* | --program-tran=* \
-  | --progr-tra=* | --program-tr=* | --program-t=*)
-    program_transform_name=$ac_optarg ;;
-
-  -pdfdir | --pdfdir | --pdfdi | --pdfd | --pdf | --pd)
-    ac_prev=pdfdir ;;
-  -pdfdir=* | --pdfdir=* | --pdfdi=* | --pdfd=* | --pdf=* | --pd=*)
-    pdfdir=$ac_optarg ;;
-
-  -psdir | --psdir | --psdi | --psd | --ps)
-    ac_prev=psdir ;;
-  -psdir=* | --psdir=* | --psdi=* | --psd=* | --ps=*)
-    psdir=$ac_optarg ;;
-
-  -q | -quiet | --quiet | --quie | --qui | --qu | --q \
-  | -silent | --silent | --silen | --sile | --sil)
-    silent=yes ;;
-
-  -sbindir | --sbindir | --sbindi | --sbind | --sbin | --sbi | --sb)
-    ac_prev=sbindir ;;
-  -sbindir=* | --sbindir=* | --sbindi=* | --sbind=* | --sbin=* \
-  | --sbi=* | --sb=*)
-    sbindir=$ac_optarg ;;
-
-  -sharedstatedir | --sharedstatedir | --sharedstatedi \
-  | --sharedstated | --sharedstate | --sharedstat | --sharedsta \
-  | --sharedst | --shareds | --shared | --share | --shar \
-  | --sha | --sh)
-    ac_prev=sharedstatedir ;;
-  -sharedstatedir=* | --sharedstatedir=* | --sharedstatedi=* \
-  | --sharedstated=* | --sharedstate=* | --sharedstat=* | --sharedsta=* \
-  | --sharedst=* | --shareds=* | --shared=* | --share=* | --shar=* \
-  | --sha=* | --sh=*)
-    sharedstatedir=$ac_optarg ;;
-
-  -site | --site | --sit)
-    ac_prev=site ;;
-  -site=* | --site=* | --sit=*)
-    site=$ac_optarg ;;
-
-  -srcdir | --srcdir | --srcdi | --srcd | --src | --sr)
-    ac_prev=srcdir ;;
-  -srcdir=* | --srcdir=* | --srcdi=* | --srcd=* | --src=* | --sr=*)
-    srcdir=$ac_optarg ;;
-
-  -sysconfdir | --sysconfdir | --sysconfdi | --sysconfd | --sysconf \
-  | --syscon | --sysco | --sysc | --sys | --sy)
-    ac_prev=sysconfdir ;;
-  -sysconfdir=* | --sysconfdir=* | --sysconfdi=* | --sysconfd=* | --sysconf=* \
-  | --syscon=* | --sysco=* | --sysc=* | --sys=* | --sy=*)
-    sysconfdir=$ac_optarg ;;
-
-  -target | --target | --targe | --targ | --tar | --ta | --t)
-    ac_prev=target_alias ;;
-  -target=* | --target=* | --targe=* | --targ=* | --tar=* | --ta=* | --t=*)
-    target_alias=$ac_optarg ;;
-
-  -v | -verbose | --verbose | --verbos | --verbo | --verb)
-    verbose=yes ;;
-
-  -version | --version | --versio | --versi | --vers | -V)
-    ac_init_version=: ;;
-
-  -with-* | --with-*)
-    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"
-    ac_useropt_orig=$ac_useropt
-    ac_useropt=`$as_echo "$ac_useropt" | sed 's/[-+.]/_/g'`
-    case $ac_user_opts in
-      *"
-"with_$ac_useropt"
-"*) ;;
-      *) ac_unrecognized_opts="$ac_unrecognized_opts$ac_unrecognized_sep--with-$ac_useropt_orig"
-	 ac_unrecognized_sep=', ';;
-    esac
-    eval with_$ac_useropt=\$ac_optarg ;;
-
-  -without-* | --without-*)
-    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"
-    ac_useropt_orig=$ac_useropt
-    ac_useropt=`$as_echo "$ac_useropt" | sed 's/[-+.]/_/g'`
-    case $ac_user_opts in
-      *"
-"with_$ac_useropt"
-"*) ;;
-      *) ac_unrecognized_opts="$ac_unrecognized_opts$ac_unrecognized_sep--without-$ac_useropt_orig"
-	 ac_unrecognized_sep=', ';;
-    esac
-    eval with_$ac_useropt=no ;;
-
-  --x)
-    # Obsolete; use --with-x.
-    with_x=yes ;;
-
-  -x-includes | --x-includes | --x-include | --x-includ | --x-inclu \
-  | --x-incl | --x-inc | --x-in | --x-i)
-    ac_prev=x_includes ;;
-  -x-includes=* | --x-includes=* | --x-include=* | --x-includ=* | --x-inclu=* \
-  | --x-incl=* | --x-inc=* | --x-in=* | --x-i=*)
-    x_includes=$ac_optarg ;;
-
-  -x-libraries | --x-libraries | --x-librarie | --x-librari \
-  | --x-librar | --x-libra | --x-libr | --x-lib | --x-li | --x-l)
-    ac_prev=x_libraries ;;
-  -x-libraries=* | --x-libraries=* | --x-librarie=* | --x-librari=* \
-  | --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"
-    ;;
-
-  *=*)
-    ac_envvar=`expr "x$ac_option" : 'x\([^=]*\)='`
-    # 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'" ;;
-    esac
-    eval $ac_envvar=\$ac_optarg
-    export $ac_envvar ;;
-
-  *)
-    # FIXME: should be removed in autoconf 3.0.
-    $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}"
-    ;;
-
-  esac
-done
-
-if test -n "$ac_prev"; then
-  ac_option=--`echo $ac_prev | sed 's/_/-/g'`
-  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" ;;
-    *)     $as_echo "$as_me: WARNING: unrecognized options: $ac_unrecognized_opts" >&2 ;;
-  esac
-fi
-
-# Check all directory arguments for consistency.
-for ac_var in	exec_prefix prefix bindir sbindir libexecdir datarootdir \
-		datadir sysconfdir sharedstatedir localstatedir includedir \
-		oldincludedir docdir infodir htmldir dvidir pdfdir psdir \
-		libdir localedir mandir
-do
-  eval ac_val=\$$ac_var
-  # Remove trailing slashes.
-  case $ac_val in
-    */ )
-      ac_val=`expr "X$ac_val" : 'X\(.*[^/]\)' \| "X$ac_val" : 'X\(.*\)'`
-      eval $ac_var=\$ac_val;;
-  esac
-  # Be sure to have absolute directory names.
-  case $ac_val in
-    [\\/$]* | ?:[\\/]* )  continue;;
-    NONE | '' ) case $ac_var in *prefix ) continue;; esac;;
-  esac
-  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'
-# used to hold the argument of --host etc.
-# FIXME: To remove some day.
-build=$build_alias
-host=$host_alias
-target=$target_alias
-
-# FIXME: To remove some day.
-if test "x$host_alias" != x; then
-  if test "x$build_alias" = x; then
-    cross_compiling=maybe
-  elif test "x$build_alias" != "x$host_alias"; then
-    cross_compiling=yes
-  fi
-fi
-
-ac_tool_prefix=
-test -n "$host_alias" && ac_tool_prefix=$host_alias-
-
-test "$silent" = yes && exec 6>/dev/null
-
-
-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"
-test "X$ac_ls_di" = "X$ac_pwd_ls_di" ||
-  as_fn_error $? "pwd does not report name of working directory"
-
-
-# Find the source files, if location was not specified.
-if test -z "$srcdir"; then
-  ac_srcdir_defaulted=yes
-  # Try the directory containing this script, then the parent directory.
-  ac_confdir=`$as_dirname -- "$as_myself" ||
-$as_expr X"$as_myself" : 'X\(.*[^/]\)//*[^/][^/]*/*$' \| \
-	 X"$as_myself" : 'X\(//\)[^/]' \| \
-	 X"$as_myself" : 'X\(//\)$' \| \
-	 X"$as_myself" : 'X\(/\)' \| . 2>/dev/null ||
-$as_echo X"$as_myself" |
-    sed '/^X\(.*[^/]\)\/\/*[^/][^/]*\/*$/{
-	    s//\1/
-	    q
-	  }
-	  /^X\(\/\/\)[^/].*/{
-	    s//\1/
-	    q
-	  }
-	  /^X\(\/\/\)$/{
-	    s//\1/
-	    q
-	  }
-	  /^X\(\/\).*/{
-	    s//\1/
-	    q
-	  }
-	  s/.*/./; q'`
-  srcdir=$ac_confdir
-  if test ! -r "$srcdir/$ac_unique_file"; then
-    srcdir=..
-  fi
-else
-  ac_srcdir_defaulted=no
-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"
-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"
-	pwd)`
-# When building in place, set srcdir=.
-if test "$ac_abs_confdir" = "$ac_pwd"; then
-  srcdir=.
-fi
-# Remove unnecessary trailing slashes from srcdir.
-# Double slashes in file names in object file debugging info
-# mess up M-x gdb in Emacs.
-case $srcdir in
-*/) srcdir=`expr "X$srcdir" : 'X\(.*[^/]\)' \| "X$srcdir" : 'X\(.*\)'`;;
-esac
-for ac_var in $ac_precious_vars; do
-  eval ac_env_${ac_var}_set=\${${ac_var}+set}
-  eval ac_env_${ac_var}_value=\$${ac_var}
-  eval ac_cv_env_${ac_var}_set=\${${ac_var}+set}
-  eval ac_cv_env_${ac_var}_value=\$${ac_var}
-done
-
-#
-# Report the --help message.
-#
-if test "$ac_init_help" = "long"; then
-  # 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.14.0.0 to adapt to many kinds of systems.
-
-Usage: $0 [OPTION]... [VAR=VALUE]...
-
-To assign environment variables (e.g., CC, CFLAGS...), specify them as
-VAR=VALUE.  See below for descriptions of some of the useful variables.
-
-Defaults for the options are specified in brackets.
-
-Configuration:
-  -h, --help              display this help and exit
-      --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
-      --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
-      --srcdir=DIR        find the sources in DIR [configure dir or \`..']
-
-Installation directories:
-  --prefix=PREFIX         install architecture-independent files in PREFIX
-                          [$ac_default_prefix]
-  --exec-prefix=EPREFIX   install architecture-dependent files in EPREFIX
-                          [PREFIX]
-
-By default, \`make install' will install all the files in
-\`$ac_default_prefix/bin', \`$ac_default_prefix/lib' etc.  You can specify
-an installation prefix other than \`$ac_default_prefix' using \`--prefix',
-for instance \`--prefix=\$HOME'.
-
-For better control, use the options below.
-
-Fine tuning of the installation directories:
-  --bindir=DIR            user executables [EPREFIX/bin]
-  --sbindir=DIR           system admin executables [EPREFIX/sbin]
-  --libexecdir=DIR        program executables [EPREFIX/libexec]
-  --sysconfdir=DIR        read-only single-machine data [PREFIX/etc]
-  --sharedstatedir=DIR    modifiable architecture-independent data [PREFIX/com]
-  --localstatedir=DIR     modifiable single-machine data [PREFIX/var]
-  --libdir=DIR            object code libraries [EPREFIX/lib]
-  --includedir=DIR        C header files [PREFIX/include]
-  --oldincludedir=DIR     C header files for non-gcc [/usr/include]
-  --datarootdir=DIR       read-only arch.-independent data root [PREFIX/share]
-  --datadir=DIR           read-only architecture-independent data [DATAROOTDIR]
-  --infodir=DIR           info documentation [DATAROOTDIR/info]
-  --localedir=DIR         locale-dependent data [DATAROOTDIR/locale]
-  --mandir=DIR            man documentation [DATAROOTDIR/man]
-  --docdir=DIR            documentation root [DATAROOTDIR/doc/accelerate-cuda]
-  --htmldir=DIR           html documentation [DOCDIR]
-  --dvidir=DIR            dvi documentation [DOCDIR]
-  --pdfdir=DIR            pdf documentation [DOCDIR]
-  --psdir=DIR             ps documentation [DOCDIR]
-_ACEOF
-
-  cat <<\_ACEOF
-_ACEOF
-fi
-
-if test -n "$ac_init_help"; then
-  case $ac_init_help in
-     short | recursive ) echo "Configuration of accelerate-cuda 0.14.0.0:";;
-   esac
-  cat <<\_ACEOF
-
-Optional Packages:
-  --with-PACKAGE[=ARG]    use PACKAGE [ARG=yes]
-  --without-PACKAGE       do not use PACKAGE (same as --with-PACKAGE=no)
-Haskell compiler
-CUDA compiler
-C compiler
-
-Some influential environment variables:
-  CXX         C++ compiler command
-  CXXFLAGS    C++ compiler flags
-  LDFLAGS     linker flags, e.g. -L<lib dir> if you have libraries in a
-              nonstandard directory <lib dir>
-  LIBS        libraries to pass to the linker, e.g. -l<library>
-  CPPFLAGS    (Objective) C/C++ preprocessor flags, e.g. -I<include dir> if
-              you have headers in a nonstandard directory <include dir>
-
-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-haskell@googlegroups.com>.
-_ACEOF
-ac_status=$?
-fi
-
-if test "$ac_init_help" = "recursive"; then
-  # If there are subdirs, report their specific --help.
-  for ac_dir in : $ac_subdirs_all; do test "x$ac_dir" = x: && continue
-    test -d "$ac_dir" ||
-      { cd "$srcdir" && ac_pwd=`pwd` && srcdir=. && test -d "$ac_dir"; } ||
-      continue
-    ac_builddir=.
-
-case "$ac_dir" in
-.) ac_dir_suffix= ac_top_builddir_sub=. ac_top_build_prefix= ;;
-*)
-  ac_dir_suffix=/`$as_echo "$ac_dir" | sed 's|^\.[\\/]||'`
-  # A ".." for each directory in $ac_dir_suffix.
-  ac_top_builddir_sub=`$as_echo "$ac_dir_suffix" | sed 's|/[^\\/]*|/..|g;s|/||'`
-  case $ac_top_builddir_sub in
-  "") ac_top_builddir_sub=. ac_top_build_prefix= ;;
-  *)  ac_top_build_prefix=$ac_top_builddir_sub/ ;;
-  esac ;;
-esac
-ac_abs_top_builddir=$ac_pwd
-ac_abs_builddir=$ac_pwd$ac_dir_suffix
-# for backward compatibility:
-ac_top_builddir=$ac_top_build_prefix
-
-case $srcdir in
-  .)  # We are building in place.
-    ac_srcdir=.
-    ac_top_srcdir=$ac_top_builddir_sub
-    ac_abs_top_srcdir=$ac_pwd ;;
-  [\\/]* | ?:[\\/]* )  # Absolute name.
-    ac_srcdir=$srcdir$ac_dir_suffix;
-    ac_top_srcdir=$srcdir
-    ac_abs_top_srcdir=$srcdir ;;
-  *) # Relative name.
-    ac_srcdir=$ac_top_build_prefix$srcdir$ac_dir_suffix
-    ac_top_srcdir=$ac_top_build_prefix$srcdir
-    ac_abs_top_srcdir=$ac_pwd/$srcdir ;;
-esac
-ac_abs_srcdir=$ac_abs_top_srcdir$ac_dir_suffix
-
-    cd "$ac_dir" || { ac_status=$?; continue; }
-    # Check for guested configure.
-    if test -f "$ac_srcdir/configure.gnu"; then
-      echo &&
-      $SHELL "$ac_srcdir/configure.gnu" --help=recursive
-    elif test -f "$ac_srcdir/configure"; then
-      echo &&
-      $SHELL "$ac_srcdir/configure" --help=recursive
-    else
-      $as_echo "$as_me: WARNING: no configuration information is in $ac_dir" >&2
-    fi || ac_status=$?
-    cd "$ac_pwd" || { ac_status=$?; break; }
-  done
-fi
-
-test -n "$ac_init_help" && exit $ac_status
-if $ac_init_version; then
-  cat <<\_ACEOF
-accelerate-cuda configure 0.14.0.0
-generated by GNU Autoconf 2.69
-
-Copyright (C) 2012 Free Software Foundation, Inc.
-This configure script is free software; the Free Software Foundation
-gives unlimited permission to copy, distribute and modify it.
-_ACEOF
-  exit
-fi
-
-## ------------------------ ##
-## Autoconf initialization. ##
-## ------------------------ ##
-
-# ac_fn_cxx_try_compile LINENO
-# ----------------------------
-# Try to compile conftest.$ac_ext, and return whether this succeeded.
-ac_fn_cxx_try_compile ()
-{
-  as_lineno=${as_lineno-"$1"} as_lineno_stack=as_lineno_stack=$as_lineno_stack
-  rm -f conftest.$ac_objext
-  if { { ac_try="$ac_compile"
-case "(($ac_try" in
-  *\"* | *\`* | *\\*) ac_try_echo=\$ac_try;;
-  *) ac_try_echo=$ac_try;;
-esac
-eval ac_try_echo="\"\$as_me:${as_lineno-$LINENO}: $ac_try_echo\""
-$as_echo "$ac_try_echo"; } >&5
-  (eval "$ac_compile") 2>conftest.err
-  ac_status=$?
-  if test -s conftest.err; then
-    grep -v '^ *+' conftest.err >conftest.er1
-    cat conftest.er1 >&5
-    mv -f conftest.er1 conftest.err
-  fi
-  $as_echo "$as_me:${as_lineno-$LINENO}: \$? = $ac_status" >&5
-  test $ac_status = 0; } && {
-	 test -z "$ac_cxx_werror_flag" ||
-	 test ! -s conftest.err
-       } && test -s conftest.$ac_objext; then :
-  ac_retval=0
-else
-  $as_echo "$as_me: failed program was:" >&5
-sed 's/^/| /' conftest.$ac_ext >&5
-
-	ac_retval=1
-fi
-  eval $as_lineno_stack; ${as_lineno_stack:+:} unset as_lineno
-  as_fn_set_status $ac_retval
-
-} # ac_fn_cxx_try_compile
-cat >config.log <<_ACEOF
-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.14.0.0, which was
-generated by GNU Autoconf 2.69.  Invocation command line was
-
-  $ $0 $@
-
-_ACEOF
-exec 5>>config.log
-{
-cat <<_ASUNAME
-## --------- ##
-## Platform. ##
-## --------- ##
-
-hostname = `(hostname || uname -n) 2>/dev/null | sed 1q`
-uname -m = `(uname -m) 2>/dev/null || echo unknown`
-uname -r = `(uname -r) 2>/dev/null || echo unknown`
-uname -s = `(uname -s) 2>/dev/null || echo unknown`
-uname -v = `(uname -v) 2>/dev/null || echo unknown`
-
-/usr/bin/uname -p = `(/usr/bin/uname -p) 2>/dev/null || echo unknown`
-/bin/uname -X     = `(/bin/uname -X) 2>/dev/null     || echo unknown`
-
-/bin/arch              = `(/bin/arch) 2>/dev/null              || echo unknown`
-/usr/bin/arch -k       = `(/usr/bin/arch -k) 2>/dev/null       || echo unknown`
-/usr/convex/getsysinfo = `(/usr/convex/getsysinfo) 2>/dev/null || echo unknown`
-/usr/bin/hostinfo      = `(/usr/bin/hostinfo) 2>/dev/null      || echo unknown`
-/bin/machine           = `(/bin/machine) 2>/dev/null           || echo unknown`
-/usr/bin/oslevel       = `(/usr/bin/oslevel) 2>/dev/null       || echo unknown`
-/bin/universe          = `(/bin/universe) 2>/dev/null          || echo unknown`
-
-_ASUNAME
-
-as_save_IFS=$IFS; IFS=$PATH_SEPARATOR
-for as_dir in $PATH
-do
-  IFS=$as_save_IFS
-  test -z "$as_dir" && as_dir=.
-    $as_echo "PATH: $as_dir"
-  done
-IFS=$as_save_IFS
-
-} >&5
-
-cat >&5 <<_ACEOF
-
-
-## ----------- ##
-## Core tests. ##
-## ----------- ##
-
-_ACEOF
-
-
-# Keep a trace of the command line.
-# Strip out --no-create and --no-recursion so they do not pile up.
-# Strip out --silent because we don't want to record it for future runs.
-# Also quote any args containing shell meta-characters.
-# Make two passes to allow for proper duplicate-argument suppression.
-ac_configure_args=
-ac_configure_args0=
-ac_configure_args1=
-ac_must_keep_next=false
-for ac_pass in 1 2
-do
-  for ac_arg
-  do
-    case $ac_arg in
-    -no-create | --no-c* | -n | -no-recursion | --no-r*) continue ;;
-    -q | -quiet | --quiet | --quie | --qui | --qu | --q \
-    | -silent | --silent | --silen | --sile | --sil)
-      continue ;;
-    *\'*)
-      ac_arg=`$as_echo "$ac_arg" | sed "s/'/'\\\\\\\\''/g"` ;;
-    esac
-    case $ac_pass in
-    1) as_fn_append ac_configure_args0 " '$ac_arg'" ;;
-    2)
-      as_fn_append ac_configure_args1 " '$ac_arg'"
-      if test $ac_must_keep_next = true; then
-	ac_must_keep_next=false # Got value, back to normal.
-      else
-	case $ac_arg in
-	  *=* | --config-cache | -C | -disable-* | --disable-* \
-	  | -enable-* | --enable-* | -gas | --g* | -nfp | --nf* \
-	  | -q | -quiet | --q* | -silent | --sil* | -v | -verb* \
-	  | -with-* | --with-* | -without-* | --without-* | --x)
-	    case "$ac_configure_args0 " in
-	      "$ac_configure_args1"*" '$ac_arg' "* ) continue ;;
-	    esac
-	    ;;
-	  -* ) ac_must_keep_next=true ;;
-	esac
-      fi
-      as_fn_append ac_configure_args " '$ac_arg'"
-      ;;
-    esac
-  done
-done
-{ ac_configure_args0=; unset ac_configure_args0;}
-{ ac_configure_args1=; unset ac_configure_args1;}
-
-# When interrupted or exit'd, cleanup temporary files, and complete
-# config.log.  We remove comments because anyway the quotes in there
-# would cause problems or look ugly.
-# WARNING: Use '\'' to represent an apostrophe within the trap.
-# WARNING: Do not start the trap code with a newline, due to a FreeBSD 4.0 bug.
-trap 'exit_status=$?
-  # Save into config.log some information that might help in debugging.
-  {
-    echo
-
-    $as_echo "## ---------------- ##
-## Cache variables. ##
-## ---------------- ##"
-    echo
-    # The following way of writing the cache mishandles newlines in values,
-(
-  for ac_var in `(set) 2>&1 | sed -n '\''s/^\([a-zA-Z_][a-zA-Z0-9_]*\)=.*/\1/p'\''`; do
-    eval ac_val=\$$ac_var
-    case $ac_val in #(
-    *${as_nl}*)
-      case $ac_var in #(
-      *_cv_*) { $as_echo "$as_me:${as_lineno-$LINENO}: WARNING: cache variable $ac_var contains a newline" >&5
-$as_echo "$as_me: WARNING: cache variable $ac_var contains a newline" >&2;} ;;
-      esac
-      case $ac_var in #(
-      _ | IFS | as_nl) ;; #(
-      BASH_ARGV | BASH_SOURCE) eval $ac_var= ;; #(
-      *) { eval $ac_var=; unset $ac_var;} ;;
-      esac ;;
-    esac
-  done
-  (set) 2>&1 |
-    case $as_nl`(ac_space='\'' '\''; set) 2>&1` in #(
-    *${as_nl}ac_space=\ *)
-      sed -n \
-	"s/'\''/'\''\\\\'\'''\''/g;
-	  s/^\\([_$as_cr_alnum]*_cv_[_$as_cr_alnum]*\\)=\\(.*\\)/\\1='\''\\2'\''/p"
-      ;; #(
-    *)
-      sed -n "/^[_$as_cr_alnum]*_cv_[_$as_cr_alnum]*=/p"
-      ;;
-    esac |
-    sort
-)
-    echo
-
-    $as_echo "## ----------------- ##
-## Output variables. ##
-## ----------------- ##"
-    echo
-    for ac_var in $ac_subst_vars
-    do
-      eval ac_val=\$$ac_var
-      case $ac_val in
-      *\'\''*) ac_val=`$as_echo "$ac_val" | sed "s/'\''/'\''\\\\\\\\'\'''\''/g"`;;
-      esac
-      $as_echo "$ac_var='\''$ac_val'\''"
-    done | sort
-    echo
-
-    if test -n "$ac_subst_files"; then
-      $as_echo "## ------------------- ##
-## File substitutions. ##
-## ------------------- ##"
-      echo
-      for ac_var in $ac_subst_files
-      do
-	eval ac_val=\$$ac_var
-	case $ac_val in
-	*\'\''*) ac_val=`$as_echo "$ac_val" | sed "s/'\''/'\''\\\\\\\\'\'''\''/g"`;;
-	esac
-	$as_echo "$ac_var='\''$ac_val'\''"
-      done | sort
-      echo
-    fi
-
-    if test -s confdefs.h; then
-      $as_echo "## ----------- ##
-## confdefs.h. ##
-## ----------- ##"
-      echo
-      cat confdefs.h
-      echo
-    fi
-    test "$ac_signal" != 0 &&
-      $as_echo "$as_me: caught signal $ac_signal"
-    $as_echo "$as_me: exit $exit_status"
-  } >&5
-  rm -f core *.core core.conftest.* &&
-    rm -f -r conftest* confdefs* conf$$* $ac_clean_files &&
-    exit $exit_status
-' 0
-for ac_signal in 1 2 13 15; do
-  trap 'ac_signal='$ac_signal'; as_fn_exit 1' $ac_signal
-done
-ac_signal=0
-
-# confdefs.h avoids OS command line length limits that DEFS can exceed.
-rm -f -r conftest* confdefs.h
-
-$as_echo "/* confdefs.h */" > confdefs.h
-
-# Predefined preprocessor variables.
-
-cat >>confdefs.h <<_ACEOF
-#define PACKAGE_NAME "$PACKAGE_NAME"
-_ACEOF
-
-cat >>confdefs.h <<_ACEOF
-#define PACKAGE_TARNAME "$PACKAGE_TARNAME"
-_ACEOF
-
-cat >>confdefs.h <<_ACEOF
-#define PACKAGE_VERSION "$PACKAGE_VERSION"
-_ACEOF
-
-cat >>confdefs.h <<_ACEOF
-#define PACKAGE_STRING "$PACKAGE_STRING"
-_ACEOF
-
-cat >>confdefs.h <<_ACEOF
-#define PACKAGE_BUGREPORT "$PACKAGE_BUGREPORT"
-_ACEOF
-
-cat >>confdefs.h <<_ACEOF
-#define PACKAGE_URL "$PACKAGE_URL"
-_ACEOF
-
-
-# Let the site file select an alternate cache file if it wants to.
-# Prefer an explicitly selected file to automatically selected ones.
-ac_site_file1=NONE
-ac_site_file2=NONE
-if test -n "$CONFIG_SITE"; then
-  # 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
-else
-  ac_site_file1=$ac_default_prefix/share/config.site
-  ac_site_file2=$ac_default_prefix/etc/config.site
-fi
-for ac_site_file in "$ac_site_file1" "$ac_site_file2"
-do
-  test "x$ac_site_file" = xNONE && continue
-  if test /dev/null != "$ac_site_file" && test -r "$ac_site_file"; then
-    { $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" \
-      || { { $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
-
-if test -r "$cache_file"; then
-  # Some versions of bash will fail to source /dev/null (special files
-  # actually), so we avoid doing that.  DJGPP emulates it as a regular file.
-  if test /dev/null != "$cache_file" && test -f "$cache_file"; then
-    { $as_echo "$as_me:${as_lineno-$LINENO}: loading cache $cache_file" >&5
-$as_echo "$as_me: loading cache $cache_file" >&6;}
-    case $cache_file in
-      [\\/]* | ?:[\\/]* ) . "$cache_file";;
-      *)                      . "./$cache_file";;
-    esac
-  fi
-else
-  { $as_echo "$as_me:${as_lineno-$LINENO}: creating cache $cache_file" >&5
-$as_echo "$as_me: creating cache $cache_file" >&6;}
-  >$cache_file
-fi
-
-# Check that the precious variables saved in the cache have kept the same
-# value.
-ac_cache_corrupted=false
-for ac_var in $ac_precious_vars; do
-  eval ac_old_set=\$ac_cv_env_${ac_var}_set
-  eval ac_new_set=\$ac_env_${ac_var}_set
-  eval ac_old_val=\$ac_cv_env_${ac_var}_value
-  eval ac_new_val=\$ac_env_${ac_var}_value
-  case $ac_old_set,$ac_new_set in
-    set,)
-      { $as_echo "$as_me:${as_lineno-$LINENO}: error: \`$ac_var' was set to \`$ac_old_val' in the previous run" >&5
-$as_echo "$as_me: error: \`$ac_var' was set to \`$ac_old_val' in the previous run" >&2;}
-      ac_cache_corrupted=: ;;
-    ,set)
-      { $as_echo "$as_me:${as_lineno-$LINENO}: error: \`$ac_var' was not set in the previous run" >&5
-$as_echo "$as_me: error: \`$ac_var' was not set in the previous run" >&2;}
-      ac_cache_corrupted=: ;;
-    ,);;
-    *)
-      if test "x$ac_old_val" != "x$ac_new_val"; then
-	# differences in whitespace do not lead to failure.
-	ac_old_val_w=`echo x $ac_old_val`
-	ac_new_val_w=`echo x $ac_new_val`
-	if test "$ac_old_val_w" != "$ac_new_val_w"; then
-	  { $as_echo "$as_me:${as_lineno-$LINENO}: error: \`$ac_var' has changed since the previous run:" >&5
-$as_echo "$as_me: error: \`$ac_var' has changed since the previous run:" >&2;}
-	  ac_cache_corrupted=:
-	else
-	  { $as_echo "$as_me:${as_lineno-$LINENO}: warning: ignoring whitespace changes in \`$ac_var' since the previous run:" >&5
-$as_echo "$as_me: warning: ignoring whitespace changes in \`$ac_var' since the previous run:" >&2;}
-	  eval $ac_var=\$ac_old_val
-	fi
-	{ $as_echo "$as_me:${as_lineno-$LINENO}:   former value:  \`$ac_old_val'" >&5
-$as_echo "$as_me:   former value:  \`$ac_old_val'" >&2;}
-	{ $as_echo "$as_me:${as_lineno-$LINENO}:   current value: \`$ac_new_val'" >&5
-$as_echo "$as_me:   current value: \`$ac_new_val'" >&2;}
-      fi;;
-  esac
-  # Pass precious variables to config.status.
-  if test "$ac_new_set" = set; then
-    case $ac_new_val in
-    *\'*) ac_arg=$ac_var=`$as_echo "$ac_new_val" | sed "s/'/'\\\\\\\\''/g"` ;;
-    *) ac_arg=$ac_var=$ac_new_val ;;
-    esac
-    case " $ac_configure_args " in
-      *" '$ac_arg' "*) ;; # Avoid dups.  Use of quotes ensures accuracy.
-      *) as_fn_append ac_configure_args " '$ac_arg'" ;;
-    esac
-  fi
-done
-if $ac_cache_corrupted; then
-  { $as_echo "$as_me:${as_lineno-$LINENO}: error: in \`$ac_pwd':" >&5
-$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
-fi
-## -------------------- ##
-## Main body of script. ##
-## -------------------- ##
-
-ac_ext=c
-ac_cpp='$CPP $CPPFLAGS'
-ac_compile='$CC -c $CFLAGS $CPPFLAGS conftest.$ac_ext >&5'
-ac_link='$CC -o conftest$ac_exeext $CFLAGS $CPPFLAGS $LDFLAGS conftest.$ac_ext $LIBS >&5'
-ac_compiler_gnu=$ac_cv_c_compiler_gnu
-
-
-
-ac_config_files="$ac_config_files accelerate-cuda.buildinfo"
-
-
-
-# Check whether --with-compiler was given.
-if test "${with_compiler+set}" = set; then :
-  withval=$with_compiler; GHC=$withval
-else
-  # Extract the first word of "ghc", so it can be a program name with args.
-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 ${ac_cv_path_GHC+:} false; then :
-  $as_echo_n "(cached) " >&6
-else
-  case $GHC in
-  [\\/]* | ?:[\\/]*)
-  ac_cv_path_GHC="$GHC" # Let the user override the test with a path.
-  ;;
-  *)
-  as_save_IFS=$IFS; IFS=$PATH_SEPARATOR
-for as_dir in $PATH
-do
-  IFS=$as_save_IFS
-  test -z "$as_dir" && as_dir=.
-    for ac_exec_ext in '' $ac_executable_extensions; do
-  if as_fn_executable_p "$as_dir/$ac_word$ac_exec_ext"; then
-    ac_cv_path_GHC="$as_dir/$ac_word$ac_exec_ext"
-    $as_echo "$as_me:${as_lineno-$LINENO}: found $as_dir/$ac_word$ac_exec_ext" >&5
-    break 2
-  fi
-done
-  done
-IFS=$as_save_IFS
-
-  ;;
-esac
-fi
-GHC=$ac_cv_path_GHC
-if test -n "$GHC"; then
-  { $as_echo "$as_me:${as_lineno-$LINENO}: result: $GHC" >&5
-$as_echo "$GHC" >&6; }
-else
-  { $as_echo "$as_me:${as_lineno-$LINENO}: result: no" >&5
-$as_echo "no" >&6; }
-fi
-
-
-fi
-
-
-# Check whether --with-nvcc was given.
-if test "${with_nvcc+set}" = set; then :
-  withval=$with_nvcc; NVCC=$withval
-else
-  ac_ext=cpp
-ac_cpp='$CXXCPP $CPPFLAGS'
-ac_compile='$CXX -c $CXXFLAGS $CPPFLAGS conftest.$ac_ext >&5'
-ac_link='$CXX -o conftest$ac_exeext $CXXFLAGS $CPPFLAGS $LDFLAGS conftest.$ac_ext $LIBS >&5'
-ac_compiler_gnu=$ac_cv_cxx_compiler_gnu
-if test -z "$CXX"; then
-  if test -n "$CCC"; then
-    CXX=$CCC
-  else
-    if test -n "$ac_tool_prefix"; then
-  for ac_prog in nvcc
-  do
-    # Extract the first word of "$ac_tool_prefix$ac_prog", so it can be a program name with args.
-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 ${ac_cv_prog_CXX+:} false; then :
-  $as_echo_n "(cached) " >&6
-else
-  if test -n "$CXX"; then
-  ac_cv_prog_CXX="$CXX" # Let the user override the test.
-else
-as_save_IFS=$IFS; IFS=$PATH_SEPARATOR
-for as_dir in $PATH
-do
-  IFS=$as_save_IFS
-  test -z "$as_dir" && as_dir=.
-    for ac_exec_ext in '' $ac_executable_extensions; do
-  if as_fn_executable_p "$as_dir/$ac_word$ac_exec_ext"; then
-    ac_cv_prog_CXX="$ac_tool_prefix$ac_prog"
-    $as_echo "$as_me:${as_lineno-$LINENO}: found $as_dir/$ac_word$ac_exec_ext" >&5
-    break 2
-  fi
-done
-  done
-IFS=$as_save_IFS
-
-fi
-fi
-CXX=$ac_cv_prog_CXX
-if test -n "$CXX"; then
-  { $as_echo "$as_me:${as_lineno-$LINENO}: result: $CXX" >&5
-$as_echo "$CXX" >&6; }
-else
-  { $as_echo "$as_me:${as_lineno-$LINENO}: result: no" >&5
-$as_echo "no" >&6; }
-fi
-
-
-    test -n "$CXX" && break
-  done
-fi
-if test -z "$CXX"; then
-  ac_ct_CXX=$CXX
-  for ac_prog in nvcc
-do
-  # Extract the first word of "$ac_prog", so it can be a program name with args.
-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 ${ac_cv_prog_ac_ct_CXX+:} false; then :
-  $as_echo_n "(cached) " >&6
-else
-  if test -n "$ac_ct_CXX"; then
-  ac_cv_prog_ac_ct_CXX="$ac_ct_CXX" # Let the user override the test.
-else
-as_save_IFS=$IFS; IFS=$PATH_SEPARATOR
-for as_dir in $PATH
-do
-  IFS=$as_save_IFS
-  test -z "$as_dir" && as_dir=.
-    for ac_exec_ext in '' $ac_executable_extensions; do
-  if as_fn_executable_p "$as_dir/$ac_word$ac_exec_ext"; then
-    ac_cv_prog_ac_ct_CXX="$ac_prog"
-    $as_echo "$as_me:${as_lineno-$LINENO}: found $as_dir/$ac_word$ac_exec_ext" >&5
-    break 2
-  fi
-done
-  done
-IFS=$as_save_IFS
-
-fi
-fi
-ac_ct_CXX=$ac_cv_prog_ac_ct_CXX
-if test -n "$ac_ct_CXX"; then
-  { $as_echo "$as_me:${as_lineno-$LINENO}: result: $ac_ct_CXX" >&5
-$as_echo "$ac_ct_CXX" >&6; }
-else
-  { $as_echo "$as_me:${as_lineno-$LINENO}: result: no" >&5
-$as_echo "no" >&6; }
-fi
-
-
-  test -n "$ac_ct_CXX" && break
-done
-
-  if test "x$ac_ct_CXX" = x; then
-    CXX="g++"
-  else
-    case $cross_compiling:$ac_tool_warned in
-yes:)
-{ $as_echo "$as_me:${as_lineno-$LINENO}: WARNING: using cross tools not prefixed with host triplet" >&5
-$as_echo "$as_me: WARNING: using cross tools not prefixed with host triplet" >&2;}
-ac_tool_warned=yes ;;
-esac
-    CXX=$ac_ct_CXX
-  fi
-fi
-
-  fi
-fi
-# Provide some information about the compiler.
-$as_echo "$as_me:${as_lineno-$LINENO}: checking for C++ compiler version" >&5
-set X $ac_compile
-ac_compiler=$2
-for ac_option in --version -v -V -qversion; do
-  { { ac_try="$ac_compiler $ac_option >&5"
-case "(($ac_try" in
-  *\"* | *\`* | *\\*) ac_try_echo=\$ac_try;;
-  *) ac_try_echo=$ac_try;;
-esac
-eval ac_try_echo="\"\$as_me:${as_lineno-$LINENO}: $ac_try_echo\""
-$as_echo "$ac_try_echo"; } >&5
-  (eval "$ac_compiler $ac_option >&5") 2>conftest.err
-  ac_status=$?
-  if test -s conftest.err; then
-    sed '10a\
-... rest of stderr output deleted ...
-         10q' conftest.err >conftest.er1
-    cat conftest.er1 >&5
-  fi
-  rm -f conftest.er1 conftest.err
-  $as_echo "$as_me:${as_lineno-$LINENO}: \$? = $ac_status" >&5
-  test $ac_status = 0; }
-done
-
-cat confdefs.h - <<_ACEOF >conftest.$ac_ext
-/* end confdefs.h.  */
-
-int
-main ()
-{
-
-  ;
-  return 0;
-}
-_ACEOF
-ac_clean_files_save=$ac_clean_files
-ac_clean_files="$ac_clean_files a.out a.out.dSYM a.exe b.out"
-# Try to create an executable without -o first, disregard a.out.
-# It will help us diagnose broken compilers, and finding out an intuition
-# of exeext.
-{ $as_echo "$as_me:${as_lineno-$LINENO}: checking whether the C++ compiler works" >&5
-$as_echo_n "checking whether the C++ compiler works... " >&6; }
-ac_link_default=`$as_echo "$ac_link" | sed 's/ -o *conftest[^ ]*//'`
-
-# The possible output files:
-ac_files="a.out conftest.exe conftest a.exe a_out.exe b.out conftest.*"
-
-ac_rmfiles=
-for ac_file in $ac_files
-do
-  case $ac_file in
-    *.$ac_ext | *.xcoff | *.tds | *.d | *.pdb | *.xSYM | *.bb | *.bbg | *.map | *.inf | *.dSYM | *.o | *.obj ) ;;
-    * ) ac_rmfiles="$ac_rmfiles $ac_file";;
-  esac
-done
-rm -f $ac_rmfiles
-
-if { { ac_try="$ac_link_default"
-case "(($ac_try" in
-  *\"* | *\`* | *\\*) ac_try_echo=\$ac_try;;
-  *) ac_try_echo=$ac_try;;
-esac
-eval ac_try_echo="\"\$as_me:${as_lineno-$LINENO}: $ac_try_echo\""
-$as_echo "$ac_try_echo"; } >&5
-  (eval "$ac_link_default") 2>&5
-  ac_status=$?
-  $as_echo "$as_me:${as_lineno-$LINENO}: \$? = $ac_status" >&5
-  test $ac_status = 0; }; then :
-  # Autoconf-2.13 could set the ac_cv_exeext variable to `no'.
-# So ignore a value of `no', otherwise this would lead to `EXEEXT = no'
-# in a Makefile.  We should not override ac_cv_exeext if it was cached,
-# so that the user can short-circuit this test for compilers unknown to
-# Autoconf.
-for ac_file in $ac_files ''
-do
-  test -f "$ac_file" || continue
-  case $ac_file in
-    *.$ac_ext | *.xcoff | *.tds | *.d | *.pdb | *.xSYM | *.bb | *.bbg | *.map | *.inf | *.dSYM | *.o | *.obj )
-	;;
-    [ab].out )
-	# We found the default executable, but exeext='' is most
-	# certainly right.
-	break;;
-    *.* )
-	if test "${ac_cv_exeext+set}" = set && test "$ac_cv_exeext" != no;
-	then :; else
-	   ac_cv_exeext=`expr "$ac_file" : '[^.]*\(\..*\)'`
-	fi
-	# We set ac_cv_exeext here because the later test for it is not
-	# safe: cross compilers may not add the suffix if given an `-o'
-	# argument, so we may need to know it at that point already.
-	# Even if this section looks crufty: it has the advantage of
-	# actually working.
-	break;;
-    * )
-	break;;
-  esac
-done
-test "$ac_cv_exeext" = no && ac_cv_exeext=
-
-else
-  ac_file=''
-fi
-if test -z "$ac_file"; then :
-  { $as_echo "$as_me:${as_lineno-$LINENO}: result: no" >&5
-$as_echo "no" >&6; }
-$as_echo "$as_me: failed program was:" >&5
-sed 's/^/| /' conftest.$ac_ext >&5
-
-{ { $as_echo "$as_me:${as_lineno-$LINENO}: error: in \`$ac_pwd':" >&5
-$as_echo "$as_me: error: in \`$ac_pwd':" >&2;}
-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; }
-fi
-{ $as_echo "$as_me:${as_lineno-$LINENO}: checking for C++ compiler default output file name" >&5
-$as_echo_n "checking for C++ compiler default output file name... " >&6; }
-{ $as_echo "$as_me:${as_lineno-$LINENO}: result: $ac_file" >&5
-$as_echo "$ac_file" >&6; }
-ac_exeext=$ac_cv_exeext
-
-rm -f -r a.out a.out.dSYM a.exe conftest$ac_cv_exeext b.out
-ac_clean_files=$ac_clean_files_save
-{ $as_echo "$as_me:${as_lineno-$LINENO}: checking for suffix of executables" >&5
-$as_echo_n "checking for suffix of executables... " >&6; }
-if { { ac_try="$ac_link"
-case "(($ac_try" in
-  *\"* | *\`* | *\\*) ac_try_echo=\$ac_try;;
-  *) ac_try_echo=$ac_try;;
-esac
-eval ac_try_echo="\"\$as_me:${as_lineno-$LINENO}: $ac_try_echo\""
-$as_echo "$ac_try_echo"; } >&5
-  (eval "$ac_link") 2>&5
-  ac_status=$?
-  $as_echo "$as_me:${as_lineno-$LINENO}: \$? = $ac_status" >&5
-  test $ac_status = 0; }; then :
-  # If both `conftest.exe' and `conftest' are `present' (well, observable)
-# catch `conftest.exe'.  For instance with Cygwin, `ls conftest' will
-# work properly (i.e., refer to `conftest.exe'), while it won't with
-# `rm'.
-for ac_file in conftest.exe conftest conftest.*; do
-  test -f "$ac_file" || continue
-  case $ac_file in
-    *.$ac_ext | *.xcoff | *.tds | *.d | *.pdb | *.xSYM | *.bb | *.bbg | *.map | *.inf | *.dSYM | *.o | *.obj ) ;;
-    *.* ) ac_cv_exeext=`expr "$ac_file" : '[^.]*\(\..*\)'`
-	  break;;
-    * ) break;;
-  esac
-done
-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; }
-fi
-rm -f conftest conftest$ac_cv_exeext
-{ $as_echo "$as_me:${as_lineno-$LINENO}: result: $ac_cv_exeext" >&5
-$as_echo "$ac_cv_exeext" >&6; }
-
-rm -f conftest.$ac_ext
-EXEEXT=$ac_cv_exeext
-ac_exeext=$EXEEXT
-cat confdefs.h - <<_ACEOF >conftest.$ac_ext
-/* end confdefs.h.  */
-#include <stdio.h>
-int
-main ()
-{
-FILE *f = fopen ("conftest.out", "w");
- return ferror (f) || fclose (f) != 0;
-
-  ;
-  return 0;
-}
-_ACEOF
-ac_clean_files="$ac_clean_files conftest.out"
-# Check that the compiler produces executables we can run.  If not, either
-# the compiler is broken, or we cross compile.
-{ $as_echo "$as_me:${as_lineno-$LINENO}: checking whether we are cross compiling" >&5
-$as_echo_n "checking whether we are cross compiling... " >&6; }
-if test "$cross_compiling" != yes; then
-  { { ac_try="$ac_link"
-case "(($ac_try" in
-  *\"* | *\`* | *\\*) ac_try_echo=\$ac_try;;
-  *) ac_try_echo=$ac_try;;
-esac
-eval ac_try_echo="\"\$as_me:${as_lineno-$LINENO}: $ac_try_echo\""
-$as_echo "$ac_try_echo"; } >&5
-  (eval "$ac_link") 2>&5
-  ac_status=$?
-  $as_echo "$as_me:${as_lineno-$LINENO}: \$? = $ac_status" >&5
-  test $ac_status = 0; }
-  if { ac_try='./conftest$ac_cv_exeext'
-  { { case "(($ac_try" in
-  *\"* | *\`* | *\\*) ac_try_echo=\$ac_try;;
-  *) ac_try_echo=$ac_try;;
-esac
-eval ac_try_echo="\"\$as_me:${as_lineno-$LINENO}: $ac_try_echo\""
-$as_echo "$ac_try_echo"; } >&5
-  (eval "$ac_try") 2>&5
-  ac_status=$?
-  $as_echo "$as_me:${as_lineno-$LINENO}: \$? = $ac_status" >&5
-  test $ac_status = 0; }; }; then
-    cross_compiling=no
-  else
-    if test "$cross_compiling" = maybe; then
-	cross_compiling=yes
-    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.
-If you meant to cross compile, use \`--host'.
-See \`config.log' for more details" "$LINENO" 5; }
-    fi
-  fi
-fi
-{ $as_echo "$as_me:${as_lineno-$LINENO}: result: $cross_compiling" >&5
-$as_echo "$cross_compiling" >&6; }
-
-rm -f conftest.$ac_ext conftest$ac_cv_exeext conftest.out
-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 ${ac_cv_objext+:} false; then :
-  $as_echo_n "(cached) " >&6
-else
-  cat confdefs.h - <<_ACEOF >conftest.$ac_ext
-/* end confdefs.h.  */
-
-int
-main ()
-{
-
-  ;
-  return 0;
-}
-_ACEOF
-rm -f conftest.o conftest.obj
-if { { ac_try="$ac_compile"
-case "(($ac_try" in
-  *\"* | *\`* | *\\*) ac_try_echo=\$ac_try;;
-  *) ac_try_echo=$ac_try;;
-esac
-eval ac_try_echo="\"\$as_me:${as_lineno-$LINENO}: $ac_try_echo\""
-$as_echo "$ac_try_echo"; } >&5
-  (eval "$ac_compile") 2>&5
-  ac_status=$?
-  $as_echo "$as_me:${as_lineno-$LINENO}: \$? = $ac_status" >&5
-  test $ac_status = 0; }; then :
-  for ac_file in conftest.o conftest.obj conftest.*; do
-  test -f "$ac_file" || continue;
-  case $ac_file in
-    *.$ac_ext | *.xcoff | *.tds | *.d | *.pdb | *.xSYM | *.bb | *.bbg | *.map | *.inf | *.dSYM ) ;;
-    *) ac_cv_objext=`expr "$ac_file" : '.*\.\(.*\)'`
-       break;;
-  esac
-done
-else
-  $as_echo "$as_me: failed program was:" >&5
-sed 's/^/| /' conftest.$ac_ext >&5
-
-{ { $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; }
-fi
-rm -f conftest.$ac_cv_objext conftest.$ac_ext
-fi
-{ $as_echo "$as_me:${as_lineno-$LINENO}: result: $ac_cv_objext" >&5
-$as_echo "$ac_cv_objext" >&6; }
-OBJEXT=$ac_cv_objext
-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 ${ac_cv_cxx_compiler_gnu+:} false; then :
-  $as_echo_n "(cached) " >&6
-else
-  cat confdefs.h - <<_ACEOF >conftest.$ac_ext
-/* end confdefs.h.  */
-
-int
-main ()
-{
-#ifndef __GNUC__
-       choke me
-#endif
-
-  ;
-  return 0;
-}
-_ACEOF
-if ac_fn_cxx_try_compile "$LINENO"; then :
-  ac_compiler_gnu=yes
-else
-  ac_compiler_gnu=no
-fi
-rm -f core conftest.err conftest.$ac_objext conftest.$ac_ext
-ac_cv_cxx_compiler_gnu=$ac_compiler_gnu
-
-fi
-{ $as_echo "$as_me:${as_lineno-$LINENO}: result: $ac_cv_cxx_compiler_gnu" >&5
-$as_echo "$ac_cv_cxx_compiler_gnu" >&6; }
-if test $ac_compiler_gnu = yes; then
-  GXX=yes
-else
-  GXX=
-fi
-ac_test_CXXFLAGS=${CXXFLAGS+set}
-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 ${ac_cv_prog_cxx_g+:} false; then :
-  $as_echo_n "(cached) " >&6
-else
-  ac_save_cxx_werror_flag=$ac_cxx_werror_flag
-   ac_cxx_werror_flag=yes
-   ac_cv_prog_cxx_g=no
-   CXXFLAGS="-g"
-   cat confdefs.h - <<_ACEOF >conftest.$ac_ext
-/* end confdefs.h.  */
-
-int
-main ()
-{
-
-  ;
-  return 0;
-}
-_ACEOF
-if ac_fn_cxx_try_compile "$LINENO"; then :
-  ac_cv_prog_cxx_g=yes
-else
-  CXXFLAGS=""
-      cat confdefs.h - <<_ACEOF >conftest.$ac_ext
-/* end confdefs.h.  */
-
-int
-main ()
-{
-
-  ;
-  return 0;
-}
-_ACEOF
-if ac_fn_cxx_try_compile "$LINENO"; then :
-
-else
-  ac_cxx_werror_flag=$ac_save_cxx_werror_flag
-	 CXXFLAGS="-g"
-	 cat confdefs.h - <<_ACEOF >conftest.$ac_ext
-/* end confdefs.h.  */
-
-int
-main ()
-{
-
-  ;
-  return 0;
-}
-_ACEOF
-if ac_fn_cxx_try_compile "$LINENO"; then :
-  ac_cv_prog_cxx_g=yes
-fi
-rm -f core conftest.err conftest.$ac_objext conftest.$ac_ext
-fi
-rm -f core conftest.err conftest.$ac_objext conftest.$ac_ext
-fi
-rm -f core conftest.err conftest.$ac_objext conftest.$ac_ext
-   ac_cxx_werror_flag=$ac_save_cxx_werror_flag
-fi
-{ $as_echo "$as_me:${as_lineno-$LINENO}: result: $ac_cv_prog_cxx_g" >&5
-$as_echo "$ac_cv_prog_cxx_g" >&6; }
-if test "$ac_test_CXXFLAGS" = set; then
-  CXXFLAGS=$ac_save_CXXFLAGS
-elif test $ac_cv_prog_cxx_g = yes; then
-  if test "$GXX" = yes; then
-    CXXFLAGS="-g -O2"
-  else
-    CXXFLAGS="-g"
-  fi
-else
-  if test "$GXX" = yes; then
-    CXXFLAGS="-O2"
-  else
-    CXXFLAGS=
-  fi
-fi
-ac_ext=c
-ac_cpp='$CPP $CPPFLAGS'
-ac_compile='$CC -c $CFLAGS $CPPFLAGS conftest.$ac_ext >&5'
-ac_link='$CC -o conftest$ac_exeext $CFLAGS $CPPFLAGS $LDFLAGS conftest.$ac_ext $LIBS >&5'
-ac_compiler_gnu=$ac_cv_c_compiler_gnu
-
-fi
-
-
-# Check whether --with-gcc was given.
-if test "${with_gcc+set}" = set; then :
-  withval=$with_gcc; CC=$withval
-fi
-
-
-
-
-
-    { $as_echo "$as_me:${as_lineno-$LINENO}: checking size of Int" >&5
-$as_echo_n "checking size of Int... " >&6; }
-
-    sizeof_hs_Int=`$GHC -w -ignore-dot-ghci -e "import Foreign" -e "import Foreign.C" -e "putStr . show $ sizeOf (undefined::Int)"`
-    { $as_echo "$as_me:${as_lineno-$LINENO}: result: $sizeof_hs_Int" >&5
-$as_echo "$sizeof_hs_Int" >&6; }
-
-    case $sizeof_hs_Int in
-        4) type_hs_int=Int32 ;;
-        8) type_hs_int=Int64 ;;
-    esac
-
-    # The type name
-    def="-DHTYPE_INT=$type_hs_int"
-    cpp_flags="$cpp_flags $def"
-    ghc_flags="$ghc_flags -optP$def"
-
-    # And size
-    def="-DSIZEOF_HTYPE_INT=$sizeof_hs_Int"
-    cpp_flags="$cpp_flags $def"
-    ghc_flags="$ghc_flags -optP$def"
-
-
-    { $as_echo "$as_me:${as_lineno-$LINENO}: checking size of Word" >&5
-$as_echo_n "checking size of Word... " >&6; }
-
-    sizeof_hs_Word=`$GHC -w -ignore-dot-ghci -e "import Foreign" -e "import Foreign.C" -e "putStr . show $ sizeOf (undefined::Word)"`
-    { $as_echo "$as_me:${as_lineno-$LINENO}: result: $sizeof_hs_Word" >&5
-$as_echo "$sizeof_hs_Word" >&6; }
-
-    case $sizeof_hs_Word in
-        4) type_hs_word=Word32 ;;
-        8) type_hs_word=Word64 ;;
-    esac
-
-    # The type name
-    def="-DHTYPE_WORD=$type_hs_word"
-    cpp_flags="$cpp_flags $def"
-    ghc_flags="$ghc_flags -optP$def"
-
-    # And size
-    def="-DSIZEOF_HTYPE_WORD=$sizeof_hs_Word"
-    cpp_flags="$cpp_flags $def"
-    ghc_flags="$ghc_flags -optP$def"
-
-
-    { $as_echo "$as_me:${as_lineno-$LINENO}: checking size of CLong" >&5
-$as_echo_n "checking size of CLong... " >&6; }
-
-    sizeof_hs_CLong=`$GHC -w -ignore-dot-ghci -e "import Foreign" -e "import Foreign.C" -e "putStr . show $ sizeOf (undefined::CLong)"`
-    { $as_echo "$as_me:${as_lineno-$LINENO}: result: $sizeof_hs_CLong" >&5
-$as_echo "$sizeof_hs_CLong" >&6; }
-
-    case $sizeof_hs_CLong in
-        4) type_hs_long=Int32 ;;
-        8) type_hs_long=Int64 ;;
-    esac
-
-    # The type name
-    def="-DHTYPE_LONG=$type_hs_long"
-    cpp_flags="$cpp_flags $def"
-    ghc_flags="$ghc_flags -optP$def"
-
-    # And size
-    def="-DSIZEOF_HTYPE_LONG=$sizeof_hs_CLong"
-    cpp_flags="$cpp_flags $def"
-    ghc_flags="$ghc_flags -optP$def"
-
-
-    { $as_echo "$as_me:${as_lineno-$LINENO}: checking size of CULong" >&5
-$as_echo_n "checking size of CULong... " >&6; }
-
-    sizeof_hs_CULong=`$GHC -w -ignore-dot-ghci -e "import Foreign" -e "import Foreign.C" -e "putStr . show $ sizeOf (undefined::CULong)"`
-    { $as_echo "$as_me:${as_lineno-$LINENO}: result: $sizeof_hs_CULong" >&5
-$as_echo "$sizeof_hs_CULong" >&6; }
-
-    case $sizeof_hs_CULong in
-        4) type_hs_unsigned_long=Word32 ;;
-        8) type_hs_unsigned_long=Word64 ;;
-    esac
-
-    # The type name
-    def="-DHTYPE_UNSIGNED_LONG=$type_hs_unsigned_long"
-    cpp_flags="$cpp_flags $def"
-    ghc_flags="$ghc_flags -optP$def"
-
-    # And size
-    def="-DSIZEOF_HTYPE_UNSIGNED_LONG=$sizeof_hs_CULong"
-    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
-# scripts and configure runs, see configure's option --config-cache.
-# It is not useful on other systems.  If it contains results you don't
-# want to keep, you may remove or edit it.
-#
-# config.status only pays attention to the cache file if you give it
-# the --recheck option to rerun configure.
-#
-# `ac_cv_env_foo' variables (set or unset) will be overridden when
-# loading this file, other *unset* `ac_cv_foo' will be assigned the
-# following values.
-
-_ACEOF
-
-# The following way of writing the cache mishandles newlines in values,
-# but we know of no workaround that is simple, portable, and efficient.
-# So, we kill variables containing newlines.
-# Ultrix sh set writes to stderr and can't be redirected directly,
-# and sets the high bit in the cache file unless we assign to the vars.
-(
-  for ac_var in `(set) 2>&1 | sed -n 's/^\([a-zA-Z_][a-zA-Z0-9_]*\)=.*/\1/p'`; do
-    eval ac_val=\$$ac_var
-    case $ac_val in #(
-    *${as_nl}*)
-      case $ac_var in #(
-      *_cv_*) { $as_echo "$as_me:${as_lineno-$LINENO}: WARNING: cache variable $ac_var contains a newline" >&5
-$as_echo "$as_me: WARNING: cache variable $ac_var contains a newline" >&2;} ;;
-      esac
-      case $ac_var in #(
-      _ | IFS | as_nl) ;; #(
-      BASH_ARGV | BASH_SOURCE) eval $ac_var= ;; #(
-      *) { eval $ac_var=; unset $ac_var;} ;;
-      esac ;;
-    esac
-  done
-
-  (set) 2>&1 |
-    case $as_nl`(ac_space=' '; set) 2>&1` in #(
-    *${as_nl}ac_space=\ *)
-      # `set' does not quote correctly, so add quotes: double-quote
-      # substitution turns \\\\ into \\, and sed turns \\ into \.
-      sed -n \
-	"s/'/'\\\\''/g;
-	  s/^\\([_$as_cr_alnum]*_cv_[_$as_cr_alnum]*\\)=\\(.*\\)/\\1='\\2'/p"
-      ;; #(
-    *)
-      # `set' quotes correctly as required by POSIX, so do not add quotes.
-      sed -n "/^[_$as_cr_alnum]*_cv_[_$as_cr_alnum]*=/p"
-      ;;
-    esac |
-    sort
-) |
-  sed '
-     /^ac_cv_env_/b end
-     t clear
-     :clear
-     s/^\([^=]*\)=\(.*[{}].*\)$/test "${\1+set}" = set || &/
-     t end
-     s/^\([^=]*\)=\(.*\)$/\1=${\1=\2}/
-     :end' >>confcache
-if diff "$cache_file" confcache >/dev/null 2>&1; then :; else
-  if test -w "$cache_file"; then
-    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;}
-      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;}
-  fi
-fi
-rm -f confcache
-
-test "x$prefix" = xNONE && prefix=$ac_default_prefix
-# Let make expand exec_prefix.
-test "x$exec_prefix" = xNONE && exec_prefix='${prefix}'
-
-# Transform confdefs.h into DEFS.
-# Protect against shell expansion while executing Makefile rules.
-# Protect against Makefile macro expansion.
-#
-# If the first sed substitution is executed (which looks for macros that
-# take arguments), then branch to the quote section.  Otherwise,
-# look for a macro that doesn't take arguments.
-ac_script='
-:mline
-/\\$/{
- N
- s,\\\n,,
- b mline
-}
-t clear
-:clear
-s/^[	 ]*#[	 ]*define[	 ][	 ]*\([^	 (][^	 (]*([^)]*)\)[	 ]*\(.*\)/-D\1=\2/g
-t quote
-s/^[	 ]*#[	 ]*define[	 ][	 ]*\([^	 ][^	 ]*\)[	 ]*\(.*\)/-D\1=\2/g
-t quote
-b any
-:quote
-s/[	 `~#$^&*(){}\\|;'\''"<>?]/\\&/g
-s/\[/\\&/g
-s/\]/\\&/g
-s/\$/$$/g
-H
-:any
-${
-	g
-	s/^\n//
-	s/\n/ /g
-	p
-}
-'
-DEFS=`sed -n "$ac_script" confdefs.h`
-
-
-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$//'
-  ac_i=`$as_echo "$ac_i" | sed "$ac_script"`
-  # 2. Prepend LIBOBJDIR.  When used with automake>=1.10 LIBOBJDIR
-  #    will be set to the directory where LIBOBJS objects are built.
-  as_fn_append ac_libobjs " \${LIBOBJDIR}$ac_i\$U.$ac_objext"
-  as_fn_append ac_ltlibobjs " \${LIBOBJDIR}$ac_i"'$U.lo'
-done
-LIBOBJS=$ac_libobjs
-
-LTLIBOBJS=$ac_ltlibobjs
-
-
-
-: "${CONFIG_STATUS=./config.status}"
-ac_write_fail=0
-ac_clean_files_save=$ac_clean_files
-ac_clean_files="$ac_clean_files $CONFIG_STATUS"
-{ $as_echo "$as_me:${as_lineno-$LINENO}: creating $CONFIG_STATUS" >&5
-$as_echo "$as_me: creating $CONFIG_STATUS" >&6;}
-as_write_fail=0
-cat >$CONFIG_STATUS <<_ASEOF || as_write_fail=1
-#! $SHELL
-# Generated by $as_me.
-# Run this file to recreate the current configuration.
-# Compiler output produced by configure, useful for debugging
-# configure, is in config.log if it exists.
-
-debug=false
-ac_cs_recheck=false
-ac_cs_silent=false
-
-SHELL=\${CONFIG_SHELL-$SHELL}
-export SHELL
-_ASEOF
-cat >>$CONFIG_STATUS <<\_ASEOF || as_write_fail=1
-## -------------------- ##
-## M4sh Initialization. ##
-## -------------------- ##
-
-# Be more Bourne compatible
-DUALCASE=1; export DUALCASE # for MKS sh
-if test -n "${ZSH_VERSION+set}" && (emulate sh) >/dev/null 2>&1; then :
-  emulate sh
-  NULLCMD=:
-  # Pre-4.2 versions of Zsh do word splitting on ${1+"$@"}, which
-  # is contrary to our usage.  Disable this feature.
-  alias -g '${1+"$@"}'='"$@"'
-  setopt NO_GLOB_SUBST
-else
-  case `(set -o) 2>/dev/null` in #(
-  *posix*) :
-    set -o posix ;; #(
-  *) :
-     ;;
-esac
-fi
-
-
-as_nl='
-'
-export as_nl
-# Printing a long string crashes Solaris 7 /usr/bin/printf.
-as_echo='\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\\'
-as_echo=$as_echo$as_echo$as_echo$as_echo$as_echo
-as_echo=$as_echo$as_echo$as_echo$as_echo$as_echo$as_echo
-# Prefer a ksh shell builtin over an external printf program on Solaris,
-# but without wasting forks for bash or zsh.
-if test -z "$BASH_VERSION$ZSH_VERSION" \
-    && (test "X`print -r -- $as_echo`" = "X$as_echo") 2>/dev/null; then
-  as_echo='print -r --'
-  as_echo_n='print -rn --'
-elif (test "X`printf %s $as_echo`" = "X$as_echo") 2>/dev/null; then
-  as_echo='printf %s\n'
-  as_echo_n='printf %s'
-else
-  if test "X`(/usr/ucb/echo -n -n $as_echo) 2>/dev/null`" = "X-n $as_echo"; then
-    as_echo_body='eval /usr/ucb/echo -n "$1$as_nl"'
-    as_echo_n='/usr/ucb/echo -n'
-  else
-    as_echo_body='eval expr "X$1" : "X\\(.*\\)"'
-    as_echo_n_body='eval
-      arg=$1;
-      case $arg in #(
-      *"$as_nl"*)
-	expr "X$arg" : "X\\(.*\\)$as_nl";
-	arg=`expr "X$arg" : ".*$as_nl\\(.*\\)"`;;
-      esac;
-      expr "X$arg" : "X\\(.*\\)" | tr -d "$as_nl"
-    '
-    export as_echo_n_body
-    as_echo_n='sh -c $as_echo_n_body as_echo'
-  fi
-  export as_echo_body
-  as_echo='sh -c $as_echo_body as_echo'
-fi
-
-# The user is always right.
-if test "${PATH_SEPARATOR+set}" != set; then
-  PATH_SEPARATOR=:
-  (PATH='/bin;/bin'; FPATH=$PATH; sh -c :) >/dev/null 2>&1 && {
-    (PATH='/bin:/bin'; FPATH=$PATH; sh -c :) >/dev/null 2>&1 ||
-      PATH_SEPARATOR=';'
-  }
-fi
-
-
-# IFS
-# We need space, tab and new line, in precisely that order.  Quoting is
-# there to prevent editors from complaining about space-tab.
-# (If _AS_PATH_WALK were called with IFS unset, it would disable word
-# splitting by setting IFS to empty value.)
-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
-for as_dir in $PATH
-do
-  IFS=$as_save_IFS
-  test -z "$as_dir" && as_dir=.
-    test -r "$as_dir/$0" && as_myself=$as_dir/$0 && break
-  done
-IFS=$as_save_IFS
-
-     ;;
-esac
-# We did not find ourselves, most probably we were run as `sh COMMAND'
-# in which case we are not to be found in the path.
-if test "x$as_myself" = x; then
-  as_myself=$0
-fi
-if test ! -f "$as_myself"; then
-  $as_echo "$as_myself: error: cannot find myself; rerun with an absolute file name" >&2
-  exit 1
-fi
-
-# Unset variables that we do not need and which cause bugs (e.g. in
-# pre-3.0 UWIN ksh).  But do not cause bugs in bash 2.01; the "|| exit 1"
-# suppresses any "Segmentation fault" message there.  '((' could
-# trigger a bug in pdksh 5.2.14.
-for as_var in BASH_ENV ENV MAIL MAILPATH
-do eval test x\${$as_var+set} = xset \
-  && ( (unset $as_var) || exit 1) >/dev/null 2>&1 && unset $as_var || :
-done
-PS1='$ '
-PS2='> '
-PS4='+ '
-
-# NLS nuisances.
-LC_ALL=C
-export LC_ALL
-LANGUAGE=C
-export LANGUAGE
-
-# CDPATH.
-(unset CDPATH) >/dev/null 2>&1 && unset CDPATH
-
-
-# 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.
-as_fn_error ()
-{
-  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: $2" >&2
-  as_fn_exit $as_status
-} # as_fn_error
-
-
-# as_fn_set_status STATUS
-# -----------------------
-# Set $? to STATUS, without forking.
-as_fn_set_status ()
-{
-  return $1
-} # as_fn_set_status
-
-# as_fn_exit STATUS
-# -----------------
-# Exit the shell with STATUS, even in a "trap 0" or "set -e" context.
-as_fn_exit ()
-{
-  set +e
-  as_fn_set_status $1
-  exit $1
-} # as_fn_exit
-
-# as_fn_unset VAR
-# ---------------
-# Portably unset VAR.
-as_fn_unset ()
-{
-  { eval $1=; unset $1;}
-}
-as_unset=as_fn_unset
-# as_fn_append VAR VALUE
-# ----------------------
-# Append the text in VALUE to the end of the definition contained in VAR. Take
-# advantage of any shell optimizations that allow amortized linear growth over
-# repeated appends, instead of the typical quadratic growth present in naive
-# implementations.
-if (eval "as_var=1; as_var+=2; test x\$as_var = x12") 2>/dev/null; then :
-  eval 'as_fn_append ()
-  {
-    eval $1+=\$2
-  }'
-else
-  as_fn_append ()
-  {
-    eval $1=\$$1\$2
-  }
-fi # as_fn_append
-
-# as_fn_arith ARG...
-# ------------------
-# Perform arithmetic evaluation on the ARGs, and store the result in the
-# global $as_val. Take advantage of shells that can avoid forks. The arguments
-# must be portable across $(()) and expr.
-if (eval "test \$(( 1 + 1 )) = 2") 2>/dev/null; then :
-  eval 'as_fn_arith ()
-  {
-    as_val=$(( $* ))
-  }'
-else
-  as_fn_arith ()
-  {
-    as_val=`expr "$@" || test $? -eq 1`
-  }
-fi # as_fn_arith
-
-
-if expr a : '\(a\)' >/dev/null 2>&1 &&
-   test "X`expr 00001 : '.*\(...\)'`" = X001; then
-  as_expr=expr
-else
-  as_expr=false
-fi
-
-if (basename -- /) >/dev/null 2>&1 && test "X`basename -- / 2>&1`" = "X/"; then
-  as_basename=basename
-else
-  as_basename=false
-fi
-
-if (as_dir=`dirname -- /` && test "X$as_dir" = X/) >/dev/null 2>&1; then
-  as_dirname=dirname
-else
-  as_dirname=false
-fi
-
-as_me=`$as_basename -- "$0" ||
-$as_expr X/"$0" : '.*/\([^/][^/]*\)/*$' \| \
-	 X"$0" : 'X\(//\)$' \| \
-	 X"$0" : 'X\(/\)' \| . 2>/dev/null ||
-$as_echo X/"$0" |
-    sed '/^.*\/\([^/][^/]*\)\/*$/{
-	    s//\1/
-	    q
-	  }
-	  /^X\/\(\/\/\)$/{
-	    s//\1/
-	    q
-	  }
-	  /^X\/\(\/\).*/{
-	    s//\1/
-	    q
-	  }
-	  s/.*/./; q'`
-
-# Avoid depending upon Character Ranges.
-as_cr_letters='abcdefghijklmnopqrstuvwxyz'
-as_cr_LETTERS='ABCDEFGHIJKLMNOPQRSTUVWXYZ'
-as_cr_Letters=$as_cr_letters$as_cr_LETTERS
-as_cr_digits='0123456789'
-as_cr_alnum=$as_cr_Letters$as_cr_digits
-
-ECHO_C= ECHO_N= ECHO_T=
-case `echo -n x` in #(((((
--n*)
-  case `echo 'xy\c'` in
-  *c*) ECHO_T='	';;	# ECHO_T is single tab character.
-  xy)  ECHO_C='\c';;
-  *)   echo `echo ksh88 bug on AIX 6.1` > /dev/null
-       ECHO_T='	';;
-  esac;;
-*)
-  ECHO_N='-n';;
-esac
-
-rm -f conf$$ conf$$.exe conf$$.file
-if test -d conf$$.dir; then
-  rm -f conf$$.dir/conf$$.file
-else
-  rm -f conf$$.dir
-  mkdir conf$$.dir 2>/dev/null
-fi
-if (echo >conf$$.file) 2>/dev/null; then
-  if ln -s conf$$.file conf$$ 2>/dev/null; then
-    as_ln_s='ln -s'
-    # ... but there are two gotchas:
-    # 1) On MSYS, both `ln -s file dir' and `ln file dir' fail.
-    # 2) DJGPP < 2.04 has no symlinks; `ln -s' creates a wrapper executable.
-    # In both cases, we have to default to `cp -pR'.
-    ln -s conf$$.file conf$$.dir 2>/dev/null && test ! -f conf$$.exe ||
-      as_ln_s='cp -pR'
-  elif ln conf$$.file conf$$ 2>/dev/null; then
-    as_ln_s=ln
-  else
-    as_ln_s='cp -pR'
-  fi
-else
-  as_ln_s='cp -pR'
-fi
-rm -f conf$$ conf$$.exe conf$$.dir/conf$$.file conf$$.file
-rmdir conf$$.dir 2>/dev/null
-
-
-# as_fn_mkdir_p
-# -------------
-# Create "$as_dir" as a directory, including parents if necessary.
-as_fn_mkdir_p ()
-{
-
-  case $as_dir in #(
-  -*) as_dir=./$as_dir;;
-  esac
-  test -d "$as_dir" || eval $as_mkdir_p || {
-    as_dirs=
-    while :; do
-      case $as_dir in #(
-      *\'*) as_qdir=`$as_echo "$as_dir" | sed "s/'/'\\\\\\\\''/g"`;; #'(
-      *) as_qdir=$as_dir;;
-      esac
-      as_dirs="'$as_qdir' $as_dirs"
-      as_dir=`$as_dirname -- "$as_dir" ||
-$as_expr X"$as_dir" : 'X\(.*[^/]\)//*[^/][^/]*/*$' \| \
-	 X"$as_dir" : 'X\(//\)[^/]' \| \
-	 X"$as_dir" : 'X\(//\)$' \| \
-	 X"$as_dir" : 'X\(/\)' \| . 2>/dev/null ||
-$as_echo X"$as_dir" |
-    sed '/^X\(.*[^/]\)\/\/*[^/][^/]*\/*$/{
-	    s//\1/
-	    q
-	  }
-	  /^X\(\/\/\)[^/].*/{
-	    s//\1/
-	    q
-	  }
-	  /^X\(\/\/\)$/{
-	    s//\1/
-	    q
-	  }
-	  /^X\(\/\).*/{
-	    s//\1/
-	    q
-	  }
-	  s/.*/./; q'`
-      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"
-
-
-} # as_fn_mkdir_p
-if mkdir -p . 2>/dev/null; then
-  as_mkdir_p='mkdir -p "$as_dir"'
-else
-  test -d ./-p && rmdir ./-p
-  as_mkdir_p=false
-fi
-
-
-# as_fn_executable_p FILE
-# -----------------------
-# Test if FILE is an executable regular file.
-as_fn_executable_p ()
-{
-  test -f "$1" && test -x "$1"
-} # as_fn_executable_p
-as_test_x='test -x'
-as_executable_p=as_fn_executable_p
-
-# Sed expression to map a string onto a valid CPP name.
-as_tr_cpp="eval sed 'y%*$as_cr_letters%P$as_cr_LETTERS%;s%[^_$as_cr_alnum]%_%g'"
-
-# Sed expression to map a string onto a valid variable name.
-as_tr_sh="eval sed 'y%*+%pp%;s%[^_$as_cr_alnum]%_%g'"
-
-
-exec 6>&1
-## ----------------------------------- ##
-## Main body of $CONFIG_STATUS script. ##
-## ----------------------------------- ##
-_ASEOF
-test $as_write_fail = 0 && chmod +x $CONFIG_STATUS || ac_write_fail=1
-
-cat >>$CONFIG_STATUS <<\_ACEOF || ac_write_fail=1
-# Save the log message, to keep $0 and so on meaningful, and to
-# 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.14.0.0, which was
-generated by GNU Autoconf 2.69.  Invocation command line was
-
-  CONFIG_FILES    = $CONFIG_FILES
-  CONFIG_HEADERS  = $CONFIG_HEADERS
-  CONFIG_LINKS    = $CONFIG_LINKS
-  CONFIG_COMMANDS = $CONFIG_COMMANDS
-  $ $0 $@
-
-on `(hostname || uname -n) 2>/dev/null | sed 1q`
-"
-
-_ACEOF
-
-case $ac_config_files in *"
-"*) set x $ac_config_files; shift; ac_config_files=$*;;
-esac
-
-
-
-cat >>$CONFIG_STATUS <<_ACEOF || ac_write_fail=1
-# Files that config.status was made for.
-config_files="$ac_config_files"
-
-_ACEOF
-
-cat >>$CONFIG_STATUS <<\_ACEOF || ac_write_fail=1
-ac_cs_usage="\
-\`$as_me' instantiates files and other configuration actions
-from templates according to the current configuration.  Unless the files
-and actions are specified as TAGs, all are instantiated by default.
-
-Usage: $0 [OPTION]... [TAG]...
-
-  -h, --help       print this help, then exit
-  -V, --version    print version number and configuration settings, then exit
-      --config     print configuration, then exit
-  -q, --quiet, --silent
-                   do not print progress messages
-  -d, --debug      don't remove temporary files
-      --recheck    update $as_me by reconfiguring in the same conditions
-      --file=FILE[:TEMPLATE]
-                   instantiate the configuration file FILE
-
-Configuration files:
-$config_files
-
-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.14.0.0
-configured by $0, generated by GNU Autoconf 2.69,
-  with options \\"\$ac_cs_config\\"
-
-Copyright (C) 2012 Free Software Foundation, Inc.
-This config.status script is free software; the Free Software Foundation
-gives unlimited permission to copy, distribute and modify it."
-
-ac_pwd='$ac_pwd'
-srcdir='$srcdir'
-test -n "\$AWK" || AWK=awk
-_ACEOF
-
-cat >>$CONFIG_STATUS <<\_ACEOF || ac_write_fail=1
-# The default lists apply if the user does not specify any file.
-ac_need_defaults=:
-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
-    ac_shift=shift
-    ;;
-  esac
-
-  case $ac_option in
-  # Handling of the options.
-  -recheck | --recheck | --rechec | --reche | --rech | --rec | --re | --r)
-    ac_cs_recheck=: ;;
-  --version | --versio | --versi | --vers | --ver | --ve | --v | -V )
-    $as_echo "$ac_cs_version"; exit ;;
-  --config | --confi | --conf | --con | --co | --c )
-    $as_echo "$ac_cs_config"; exit ;;
-  --debug | --debu | --deb | --de | --d | -d )
-    debug=: ;;
-  --file | --fil | --fi | --f )
-    $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;;
-  --he | --h |  --help | --hel | -h )
-    $as_echo "$ac_cs_usage"; exit ;;
-  -q | -quiet | --quiet | --quie | --qui | --qu | --q \
-  | -silent | --silent | --silen | --sile | --sil | --si | --s)
-    ac_cs_silent=: ;;
-
-  # This is an error.
-  -*) as_fn_error $? "unrecognized option: \`$1'
-Try \`$0 --help' for more information." ;;
-
-  *) as_fn_append ac_config_targets " $1"
-     ac_need_defaults=false ;;
-
-  esac
-  shift
-done
-
-ac_configure_extra_args=
-
-if $ac_cs_silent; then
-  exec 6>/dev/null
-  ac_configure_extra_args="$ac_configure_extra_args --silent"
-fi
-
-_ACEOF
-cat >>$CONFIG_STATUS <<_ACEOF || ac_write_fail=1
-if \$ac_cs_recheck; then
-  set X $SHELL '$0' $ac_configure_args \$ac_configure_extra_args --no-create --no-recursion
-  shift
-  \$as_echo "running CONFIG_SHELL=$SHELL \$*" >&6
-  CONFIG_SHELL='$SHELL'
-  export CONFIG_SHELL
-  exec "\$@"
-fi
-
-_ACEOF
-cat >>$CONFIG_STATUS <<\_ACEOF || ac_write_fail=1
-exec 5>>config.log
-{
-  echo
-  sed 'h;s/./-/g;s/^.../## /;s/...$/ ##/;p;x;p;x' <<_ASBOX
-## Running $as_me. ##
-_ASBOX
-  $as_echo "$ac_log"
-} >&5
-
-_ACEOF
-cat >>$CONFIG_STATUS <<_ACEOF || ac_write_fail=1
-_ACEOF
-
-cat >>$CONFIG_STATUS <<\_ACEOF || ac_write_fail=1
-
-# Handling of arguments.
-for ac_config_target in $ac_config_targets
-do
-  case $ac_config_target in
-    "accelerate-cuda.buildinfo") CONFIG_FILES="$CONFIG_FILES accelerate-cuda.buildinfo" ;;
-
-  *) as_fn_error $? "invalid argument: \`$ac_config_target'" "$LINENO" 5;;
-  esac
-done
-
-
-# If the user did not use the arguments to specify the items to instantiate,
-# then the envvar interface is used.  Set only those that are not.
-# We use the long form for the default assignment because of an extremely
-# bizarre bug on SunOS 4.1.3.
-if $ac_need_defaults; then
-  test "${CONFIG_FILES+set}" = set || CONFIG_FILES=$config_files
-fi
-
-# Have a temporary directory for convenience.  Make it in the build tree
-# simply because there is no reason against having it here, and in addition,
-# creating and moving files from /tmp can sometimes cause problems.
-# Hook for its removal unless debugging.
-# Note that there is a small window in which the directory will not be cleaned:
-# after its creation but before its name has been assigned to `$tmp'.
-$debug ||
-{
-  tmp= ac_tmp=
-  trap '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
-}
-# Create a (secure) tmp directory for tmp files.
-
-{
-  tmp=`(umask 077 && mktemp -d "./confXXXXXX") 2>/dev/null` &&
-  test -d "$tmp"
-}  ||
-{
-  tmp=./conf$$-$RANDOM
-  (umask 077 && mkdir "$tmp")
-} || 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.
-# This happens for instance with `./config.status config.h'.
-if test -n "$CONFIG_FILES"; then
-
-
-ac_cr=`echo X | tr X '\015'`
-# On cygwin, bash can eat \r inside `` if the user requested igncr.
-# But we know of no other shell where ac_cr would be empty at this
-# point, so we can use a bashism as a fallback.
-if test "x$ac_cr" = x; then
-  eval ac_cr=\$\'\\r\'
-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'
-else
-  ac_cs_awk_cr=$ac_cr
-fi
-
-echo 'BEGIN {' >"$ac_tmp/subs1.awk" &&
-_ACEOF
-
-
-{
-  echo "cat >conf$$subs.awk <<_ACEOF" &&
-  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 '^'`
-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
-
-  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
-  else
-    ac_delim="$ac_delim!$ac_delim _$ac_delim!! "
-  fi
-done
-rm -f conf$$subs.sh
-
-cat >>$CONFIG_STATUS <<_ACEOF || ac_write_fail=1
-cat >>"\$ac_tmp/subs1.awk" <<\\_ACAWK &&
-_ACEOF
-sed -n '
-h
-s/^/S["/; s/!.*/"]=/
-p
-g
-s/^[^!]*!//
-:repl
-t repl
-s/'"$ac_delim"'$//
-t delim
-:nl
-h
-s/\(.\{148\}\)..*/\1/
-t more1
-s/["\\]/\\&/g; s/^/"/; s/$/\\n"\\/
-p
-n
-b repl
-:more1
-s/["\\]/\\&/g; s/^/"/; s/$/"\\/
-p
-g
-s/.\{148\}//
-t nl
-:delim
-h
-s/\(.\{148\}\)..*/\1/
-t more2
-s/["\\]/\\&/g; s/^/"/; s/$/"/
-p
-b
-:more2
-s/["\\]/\\&/g; s/^/"/; s/$/"\\/
-p
-g
-s/.\{148\}//
-t delim
-' <conf$$subs.awk | sed '
-/^[^""]/{
-  N
-  s/\n//
-}
-' >>$CONFIG_STATUS || ac_write_fail=1
-rm -f conf$$subs.awk
-cat >>$CONFIG_STATUS <<_ACEOF || ac_write_fail=1
-_ACAWK
-cat >>"\$ac_tmp/subs1.awk" <<_ACAWK &&
-  for (key in S) S_is_set[key] = 1
-  FS = ""
-
-}
-{
-  line = $ 0
-  nfields = split(line, field, "@")
-  substed = 0
-  len = length(field[1])
-  for (i = 2; i < nfields; i++) {
-    key = field[i]
-    keylen = length(key)
-    if (S_is_set[key]) {
-      value = S[key]
-      line = substr(line, 1, len) "" value "" substr(line, len + keylen + 3)
-      len += length(value) + length(field[++i])
-      substed = 1
-    } else
-      len += 1 + keylen
-  }
-
-  print line
-}
-
-_ACAWK
-_ACEOF
-cat >>$CONFIG_STATUS <<\_ACEOF || ac_write_fail=1
-if sed "s/$ac_cr//" < /dev/null > /dev/null 2>&1; then
-  sed "s/$ac_cr\$//; s/$ac_cr/$ac_cs_awk_cr/g"
-else
-  cat
-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 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[	 ]*=[	 ]*/{
-h
-s///
-s/^/:/
-s/[	 ]*$/:/
-s/:\$(srcdir):/:/g
-s/:\${srcdir}:/:/g
-s/:@srcdir@:/:/g
-s/^:*//
-s/:*$//
-x
-s/\(=[	 ]*\).*/\1/
-G
-s/\n//
-s/^[^=]*=[	 ]*$//
-}'
-fi
-
-cat >>$CONFIG_STATUS <<\_ACEOF || ac_write_fail=1
-fi # test -n "$CONFIG_FILES"
-
-
-eval set X "  :F $CONFIG_FILES      "
-shift
-for ac_tag
-do
-  case $ac_tag in
-  :[FHLC]) ac_mode=$ac_tag; continue;;
-  esac
-  case $ac_mode$ac_tag in
-  :[FHL]*:*);;
-  :L* | :C*:*) as_fn_error $? "invalid tag \`$ac_tag'" "$LINENO" 5;;
-  :[FH]-) ac_tag=-:-;;
-  :[FH]*) ac_tag=$ac_tag:$ac_tag.in;;
-  esac
-  ac_save_IFS=$IFS
-  IFS=:
-  set x $ac_tag
-  IFS=$ac_save_IFS
-  shift
-  ac_file=$1
-  shift
-
-  case $ac_mode in
-  :L) ac_source=$1;;
-  :[FH])
-    ac_file_inputs=
-    for ac_f
-    do
-      case $ac_f in
-      -) 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 `:'.
-	 test -f "$ac_f" ||
-	   case $ac_f in
-	   [\\/$]*) false;;
-	   *) test -f "$srcdir/$ac_f" && ac_f="$srcdir/$ac_f";;
-	   esac ||
-	   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'"
-    done
-
-    # Let's still pretend it is `configure' which instantiates (i.e., don't
-    # use $as_me), people would be surprised to read:
-    #    /* config.h.  Generated by config.status.  */
-    configure_input='Generated from '`
-	  $as_echo "$*" | sed 's|^[^:]*/||;s|:[^:]*/|, |g'
-	`' by configure.'
-    if test x"$ac_file" != x-; then
-      configure_input="$ac_file.  $configure_input"
-      { $as_echo "$as_me:${as_lineno-$LINENO}: creating $ac_file" >&5
-$as_echo "$as_me: creating $ac_file" >&6;}
-    fi
-    # Neutralize special characters interpreted by sed in replacement strings.
-    case $configure_input in #(
-    *\&* | *\|* | *\\* )
-       ac_sed_conf_input=`$as_echo "$configure_input" |
-       sed 's/[\\\\&|]/\\\\&/g'`;; #(
-    *) ac_sed_conf_input=$configure_input;;
-    esac
-
-    case $ac_tag in
-    *:-:* | *:-) cat >"$ac_tmp/stdin" \
-      || as_fn_error $? "could not create $ac_file" "$LINENO" 5 ;;
-    esac
-    ;;
-  esac
-
-  ac_dir=`$as_dirname -- "$ac_file" ||
-$as_expr X"$ac_file" : 'X\(.*[^/]\)//*[^/][^/]*/*$' \| \
-	 X"$ac_file" : 'X\(//\)[^/]' \| \
-	 X"$ac_file" : 'X\(//\)$' \| \
-	 X"$ac_file" : 'X\(/\)' \| . 2>/dev/null ||
-$as_echo X"$ac_file" |
-    sed '/^X\(.*[^/]\)\/\/*[^/][^/]*\/*$/{
-	    s//\1/
-	    q
-	  }
-	  /^X\(\/\/\)[^/].*/{
-	    s//\1/
-	    q
-	  }
-	  /^X\(\/\/\)$/{
-	    s//\1/
-	    q
-	  }
-	  /^X\(\/\).*/{
-	    s//\1/
-	    q
-	  }
-	  s/.*/./; q'`
-  as_dir="$ac_dir"; as_fn_mkdir_p
-  ac_builddir=.
-
-case "$ac_dir" in
-.) ac_dir_suffix= ac_top_builddir_sub=. ac_top_build_prefix= ;;
-*)
-  ac_dir_suffix=/`$as_echo "$ac_dir" | sed 's|^\.[\\/]||'`
-  # A ".." for each directory in $ac_dir_suffix.
-  ac_top_builddir_sub=`$as_echo "$ac_dir_suffix" | sed 's|/[^\\/]*|/..|g;s|/||'`
-  case $ac_top_builddir_sub in
-  "") ac_top_builddir_sub=. ac_top_build_prefix= ;;
-  *)  ac_top_build_prefix=$ac_top_builddir_sub/ ;;
-  esac ;;
-esac
-ac_abs_top_builddir=$ac_pwd
-ac_abs_builddir=$ac_pwd$ac_dir_suffix
-# for backward compatibility:
-ac_top_builddir=$ac_top_build_prefix
-
-case $srcdir in
-  .)  # We are building in place.
-    ac_srcdir=.
-    ac_top_srcdir=$ac_top_builddir_sub
-    ac_abs_top_srcdir=$ac_pwd ;;
-  [\\/]* | ?:[\\/]* )  # Absolute name.
-    ac_srcdir=$srcdir$ac_dir_suffix;
-    ac_top_srcdir=$srcdir
-    ac_abs_top_srcdir=$srcdir ;;
-  *) # Relative name.
-    ac_srcdir=$ac_top_build_prefix$srcdir$ac_dir_suffix
-    ac_top_srcdir=$ac_top_build_prefix$srcdir
-    ac_abs_top_srcdir=$ac_pwd/$srcdir ;;
-esac
-ac_abs_srcdir=$ac_abs_top_srcdir$ac_dir_suffix
-
-
-  case $ac_mode in
-  :F)
-  #
-  # CONFIG_FILE
-  #
-
-_ACEOF
-
-cat >>$CONFIG_STATUS <<\_ACEOF || ac_write_fail=1
-# If the template does not know about datarootdir, expand it.
-# FIXME: This hack should be removed a few years after 2.60.
-ac_datarootdir_hack=; ac_datarootdir_seen=
-ac_sed_dataroot='
-/datarootdir/ {
-  p
-  q
-}
-/@datadir@/p
-/@docdir@/p
-/@infodir@/p
-/@localedir@/p
-/@mandir@/p'
-case `eval "sed -n \"\$ac_sed_dataroot\" $ac_file_inputs"` in
-*datarootdir*) ac_datarootdir_seen=yes;;
-*@datadir@*|*@docdir@*|*@infodir@*|*@localedir@*|*@mandir@*)
-  { $as_echo "$as_me:${as_lineno-$LINENO}: WARNING: $ac_file_inputs seems to ignore the --datarootdir setting" >&5
-$as_echo "$as_me: WARNING: $ac_file_inputs seems to ignore the --datarootdir setting" >&2;}
-_ACEOF
-cat >>$CONFIG_STATUS <<_ACEOF || ac_write_fail=1
-  ac_datarootdir_hack='
-  s&@datadir@&$datadir&g
-  s&@docdir@&$docdir&g
-  s&@infodir@&$infodir&g
-  s&@localedir@&$localedir&g
-  s&@mandir@&$mandir&g
-  s&\\\${datarootdir}&$datarootdir&g' ;;
-esac
-_ACEOF
-
-# Neutralize VPATH when `$srcdir' = `.'.
-# Shell code in configure.ac might set extrasub.
-# FIXME: do we really want to maintain this feature?
-cat >>$CONFIG_STATUS <<_ACEOF || ac_write_fail=1
-ac_sed_extra="$ac_vpsub
-$extrasub
-_ACEOF
-cat >>$CONFIG_STATUS <<\_ACEOF || ac_write_fail=1
-:t
-/@[a-zA-Z_][a-zA-Z_0-9]*@/!b
-s|@configure_input@|$ac_sed_conf_input|;t t
-s&@top_builddir@&$ac_top_builddir_sub&;t t
-s&@top_build_prefix@&$ac_top_build_prefix&;t t
-s&@srcdir@&$ac_srcdir&;t t
-s&@abs_srcdir@&$ac_abs_srcdir&;t t
-s&@top_srcdir@&$ac_top_srcdir&;t t
-s&@abs_top_srcdir@&$ac_abs_top_srcdir&;t t
-s&@builddir@&$ac_builddir&;t t
-s&@abs_builddir@&$ac_abs_builddir&;t t
-s&@abs_top_builddir@&$ac_abs_top_builddir&;t t
-$ac_datarootdir_hack
-"
-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' "$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
-$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;}
-
-  rm -f "$ac_tmp/stdin"
-  case $ac_file in
-  -) 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
- ;;
-
-
-
-  esac
-
-done # for ac_tag
-
-
-as_fn_exit 0
-_ACEOF
-ac_clean_files=$ac_clean_files_save
-
-test $ac_write_fail = 0 ||
-  as_fn_error $? "write failure creating $CONFIG_STATUS" "$LINENO" 5
-
-
-# configure is writing to config.log, and then calls config.status.
-# config.status does its own redirection, appending to config.log.
-# Unfortunately, on DOS this fails, as config.log is still kept open
-# by configure, so config.status won't be able to write to it; its
-# output is simply discarded.  So we exec the FD to /dev/null,
-# effectively closing config.log, so it can be properly (re)opened and
-# appended to by config.status.  When coming back to configure, we
-# need to make the FD available again.
-if test "$no_create" != yes; then
-  ac_cs_success=:
-  ac_config_status_args=
-  test "$silent" = yes &&
-    ac_config_status_args="$ac_config_status_args --quiet"
-  exec 5>/dev/null
-  $SHELL $CONFIG_STATUS $ac_config_status_args || ac_cs_success=false
-  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 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
-$as_echo "$as_me: WARNING: unrecognized options: $ac_unrecognized_opts" >&2;}
-fi
-
-
diff --git a/cubits/accelerate_cuda.h b/cubits/accelerate_cuda.h
--- a/cubits/accelerate_cuda.h
+++ b/cubits/accelerate_cuda.h
@@ -1,8 +1,8 @@
 /* -----------------------------------------------------------------------------
  *
  * Module      : Accelerate-CUDA
- * Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
- *               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+ * Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+ *               [2009..2014] Trevor L. McDonell
  * License     : BSD3
  *
  * Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
diff --git a/cubits/accelerate_cuda_assert.h b/cubits/accelerate_cuda_assert.h
--- a/cubits/accelerate_cuda_assert.h
+++ b/cubits/accelerate_cuda_assert.h
@@ -1,8 +1,8 @@
 /* -----------------------------------------------------------------------------
  *
  * Module      : Assert
- * Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
- *               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+ * Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+ *               [2009..2014] Trevor L. McDonell
  * License     : BSD3
  *
  * Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
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
@@ -1,8 +1,8 @@
 /* -----------------------------------------------------------------------------
  *
  * Module      : Function
- * Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
- *               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+ * Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+ *               [2009..2014] Trevor L. McDonell
  * License     : BSD3
  *
  * Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
@@ -49,25 +49,25 @@
 }
 
 template <>
-static __inline__ __device__ Word8 idiv(const Word8 x, const Word8 y)
+__inline__ __device__ Word8 idiv(const Word8 x, const Word8 y)
 {
     return x / y;
 }
 
 template <>
-static __inline__ __device__ Word16 idiv(const Word16 x, const Word16 y)
+__inline__ __device__ Word16 idiv(const Word16 x, const Word16 y)
 {
     return x / y;
 }
 
 template <>
-static __inline__ __device__ Word32 idiv(const Word32 x, const Word32 y)
+__inline__ __device__ Word32 idiv(const Word32 x, const Word32 y)
 {
     return x / y;
 }
 
 template <>
-static __inline__ __device__ Word64 idiv(const Word64 x, const Word64 y)
+__inline__ __device__ Word64 idiv(const Word64 x, const Word64 y)
 {
     return x / y;
 }
@@ -84,25 +84,25 @@
 }
 
 template <>
-static __inline__ __device__ Word8 mod(const Word8 x, const Word8 y)
+__inline__ __device__ Word8 mod(const Word8 x, const Word8 y)
 {
     return x % y;
 }
 
 template <>
-static __inline__ __device__ Word16 mod(const Word16 x, const Word16 y)
+__inline__ __device__ Word16 mod(const Word16 x, const Word16 y)
 {
     return x % y;
 }
 
 template <>
-static __inline__ __device__ Word32 mod(const Word32 x, const Word32 y)
+__inline__ __device__ Word32 mod(const Word32 x, const Word32 y)
 {
     return x % y;
 }
 
 template <>
-static __inline__ __device__ Word64 mod(const Word64 x, const Word64 y)
+__inline__ __device__ Word64 mod(const Word64 x, const Word64 y)
 {
     return x % y;
 }
@@ -122,13 +122,13 @@
 }
 
 template <>
-static __inline__ __device__ Word32 reinterpret32(const Word32 x)
+__inline__ __device__ Word32 reinterpret32(const Word32 x)
 {
     return x;
 }
 
 template <>
-static __inline__ __device__ Word32 reinterpret32(const float x)
+__inline__ __device__ Word32 reinterpret32(const float x)
 {
     return __float_as_int(x);
 }
@@ -143,14 +143,14 @@
 }
 
 template <>
-static __inline__ __device__ Word64 reinterpret64(const Word64 x)
+__inline__ __device__ Word64 reinterpret64(const Word64 x)
 {
     return x;
 }
 
 #if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 130
 template <>
-static __inline__ __device__ Word64 reinterpret64(const double x)
+__inline__ __device__ Word64 reinterpret64(const double x)
 {
     return __double_as_longlong(x);
 }
@@ -171,13 +171,13 @@
 
 #if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 110
 template <>
-static __inline__ __device__ Int32 atomicCAS32(Int32* address, Int32 compare, Int32 val)
+__inline__ __device__ Int32 atomicCAS32(Int32* address, Int32 compare, Int32 val)
 {
     return atomicCAS(address, compare, val);
 }
 
 template <>
-static __inline__ __device__ Word32 atomicCAS32(Word32* address, Word32 compare, Word32 val)
+__inline__ __device__ Word32 atomicCAS32(Word32* address, Word32 compare, Word32 val)
 {
     return atomicCAS(address, compare, val);
 }
@@ -194,7 +194,7 @@
 }
 
 template <>
-static __inline__ __device__ Word64 atomicCAS64(Word64* address, Word64 compare, Word64 val)
+__inline__ __device__ Word64 atomicCAS64(Word64* address, Word64 compare, Word64 val)
 {
     return atomicCAS(address, compare, val);
 }
@@ -217,13 +217,13 @@
 }
 
 template <>
-static __inline__ __device__ int shfl_up32(int var, unsigned int delta, int width)
+__inline__ __device__ int shfl_up32(int var, unsigned int delta, int width)
 {
     return __shfl_up(var, delta, width);
 }
 
 template <>
-static __inline__ __device__ float shfl_up32(float var, unsigned int delta, int width)
+__inline__ __device__ float shfl_up32(float var, unsigned int delta, int width)
 {
     return __shfl_up(var, delta, width);
 }
@@ -254,13 +254,13 @@
 }
 
 template <>
-static __inline__ __device__ int shfl_xor32(int var, int laneMask, int width)
+__inline__ __device__ int shfl_xor32(int var, int laneMask, int width)
 {
     return __shfl_xor(var, laneMask, width);
 }
 
 template <>
-static __inline__ __device__ float shfl_xor32(float var, int laneMask, int width)
+__inline__ __device__ float shfl_xor32(float var, int laneMask, int width)
 {
     return __shfl_xor(var, laneMask, width);
 }
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
@@ -1,8 +1,8 @@
 /* -----------------------------------------------------------------------------
  *
  * Module      : Texture
- * Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
- *               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+ * Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+ *               [2009..2014] Trevor L. McDonell
  * License     : BSD3
  *
  * Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
diff --git a/cubits/accelerate_cuda_type.h b/cubits/accelerate_cuda_type.h
--- a/cubits/accelerate_cuda_type.h
+++ b/cubits/accelerate_cuda_type.h
@@ -1,8 +1,8 @@
 /* -----------------------------------------------------------------------------
  *
  * Module      : Type
- * Copyright   : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee
- *               [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell
+ * Copyright   : [2008..2014] Manuel M T Chakravarty, Gabriele Keller
+ *               [2009..2014] Trevor L. McDonell
  * License     : BSD3
  *
  * Maintainer  : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>
diff --git a/include/accelerate.h b/include/accelerate.h
deleted file mode 100644
--- a/include/accelerate.h
+++ /dev/null
@@ -1,31 +0,0 @@
-
-#ifndef ACCELERATE_H
-#define ACCELERATE_H
-import qualified Data.Array.Accelerate.Internal.Check as Ck
-
-
-/*
- * Internal checks
- */
-#define ERROR(f)  (Ck.f __FILE__ __LINE__)
-#define ASSERT (Ck.assert __FILE__ __LINE__)
-#define ENSURE (Ck.f __FILE__ __LINE__)
-#define CHECK(f) (Ck.f __FILE__ __LINE__)
-
-#define BOUNDS_ERROR(f) (ERROR(f) Ck.Bounds)
-#define BOUNDS_ASSERT (ASSERT Ck.Bounds)
-#define BOUNDS_ENSURE (ENSURE Ck.Bounds)
-#define BOUNDS_CHECK(f) (CHECK(f) Ck.Bounds)
-
-#define UNSAFE_ERROR(f) (ERROR(f) Ck.Unsafe)
-#define UNSAFE_ASSERT (ASSERT Ck.Unsafe)
-#define UNSAFE_ENSURE (ENSURE Ck.Unsafe)
-#define UNSAFE_CHECK(f) (CHECK(f) Ck.Unsafe)
-
-#define INTERNAL_ERROR(f) (ERROR(f) Ck.Internal)
-#define INTERNAL_ASSERT (ASSERT Ck.Internal)
-#define INTERNAL_ENSURE (ENSURE Ck.Internal)
-#define INTERNAL_CHECK(f) (CHECK(f) Ck.Internal)
-
-#endif
-