accelerate-cuda 0.14.0.0 → 0.15.0.0
raw patch · 43 files changed
+731/−4264 lines, 43 filesdep ~SafeSemaphoredep ~acceleratedep ~basesetup-changed
Dependency ranges changed: SafeSemaphore, accelerate, base, binary, bytestring, cryptohash, cuda, fclabels, hashable, hashtables, language-c-quote, mainland-pretty, mtl, srcloc, text, transformers, unordered-containers
Files
- Data/Array/Accelerate/CUDA.hs +4/−5
- Data/Array/Accelerate/CUDA/AST.hs +2/−2
- Data/Array/Accelerate/CUDA/Analysis/Device.hs +2/−2
- Data/Array/Accelerate/CUDA/Analysis/Launch.hs +18/−19
- Data/Array/Accelerate/CUDA/Array/Data.hs +27/−9
- Data/Array/Accelerate/CUDA/Array/Nursery.hs +2/−2
- Data/Array/Accelerate/CUDA/Array/Prim.hs +32/−17
- Data/Array/Accelerate/CUDA/Array/Sugar.hs +2/−2
- Data/Array/Accelerate/CUDA/Array/Table.hs +6/−7
- Data/Array/Accelerate/CUDA/Async.hs +1/−1
- Data/Array/Accelerate/CUDA/CodeGen.hs +152/−103
- Data/Array/Accelerate/CUDA/CodeGen/Base.hs +30/−10
- Data/Array/Accelerate/CUDA/CodeGen/IndexSpace.hs +81/−48
- Data/Array/Accelerate/CUDA/CodeGen/Mapping.hs +3/−2
- Data/Array/Accelerate/CUDA/CodeGen/Monad.hs +2/−2
- Data/Array/Accelerate/CUDA/CodeGen/PrefixSum.hs +20/−12
- Data/Array/Accelerate/CUDA/CodeGen/Reduction.hs +52/−35
- Data/Array/Accelerate/CUDA/CodeGen/Stencil.hs +3/−2
- Data/Array/Accelerate/CUDA/CodeGen/Stencil/Extra.hs +7/−8
- Data/Array/Accelerate/CUDA/CodeGen/Type.hs +15/−15
- Data/Array/Accelerate/CUDA/Compile.hs +19/−16
- Data/Array/Accelerate/CUDA/Context.hs +2/−1
- Data/Array/Accelerate/CUDA/Debug.hs +7/−5
- Data/Array/Accelerate/CUDA/Execute.hs +36/−31
- Data/Array/Accelerate/CUDA/Execute/Event.hs +30/−21
- Data/Array/Accelerate/CUDA/Execute/Stream.hs +59/−42
- Data/Array/Accelerate/CUDA/Foreign.hs +1/−1
- Data/Array/Accelerate/CUDA/Foreign/Export.hs +2/−2
- Data/Array/Accelerate/CUDA/Foreign/Import.hs +5/−4
- Data/Array/Accelerate/CUDA/FullList.hs +2/−2
- Data/Array/Accelerate/CUDA/Persistent.hs +51/−24
- Data/Array/Accelerate/CUDA/State.hs +5/−6
- Data/Array/Accelerate/Internal/Check.hs +0/−123
- Setup.hs +1/−18
- accelerate-cuda.buildinfo.in +0/−3
- accelerate-cuda.cabal +21/−34
- configure +0/−3568
- cubits/accelerate_cuda.h +2/−2
- cubits/accelerate_cuda_assert.h +2/−2
- cubits/accelerate_cuda_function.h +21/−21
- cubits/accelerate_cuda_texture.h +2/−2
- cubits/accelerate_cuda_type.h +2/−2
- include/accelerate.h +0/−31
Data/Array/Accelerate/CUDA.hs view
@@ -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:/]
Data/Array/Accelerate/CUDA/AST.hs view
@@ -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>
Data/Array/Accelerate/CUDA/Analysis/Device.hs view
@@ -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>
Data/Array/Accelerate/CUDA/Analysis/Launch.hs view
@@ -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
Data/Array/Accelerate/CUDA/Array/Data.hs view
@@ -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)+
Data/Array/Accelerate/CUDA/Array/Nursery.hs view
@@ -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>
Data/Array/Accelerate/CUDA/Array/Prim.hs view
@@ -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
Data/Array/Accelerate/CUDA/Array/Sugar.hs view
@@ -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>
Data/Array/Accelerate/CUDA/Array/Table.hs view
@@ -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.
Data/Array/Accelerate/CUDA/Async.hs view
@@ -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>
Data/Array/Accelerate/CUDA/CodeGen.hs view
@@ -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
Data/Array/Accelerate/CUDA/CodeGen/Base.hs view
@@ -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"
Data/Array/Accelerate/CUDA/CodeGen/IndexSpace.hs view
@@ -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);+ |]+ ]
Data/Array/Accelerate/CUDA/CodeGen/Mapping.hs view
@@ -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>
Data/Array/Accelerate/CUDA/CodeGen/Monad.hs view
@@ -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>
Data/Array/Accelerate/CUDA/CodeGen/PrefixSum.hs view
@@ -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)
Data/Array/Accelerate/CUDA/CodeGen/Reduction.hs view
@@ -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
Data/Array/Accelerate/CUDA/CodeGen/Stencil.hs view
@@ -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>
Data/Array/Accelerate/CUDA/CodeGen/Stencil/Extra.hs view
@@ -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
Data/Array/Accelerate/CUDA/CodeGen/Type.hs view
@@ -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"
Data/Array/Accelerate/CUDA/Compile.hs view
@@ -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
Data/Array/Accelerate/CUDA/Context.hs view
@@ -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
Data/Array/Accelerate/CUDA/Debug.hs view
@@ -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
Data/Array/Accelerate/CUDA/Execute.hs view
@@ -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
Data/Array/Accelerate/CUDA/Execute/Event.hs view
@@ -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
Data/Array/Accelerate/CUDA/Execute/Stream.hs view
@@ -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
Data/Array/Accelerate/CUDA/Foreign.hs view
@@ -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>
Data/Array/Accelerate/CUDA/Foreign/Export.hs view
@@ -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
Data/Array/Accelerate/CUDA/Foreign/Import.hs view
@@ -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
Data/Array/Accelerate/CUDA/FullList.hs view
@@ -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>
Data/Array/Accelerate/CUDA/Persistent.hs view
@@ -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
Data/Array/Accelerate/CUDA/State.hs view
@@ -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
− Data/Array/Accelerate/Internal/Check.hs
@@ -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-
Setup.hs view
@@ -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
− accelerate-cuda.buildinfo.in
@@ -1,3 +0,0 @@-ghc-options: @ghc_flags@-cc-options: @cpp_flags@-
accelerate-cuda.cabal view
@@ -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)
− configure
@@ -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--
cubits/accelerate_cuda.h view
@@ -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>
cubits/accelerate_cuda_assert.h view
@@ -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>
cubits/accelerate_cuda_function.h view
@@ -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); }
cubits/accelerate_cuda_texture.h view
@@ -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>
cubits/accelerate_cuda_type.h view
@@ -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>
− include/accelerate.h
@@ -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-