accelerate-cuda 0.16.0.0 → 0.17.0.0
raw patch · 37 files changed
+2835/−1872 lines, 37 filesdep +containersdep ~acceleratedep ~basedep ~cuda
Dependencies added: containers
Dependency ranges changed: accelerate, base, cuda, language-c-quote, template-haskell
Files
- Data/Array/Accelerate/CUDA.hs +86/−72
- Data/Array/Accelerate/CUDA/AST.hs +107/−25
- Data/Array/Accelerate/CUDA/Analysis/Launch.hs +10/−10
- Data/Array/Accelerate/CUDA/Analysis/Shape.hs +50/−0
- Data/Array/Accelerate/CUDA/Array/Data.hs +82/−34
- Data/Array/Accelerate/CUDA/Array/Nursery.hs +0/−124
- Data/Array/Accelerate/CUDA/Array/Prim.hs +112/−75
- Data/Array/Accelerate/CUDA/Array/Remote.hs +249/−0
- Data/Array/Accelerate/CUDA/Array/Slice.hs +212/−0
- Data/Array/Accelerate/CUDA/Array/Sugar.hs +5/−5
- Data/Array/Accelerate/CUDA/Array/Table.hs +0/−303
- Data/Array/Accelerate/CUDA/Async.hs +0/−56
- Data/Array/Accelerate/CUDA/CodeGen.hs +177/−276
- Data/Array/Accelerate/CUDA/CodeGen/Arithmetic.hs +393/−0
- Data/Array/Accelerate/CUDA/CodeGen/Base.hs +27/−9
- Data/Array/Accelerate/CUDA/CodeGen/Constant.hs +93/−0
- Data/Array/Accelerate/CUDA/CodeGen/IndexSpace.hs +6/−4
- Data/Array/Accelerate/CUDA/CodeGen/PrefixSum.hs +66/−29
- Data/Array/Accelerate/CUDA/CodeGen/Reduction.hs +3/−3
- Data/Array/Accelerate/CUDA/CodeGen/Streaming.hs +84/−0
- Data/Array/Accelerate/CUDA/CodeGen/Type.hs +138/−93
- Data/Array/Accelerate/CUDA/Compile.hs +219/−109
- Data/Array/Accelerate/CUDA/Context.hs +21/−31
- Data/Array/Accelerate/CUDA/Debug.hs +32/−172
- Data/Array/Accelerate/CUDA/Execute.hs +404/−99
- Data/Array/Accelerate/CUDA/Execute/Event.hs +95/−30
- Data/Array/Accelerate/CUDA/Execute/Stream.hs +21/−25
- Data/Array/Accelerate/CUDA/Foreign/Export.hs +23/−17
- Data/Array/Accelerate/CUDA/Foreign/Import.hs +19/−16
- Data/Array/Accelerate/CUDA/FullList.hs +0/−118
- Data/Array/Accelerate/CUDA/Persistent.hs +20/−19
- Data/Array/Accelerate/CUDA/State.hs +28/−8
- accelerate-cuda.cabal +25/−14
- cubits/accelerate_cuda.h +1/−0
- cubits/accelerate_cuda_exceptional.h +26/−0
- cubits/accelerate_cuda_function.h +0/−96
- cubits/accelerate_cuda_type.h +1/−0
Data/Array/Accelerate/CUDA.hs view
@@ -186,39 +186,38 @@ Arrays, -- * Synchronous execution- run, run1, stream, runIn, run1In, streamIn,+ run, run1, runWith, run1With,+ stream, streamWith, -- * Asynchronous execution Async, wait, poll, cancel,- runAsync, run1Async, runAsyncIn, run1AsyncIn,+ runAsync, run1Async, runAsyncWith, run1AsyncWith, -- * Execution contexts Context, create, destroy,- unsafeFree, unsafeFreeIn, performGC, performGCIn,+ unsafeFree, unsafeFreeWith, performGC, performGCWith, ) where -- standard library-import Control.Exception import Control.Applicative import Control.Monad.Trans import System.IO.Unsafe import Prelude -- friends-import Data.Array.Accelerate.Trafo-import Data.Array.Accelerate.Smart ( Acc ) import Data.Array.Accelerate.Array.Sugar ( Arrays(..), ArraysR(..) )+import Data.Array.Accelerate.Smart ( Acc )+import Data.Array.Accelerate.Async+import Data.Array.Accelerate.Trafo+ import Data.Array.Accelerate.CUDA.Array.Data-import Data.Array.Accelerate.CUDA.Async import Data.Array.Accelerate.CUDA.State import Data.Array.Accelerate.CUDA.Context import Data.Array.Accelerate.CUDA.Compile import Data.Array.Accelerate.CUDA.Execute -#if ACCELERATE_DEBUG-import Data.Array.Accelerate.Debug-#endif+import Data.Array.Accelerate.Debug as Debug -- Accelerate: CUDA@@ -231,21 +230,7 @@ -- Note that it is recommended you use 'run1' whenever possible. -- run :: Arrays a => Acc a -> a-run a- = unsafePerformIO- $ evaluate (runIn defaultContext a)---- | As 'run', but allow the computation to continue running in a thread and--- return immediately without waiting for the result. The status of the--- computation can be queried using 'wait', 'poll', and 'cancel'.------ Note that a CUDA Context can be active on only one host thread at a time. If--- you want to execute multiple computations in parallel, use 'runAsyncIn'.----runAsync :: Arrays a => Acc a -> Async a-runAsync a- = unsafePerformIO- $ evaluate (runAsyncIn defaultContext a)+run = runWith defaultContext -- | As 'run', but execute using the specified device context rather than using -- the default, automatically selected device.@@ -257,24 +242,34 @@ -- 'Foreign.CUDA.Driver.Context.create' pushes the new context on top of the -- stack and makes it current with the calling thread. You should call -- 'Foreign.CUDA.Driver.Context.pop' to make the context floating before passing--- it to 'runIn', which will make it current for the duration of evaluating the+-- it to 'runWith', which will make it current for the duration of evaluating the -- expression. See the CUDA C Programming Guide (G.1) for more information. ---runIn :: Arrays a => Context -> Acc a -> a-runIn ctx a+runWith :: Arrays a => Context -> Acc a -> a+runWith ctx a = unsafePerformIO- $ evaluate (runAsyncIn ctx a) >>= wait+ $ wait =<< runAsyncWith ctx a --- | As 'runIn', but execute asynchronously. Be sure not to destroy the context,+-- | As 'run', but allow the computation to continue running in a thread and+-- return immediately without waiting for the result. The status of the+-- computation can be queried using 'wait', 'poll', and 'cancel'.+--+-- Note that a CUDA Context can be active on only one host thread at a time. If+-- you want to execute multiple computations in parallel, use 'runAsyncWith'.+--+runAsync :: Arrays a => Acc a -> IO (Async a)+runAsync = runAsyncWith defaultContext++-- | As 'runWith', but execute asynchronously. Be sure not to destroy the context, -- or attempt to attach it to a different host thread, before all outstanding -- operations have completed. ---runAsyncIn :: Arrays a => Context -> Acc a -> Async a-runAsyncIn ctx a = unsafePerformIO $ async execute+runAsyncWith :: Arrays a => Context -> Acc a -> IO (Async a)+runAsyncWith ctx a = asyncBound execute where !acc = convertAccWith config a- execute = evalCUDA ctx (compileAcc acc >>= dumpStats >>= executeAcc >>= collect)+ execute = dumpGraph acc >> evalCUDA ctx (compileAcc acc >>= dumpStats >>= executeAcc >>= collect) -- | Prepare and execute an embedded array program of one argument.@@ -310,31 +305,28 @@ -- See the programs in the 'accelerate-examples' package for examples. -- run1 :: (Arrays a, Arrays b) => (Acc a -> Acc b) -> a -> b-run1 f- = unsafePerformIO- $ evaluate (run1In defaultContext f)-+run1 = run1With defaultContext --- | As 'run1', but the computation is executed asynchronously.+-- | As 'run1', but execute in the specified context. ---run1Async :: (Arrays a, Arrays b) => (Acc a -> Acc b) -> a -> Async b-run1Async f- = unsafePerformIO- $ evaluate (run1AsyncIn defaultContext f)+run1With :: (Arrays a, Arrays b) => Context -> (Acc a -> Acc b) -> a -> b+run1With ctx f =+ let go = run1AsyncWith ctx f+ in \a -> unsafePerformIO $ wait =<< go a --- | As 'run1', but execute in the specified context.++-- | As 'run1', but the computation is executed asynchronously. ---run1In :: (Arrays a, Arrays b) => Context -> (Acc a -> Acc b) -> a -> b-run1In ctx f = let go = run1AsyncIn ctx f- in \a -> unsafePerformIO $ wait (go a)+run1Async :: (Arrays a, Arrays b) => (Acc a -> Acc b) -> a -> IO (Async b)+run1Async = run1AsyncWith defaultContext --- | As 'run1In', but execute asynchronously.+-- | As 'run1With', but execute asynchronously. ---run1AsyncIn :: (Arrays a, Arrays b) => Context -> (Acc a -> Acc b) -> a -> Async b-run1AsyncIn ctx f = \a -> unsafePerformIO $ async (execute a)+run1AsyncWith :: (Arrays a, Arrays b) => Context -> (Acc a -> Acc b) -> a -> IO (Async b)+run1AsyncWith ctx f = \a -> asyncBound (execute a) where !acc = convertAfunWith config f- !afun = unsafePerformIO $ evalCUDA ctx (compileAfun acc) >>= dumpStats+ !afun = unsafePerformIO $ dumpGraph acc >> evalCUDA ctx (compileAfun acc) >>= dumpStats execute a = evalCUDA ctx (executeAfun1 afun a >>= collect) -- TLM: We need to be very careful with run1* variants, to ensure that the@@ -345,18 +337,46 @@ -- collecting results as we go. -- stream :: (Arrays a, Arrays b) => (Acc a -> Acc b) -> [a] -> [b]-stream f arrs- = unsafePerformIO- $ evaluate (streamIn defaultContext f arrs)+stream = streamWith defaultContext -- | As 'stream', but execute in the specified context. ---streamIn :: (Arrays a, Arrays b) => Context -> (Acc a -> Acc b) -> [a] -> [b]-streamIn ctx f arrs- = let go = run1In ctx f- in map go arrs+streamWith :: (Arrays a, Arrays b) => Context -> (Acc a -> Acc b) -> [a] -> [b]+streamWith ctx f arrs = map go arrs+ where+ !go = run1With ctx f +{--+-- | Generate a lazy list from a sequence computation.+--+streamOut :: Arrays a => Seq [a] -> [a]+streamOut = streamOutWith defaultContext +streamOutWith :: forall a. Arrays a => Context -> Seq [a] -> [a]+streamOutWith ctx = exec . compile . convertSeq+ where+ compile = unsafePerformIO . evalCUDA ctx . compileSeq+ exec s = go (streamSeq s)+ where+ go !s' = case step s' of+ Nothing -> []+ Just (a, s'') -> a : go s''++ step (StreamSeq ss)+ = unsafePerformIO+ $ evalCUDA ctx+ $ do m <- ss+ case m of+ Nothing -> return Nothing+ Just (a, s') -> collect a >> return (Just (a, s'))+--}+++-- RCE: Similar to run1* variants, we need to be ultra careful with streamOut*+-- in order to make sure that the entire sequence is not reified at once.+-- The steps of the sequence computation should only be performed as needed+-- when elements of the list are forced.+ -- Copy arrays from device to host. -- collect :: forall arrs. Arrays arrs => arrs -> CIO arrs@@ -376,22 +396,16 @@ config = Phase { recoverAccSharing = True , recoverExpSharing = True+ , recoverSeqSharing = True , floatOutAccFromExp = True , enableAccFusion = True , convertOffsetOfSegment = True+ -- , vectoriseSequences = False } dumpStats :: MonadIO m => a -> m a-#if ACCELERATE_DEBUG-dumpStats next = do- stats <- liftIO simplCount- liftIO $ traceMessage dump_simpl_stats (show stats)- liftIO $ resetSimplCount- return next-#else-dumpStats next = return next-#endif+dumpStats next = dumpSimplStats >> return next -- Device memory management@@ -405,10 +419,10 @@ -- the array is currently in use. -- unsafeFree :: Arrays arrs => arrs -> IO ()-unsafeFree = unsafeFreeIn defaultContext+unsafeFree = unsafeFreeWith defaultContext -unsafeFreeIn :: forall arrs. Arrays arrs => Context -> arrs -> IO ()-unsafeFreeIn !ctx !arrs+unsafeFreeWith :: forall arrs. Arrays arrs => Context -> arrs -> IO ()+unsafeFreeWith !ctx !arrs = evalCUDA ctx $ freeR (arrays (undefined :: arrs)) (fromArr arrs) where@@ -421,8 +435,8 @@ -- Release any unused device memory -- performGC :: IO ()-performGC = performGCIn defaultContext+performGC = performGCWith defaultContext -performGCIn :: Context -> IO ()-performGCIn !ctx = evalCUDA ctx cleanupArrayData+performGCWith :: Context -> IO ()+performGCWith !ctx = evalCUDA ctx cleanupArrayData
Data/Array/Accelerate/CUDA/AST.hs view
@@ -1,3 +1,4 @@+{-# LANGUAGE FlexibleContexts #-} {-# LANGUAGE FlexibleInstances #-} {-# LANGUAGE GADTs #-} {-# LANGUAGE GeneralizedNewtypeDeriving #-}@@ -18,17 +19,19 @@ module Data.Array.Accelerate.AST, AccKernel(..), Free, Gamma(..), Idx_(..),- ExecAcc, ExecAfun, ExecOpenAcc(..),+ ExecAcc, ExecAfun, ExecOpenAfun, ExecOpenAcc(..), ExecExp, ExecFun, ExecOpenExp, ExecOpenFun,+ -- ExecSeq(..), ExecOpenSeq(..), ExecP(..), ExecC(..), freevar, makeEnvMap, ) where -- friends import Data.Array.Accelerate.AST-import Data.Array.Accelerate.Pretty+import Data.Array.Accelerate.Lifetime+import Data.Array.Accelerate.Pretty as PP import Data.Array.Accelerate.Array.Sugar ( Array, Shape, Elt )-import qualified Data.Array.Accelerate.CUDA.FullList as FL+import qualified Data.Array.Accelerate.FullList as FL import qualified Foreign.CUDA.Driver as CUDA import qualified Foreign.CUDA.Analysis as CUDA @@ -46,13 +49,13 @@ -- and execution information. -- data AccKernel a where- AccKernel :: !String -- __global__ entry function name- -> {-# UNPACK #-} !CUDA.Fun -- __global__ function object- -> {-# UNPACK #-} !CUDA.Module -- binary module- -> {-# UNPACK #-} !CUDA.Occupancy -- occupancy analysis- -> {-# UNPACK #-} !Int -- thread block size- -> {-# UNPACK #-} !Int -- shared memory per block (bytes)- -> !(Int -> Int) -- number of blocks for input problem size+ AccKernel :: !String -- __global__ entry function name+ -> {-# UNPACK #-} !CUDA.Fun -- __global__ function object+ -> {-# UNPACK #-} !(Lifetime CUDA.Module) -- binary module+ -> {-# UNPACK #-} !CUDA.Occupancy -- occupancy analysis+ -> {-# UNPACK #-} !Int -- thread block size+ -> {-# UNPACK #-} !Int -- shared memory per block (bytes)+ -> !(Int -> Int) -- number of blocks for input problem size -> AccKernel a @@ -122,41 +125,45 @@ => !(PreExp ExecOpenAcc aenv sh) -- shape of the result array, used by execution -> ExecOpenAcc aenv (Array sh e) + -- ExecSeq :: Arrays arrs+ -- => ExecOpenSeq aenv () arrs+ -- -> ExecOpenAcc aenv arrs + -- An annotated AST suitable for execution in the CUDA environment ---type ExecAcc a = ExecOpenAcc () a-type ExecAfun a = PreAfun ExecOpenAcc a+type ExecAcc a = ExecOpenAcc () a+type ExecAfun a = PreAfun ExecOpenAcc a+type ExecOpenAfun aenv a = PreOpenAfun ExecOpenAcc aenv a -type ExecOpenExp = PreOpenExp ExecOpenAcc-type ExecOpenFun = PreOpenFun ExecOpenAcc+type ExecOpenExp = PreOpenExp ExecOpenAcc+type ExecOpenFun = PreOpenFun ExecOpenAcc -type ExecExp = ExecOpenExp ()-type ExecFun = ExecOpenFun ()+type ExecExp = ExecOpenExp ()+type ExecFun = ExecOpenFun () -- Display the annotated AST -- ------------------------- -instance Show (ExecOpenAcc aenv a) where- show = render . prettyExecAcc 0 noParens+instance Show (ExecAcc a) where+ show = render . prettyExecAcc noParens PP.Empty instance Show (ExecAfun a) where- show = render . prettyExecAfun 0-+ show = render . prettyExecAfun -prettyExecAfun :: Int -> ExecAfun a -> Doc-prettyExecAfun alvl pfun = prettyPreAfun prettyExecAcc alvl pfun+prettyExecAfun :: ExecAfun a -> Doc+prettyExecAfun pfun = prettyPreOpenAfun prettyExecAcc PP.Empty pfun prettyExecAcc :: PrettyAcc ExecOpenAcc-prettyExecAcc alvl wrap exec =+prettyExecAcc wrap aenv exec = case exec of EmbedAcc sh -> wrap $ hang (text "Embedded") 2- $ sep [ prettyPreExp prettyExecAcc 0 alvl parens sh ]+ $ sep [ prettyPreExp prettyExecAcc parens aenv sh ] ExecAcc _ (Gamma fv) pacc ->- let base = prettyPreAcc prettyExecAcc alvl wrap pacc+ let base = prettyPreOpenAcc prettyExecAcc wrap aenv pacc ann = braces (freevars (Map.keys fv)) freevars = (text "fv=" <>) . brackets . hcat . punctuate comma . map (\(Idx_ ix) -> char 'a' <> int (idxToInt ix))@@ -169,4 +176,79 @@ Atuple{} -> base Aprj{} -> base _ -> ann <+> base++ -- ExecSeq _ -> text "<SequenceComputation>"++{--+data ExecSeq a where+ ExecS :: Extend ExecOpenAcc () aenv -> ExecOpenSeq aenv () a -> ExecSeq a++data ExecOpenSeq aenv lenv arrs where+ ExecP :: Arrays a => ExecP aenv lenv a -> ExecOpenSeq aenv (lenv, a) arrs -> ExecOpenSeq aenv lenv arrs+ ExecC :: (Arrays a) => ExecC aenv lenv a -> ExecOpenSeq aenv lenv a+ ExecR :: Idx lenv a -> Maybe a -> ExecOpenSeq aenv lenv [a]++data ExecP aenv lenv a where++ ExecToSeq :: (Elt slix, Shape sl, Shape sh, Elt e)+ => SliceIndex (EltRepr slix)+ (EltRepr sl)+ co+ (EltRepr sh)+ -> ExecOpenAcc aenv (Array sh e)+ -> AccKernel (Array sl e)+ -> !(Gamma aenv)+ -> [slix]+ -> ExecP aenv lenv (Array sl e)++ ExecUseLazy :: (Elt slix, Shape sl, Shape sh, Elt e)+ => SliceIndex (EltRepr slix)+ (EltRepr sl)+ co+ (EltRepr sh)+ -> Array sh e+ -> [slix]+ -> ExecP aenv lenv (Array sl e)++ ExecStreamIn :: Arrays a+ => [a]+ -> ExecP aenv lenv a++ ExecMap :: (Arrays a, Arrays b)+ => ExecOpenAfun aenv (a -> b)+ -> Idx lenv a+ -> ExecP aenv lenv b++ ExecZipWith :: (Arrays a, Arrays b, Arrays c)+ => ExecOpenAfun aenv (a -> b -> c)+ -> Idx lenv a+ -> Idx lenv b+ -> ExecP aenv lenv c++ ExecScanSeq :: Elt a+ => ExecFun aenv (a -> a -> a)+ -> ExecExp aenv a+ -> Idx lenv (Scalar a)+ -> Maybe a+ -> ExecP aenv lenv (Scalar a)++data ExecC aenv lenv a where+ ExecFoldSeq :: Elt a+ => ExecFun aenv (a -> a -> a)+ -> ExecExp aenv a+ -> Idx lenv (Scalar a)+ -> Maybe a+ -> ExecC aenv lenv (Scalar a)++ ExecFoldSeqFlatten :: (Arrays a, Shape sh, Elt e)+ => ExecOpenAfun aenv (a -> Vector sh -> Vector e -> a)+ -> ExecOpenAcc aenv a+ -> Idx lenv (Array sh e)+ -> Maybe a+ -> ExecC aenv lenv a++ ExecStuple :: (Arrays a, IsAtuple a)+ => Atuple (ExecC aenv senv) (TupleRepr a)+ -> ExecC aenv senv a+--}
Data/Array/Accelerate/CUDA/Analysis/Launch.hs view
@@ -88,7 +88,7 @@ -> (Int -> Int) -- shared memory as a function of thread block size (bytes) -> (Int, CUDA.Occupancy) blockSize dev acc lim regs smem =- CUDA.optimalBlockSizeBy dev (filter (<= lim) . strategy) (const regs) smem+ CUDA.optimalBlockSizeOf dev (filter (<= lim) (strategy dev)) (const regs) smem where strategy = case acc of Fold _ _ _ -> CUDA.incPow2@@ -101,7 +101,6 @@ Scanr1 _ _ -> CUDA.incWarp _ -> CUDA.decWarp - -- | -- Determine the number of blocks of the given size necessary to process the -- given array expression. This should understand things like #elements per@@ -116,17 +115,17 @@ -- for 1D reductions this is the total number of elements -- gridSize :: CUDA.DeviceProperties -> PreOpenAcc DelayedOpenAcc aenv a -> Int -> Int -> Int-gridSize p acc@(FoldSeg _ _ _ _) size cta = split acc (size * CUDA.warpSize p) cta-gridSize p acc@(Fold1Seg _ _ _) size cta = split acc (size * CUDA.warpSize p) cta-gridSize _ acc@(Fold _ _ _) size cta = if preAccDim delayedDim acc == 0 then split acc size cta else max 1 size-gridSize _ acc@(Fold1 _ _) size cta = if preAccDim delayedDim acc == 0 then split acc size cta else max 1 size-gridSize _ acc size cta = split acc size cta+gridSize p (FoldSeg _ _ _ _) size cta = split (size * CUDA.warpSize p) cta+gridSize p (Fold1Seg _ _ _) size cta = split (size * CUDA.warpSize p) cta+gridSize _ acc@(Fold _ _ _) size cta = if preAccDim delayedDim acc == 0 then split size cta else max 1 size+gridSize _ acc@(Fold1 _ _) size cta = if preAccDim delayedDim acc == 0 then split size cta else max 1 size+gridSize _ _ size cta = split size cta -split :: acc aenv a -> Int -> Int -> Int-split acc size cta = (size `between` eltsPerThread acc) `between` cta+split :: Int -> Int -> Int+split size cta = (size `between` eltsPerThread) `between` cta where between arr n = 1 `max` ((n + arr - 1) `div` n)- eltsPerThread _ = 1+ eltsPerThread = 1 -- |@@ -171,4 +170,5 @@ (blockDim `div` CUDA.warpSize p) * 8 + blockDim * sizeOf (delayedExpType x) -- TLM: why 8? I can't remember... sharedMem p (Fold1Seg _ a _) blockDim = (blockDim `div` CUDA.warpSize p) * 8 + blockDim * sizeOf (delayedAccType a)+-- sharedMem _ (Collect _) _ = 0
+ Data/Array/Accelerate/CUDA/Analysis/Shape.hs view
@@ -0,0 +1,50 @@+{-# LANGUAGE GADTs #-}+{-# LANGUAGE ScopedTypeVariables #-}+{-# LANGUAGE TypeOperators #-}+-- |+-- Module : Data.Array.Accelerate.CUDA.Analysis.Shape+-- Copyright : [2012..2014] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.CUDA.Analysis.Shape (++ module Data.Array.Accelerate.Analysis.Shape,+ module Data.Array.Accelerate.CUDA.Analysis.Shape,+ (:~:)(..),++) where++import Data.Typeable++import Data.Array.Accelerate.Array.Sugar+import Data.Array.Accelerate.Analysis.Shape+import Data.Array.Accelerate.Analysis.Match++import Data.Array.Accelerate.CUDA.AST+++-- | Reify dimensionality of the result type of an array computation+--+execAccDim :: AccDim ExecOpenAcc+execAccDim (ExecAcc _ _ pacc) = preAccDim execAccDim pacc+++-- Match reified shape types+--+matchShapeType+ :: forall sh sh'. (Shape sh, Shape sh')+ => sh+ -> sh'+ -> Maybe (sh :~: sh')+matchShapeType _ _+ | Just Refl <- matchTupleType (eltType (undefined::sh)) (eltType (undefined::sh'))+ = gcast Refl++matchShapeType _ _+ = Nothing+
Data/Array/Accelerate/CUDA/Array/Data.hs view
@@ -3,7 +3,9 @@ {-# LANGUAGE GADTs #-} {-# LANGUAGE ScopedTypeVariables #-} {-# LANGUAGE TemplateHaskell #-}+{-# LANGUAGE TupleSections #-} {-# LANGUAGE TypeFamilies #-}+{-# LANGUAGE TypeOperators #-} -- | -- Module : Data.Array.Accelerate.CUDA.Array.Data -- Copyright : [2008..2014] Manuel M T Chakravarty, Gabriele Keller@@ -21,17 +23,17 @@ -- * Array operations and representations mallocArray, freeArray, indexArray,- useArray, useArrayAsync,+ useArray, useArrayAsync, useArraySlice, useDevicePtrs, copyArray, copyArrayAsync, copyArrayPeer, copyArrayPeerAsync, peekArray, peekArrayAsync, pokeArray, pokeArrayAsync, marshalArrayData, marshalTextureData, marshalDevicePtrs,- devicePtrsOfArrayData, advancePtrsOfArrayData,+ withDevicePtrs, advancePtrsOfArrayData, devicePtrsFromList, devicePtrsToWordPtrs, -- * Garbage collection- cleanupArrayData+ cleanupArrayData, ) where @@ -40,18 +42,24 @@ import Control.Monad.Reader ( asks ) import Control.Monad.State ( gets ) import Control.Monad.Trans ( liftIO )+import Control.Monad.Trans.Cont import Foreign.C.Types import Foreign.Ptr import Prelude hiding ( fst, snd ) import qualified Prelude as P + -- 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 )+import Data.Array.Accelerate.Array.Sugar ( Array(..), Shape, Elt, fromElt, toElt, EltRepr )+import Data.Array.Accelerate.Array.Representation ( size, SliceIndex ) import Data.Array.Accelerate.CUDA.State-import Data.Array.Accelerate.CUDA.Array.Table+import Data.Array.Accelerate.CUDA.Array.Slice ( TransferDesc, transferDesc )+import Data.Array.Accelerate.CUDA.Array.Remote+import Data.Array.Accelerate.CUDA.Persistent ( KernelTable )+import Data.Array.Accelerate.CUDA.Execute.Event ( EventTable )+import Data.Array.Accelerate.CUDA.Execute.Stream ( Reservoir ) import qualified Data.Array.Accelerate.CUDA.Array.Prim as Prim import qualified Foreign.CUDA.Driver as CUDA import qualified Foreign.CUDA.Driver.Stream as CUDA@@ -83,6 +91,15 @@ mt <- gets memoryTable liftIO $! f ctx mt +run' :: (Context -> MemoryTable -> KernelTable -> Reservoir -> EventTable -> IO a) -> CIO a+run' f = do+ ctx <- asks activeContext+ mt <- gets memoryTable+ kt <- gets kernelTable+ rsv <- gets streamReservoir+ etbl <- gets eventTable+ liftIO $! f ctx mt kt rsv etbl+ -- CPP hackery to generate the cases where we dispatch to the worker function handling -- elementary types. --@@ -184,6 +201,27 @@ mkPrimDispatch(usePrim,Prim.useArrayAsync) +-- | Upload a slice of an existing array (eg. row of a matrix) to the+-- device. TODO : Bounds checking, generalize slices to more than just+-- inner dimension?+useArraySlice :: (Elt slix, Shape sl, Shape dim, Elt e)+ => SliceIndex (EltRepr slix) (EltRepr sl) co (EltRepr dim)+ -> slix -- Slice+ -> Array dim e -- Host array+ -> Array sl e -- Device array+ -> CIO ()+useArraySlice slix sl (Array dim !adata_host) (Array _ !adata_dev) = run doUse+ where+ tdesc = transferDesc slix (fromElt sl) dim+ doUse !ctx !mt = useR arrayElt adata_host adata_dev+ where+ useR :: ArrayEltR e -> ArrayData e -> ArrayData e -> IO ()+ useR ArrayEltRunit _ _ = return ()+ useR (ArrayEltRpair aeR1 aeR2) adh add = useR aeR1 (fst adh) (fst add) >> useR aeR2 (snd adh) (snd add)+ useR aer adh add = usePrim aer ctx mt adh add tdesc -- usePrim aer ctx mt adh add tdesc+ usePrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> ArrayData e -> TransferDesc -> IO ()+ mkPrimDispatch(usePrim,Prim.useArraySlice)+ useDevicePtrs :: (Shape sh, Elt e) => EltRepr sh -> Prim.DevicePtrs (EltRepr e) -> CIO (Array sh e) useDevicePtrs sh ptrs = run doUse where@@ -420,61 +458,71 @@ marshalDevicePtrs !adata ptrs = map (CUDA.VArg . CUDA.wordPtrToDevPtr) $ devicePtrsToWordPtrs adata ptrs -- |Wrap the device pointers corresponding to a host-side array into arguments--- that can be passed to a kernel upon invocation.+-- that can be passed to a kernel upon invocation and call the+-- supplied continuation. Any asynchronous CUDA functions called by the+-- continuation must be in the same stream as given by the 2nd argument. ---marshalArrayData :: ArrayElt e => ArrayData e -> CIO [CUDA.FunParam]-marshalArrayData !adata = run doMarshal+marshalArrayData :: ArrayElt e => ArrayData e -> Maybe CUDA.Stream -> ([CUDA.FunParam] -> CIO b) -> CIO b+marshalArrayData !adata ms f = run' doMarshal where- doMarshal !ctx !mt = marshalR arrayElt adata+ doMarshal !ctx !mt !kt !rsv !etbl = runContT (marshalR arrayElt adata) (evalCUDAState ctx mt kt rsv etbl . f) where- marshalR :: ArrayEltR e -> ArrayData e -> IO [CUDA.FunParam]+ marshalR :: ArrayEltR e -> ArrayData e -> ContT b IO [CUDA.FunParam] marshalR ArrayEltRunit _ = return [] marshalR (ArrayEltRpair aeR1 aeR2) ad = (++) <$> marshalR aeR1 (fst ad) <*> marshalR aeR2 (snd ad)- marshalR aer ad = return <$> marshalPrim aer ctx mt ad+ marshalR aer ad = do+ param <- ContT $ marshalPrim aer ctx mt ad ms+ return [param] --- marshalPrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> IO CUDA.FunParam+ marshalPrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> Maybe CUDA.Stream -> (CUDA.FunParam -> IO b) -> IO b mkPrimDispatch(marshalPrim,Prim.marshalArrayData) -- |Bind the device memory arrays to the given texture reference(s), setting--- appropriate type. The arrays are bound, and the list of textures thereby--- consumed, in projection index order --- i.e. right-to-left+-- appropriate type, and call the supplied continuation. The arrays are bound,+-- and the list of textures thereby consumed, in projection index order+-- --- i.e. right-to-left+-- The textures should only be considered bound during the execution of the+-- continuation. Any asynchronous CUDA functions called by the continuation+-- must be in the same stream as given by the 4th argument. ---marshalTextureData :: ArrayElt e => ArrayData e -> Int -> [CUDA.Texture] -> CIO ()-marshalTextureData !adata !n !texs = run doMarshal+marshalTextureData :: ArrayElt e => ArrayData e -> Int -> [CUDA.Texture] -> Maybe CUDA.Stream -> ([CUDA.Texture] -> CIO b) -> CIO b+marshalTextureData !adata !n !texs ms f = run' doMarshal where- doMarshal !ctx !mt = marshalR arrayElt adata texs >> return ()+ doMarshal !ctx !mt !kt !rsv !etbl = runContT (marshalR arrayElt adata texs) (evalCUDAState ctx mt kt rsv etbl . \(_,ts) -> f ts) where- marshalR :: ArrayEltR e -> ArrayData e -> [CUDA.Texture] -> IO Int- marshalR ArrayEltRunit _ _ = return 0+ marshalR :: ArrayEltR e -> ArrayData e -> [CUDA.Texture] -> ContT b IO (Int, [CUDA.Texture])+ marshalR ArrayEltRunit _ _ = return (0, []) marshalR (ArrayEltRpair aeR1 aeR2) ad t- = do r <- marshalR aeR2 (snd ad) t- l <- marshalR aeR1 (fst ad) (drop r t)- return (l + r)+ = do (r, rs) <- marshalR aeR2 (snd ad) t+ (l, ls) <- marshalR aeR1 (fst ad) (drop r t)+ return (l + r, ls ++ rs) marshalR aer ad t- = do marshalPrim aer ctx mt ad n (head t)- return 1+ = do param <- ContT $ marshalPrim aer ctx mt ad n (head t) ms+ return (1, [param]) --- marshalPrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> Int -> CUDA.Texture -> IO ()+ marshalPrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> Int -> CUDA.Texture -> Maybe CUDA.Stream -> (CUDA.Texture -> IO b) -> IO b mkPrimDispatch(marshalPrim,Prim.marshalTextureData) --- |Raw device pointers associated with a host-side array+-- | Perform an operation using the device pointers of the given array. Any+-- asynchronous CUDA functions called by the supplied continuation must be in+-- the same stream as given by the second argument. ---devicePtrsOfArrayData :: ArrayElt e => ArrayData e -> CIO (Prim.DevicePtrs e)-devicePtrsOfArrayData !adata = run ptrs+withDevicePtrs :: ArrayElt e => ArrayData e -> Maybe CUDA.Stream -> (Prim.DevicePtrs e -> CIO b) -> CIO b+withDevicePtrs !adata ms f = run' ptrs where- ptrs !ctx !mt = ptrsR arrayElt adata+ ptrs !ctx !mt !kt !rsv !etbl = runContT (ptrsR arrayElt adata) (evalCUDAState ctx mt kt rsv etbl . f) where- ptrsR :: ArrayEltR e -> ArrayData e -> IO (Prim.DevicePtrs e)+ ptrsR :: ArrayEltR e -> ArrayData e -> ContT b IO (Prim.DevicePtrs e) ptrsR ArrayEltRunit _ = return () ptrsR (ArrayEltRpair aeR1 aeR2) ad = (,) <$> ptrsR aeR1 (fst ad) <*> ptrsR aeR2 (snd ad)- ptrsR aer ad = ptrsPrim aer ctx mt ad+ ptrsR aer ad = ContT $ ptrsPrim aer ctx mt ad ms --- ptrsPrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> IO (Prim.DevicePtrs e)- mkPrimDispatch(ptrsPrim,Prim.devicePtrsOfArrayData)+ ptrsPrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> Maybe CUDA.Stream -> (Prim.DevicePtrs e -> IO b) -> IO b+ mkPrimDispatch(ptrsPrim,Prim.withDevicePtrs) -- |Advance a set of device pointers by the given number of elements each
− Data/Array/Accelerate/CUDA/Array/Nursery.hs
@@ -1,124 +0,0 @@-{-# LANGUAGE BangPatterns #-}-{-# LANGUAGE ViewPatterns #-}-{-# OPTIONS_GHC -fno-warn-orphans #-}--- |--- Module : Data.Array.Accelerate.CUDA.Array.Nursery--- 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>--- Stability : experimental--- Portability : non-portable (GHC extensions)-----module Data.Array.Accelerate.CUDA.Array.Nursery (-- Nursery(..), NRS, new, malloc, stash, flush,--) where---- friends-import Data.Array.Accelerate.CUDA.FullList ( FullList(..) )-import qualified Data.Array.Accelerate.CUDA.FullList as FL-import qualified Data.Array.Accelerate.CUDA.Debug as D---- libraries-import Prelude-import Data.Hashable-import Control.Exception ( bracket_ )-import Control.Concurrent.MVar ( MVar, newMVar, withMVar, mkWeakMVar )-import System.Mem.Weak ( Weak )-import Foreign.Ptr ( ptrToIntPtr )-import Foreign.CUDA.Ptr ( DevicePtr )--import qualified Foreign.CUDA.Driver as CUDA-import qualified Data.HashTable.IO as HT----- The nursery is a place to store device memory arrays that are no longer--- needed. If a new array is requested of a similar size, we might return an--- array from the nursery instead of calling into the CUDA API to allocate fresh--- memory.------ Note that pointers are also related to a specific context, so we must include--- that when looking up the map.------ Note that since there might be many arrays for the same size, each entry in--- the map keeps a (non-empty) list of device pointers.----type HashTable key val = HT.BasicHashTable key val--type NRS = MVar ( HashTable (CUDA.Context, Int) (FullList () (DevicePtr ())) )-data Nursery = Nursery {-# UNPACK #-} !NRS- {-# UNPACK #-} !(Weak NRS)--instance Hashable CUDA.Context where- {-# INLINE hashWithSalt #-}- hashWithSalt salt (CUDA.Context ctx)- = salt `hashWithSalt` (fromIntegral (ptrToIntPtr ctx) :: Int)----- Generate a fresh nursery----new :: IO Nursery-new = do- tbl <- HT.new- ref <- newMVar tbl- weak <- mkWeakMVar ref (flush tbl)- return $! Nursery ref weak----- Look for a chunk of memory in the nursery of a given size (or a little bit--- larger). If found, it is removed from the nursery and a pointer to it--- returned.----{-# INLINE malloc #-}-malloc :: Int -> CUDA.Context -> Nursery -> IO (Maybe (DevicePtr ()))-malloc !n !ctx (Nursery !ref _) = withMVar ref $ \tbl -> do- let !key = (ctx,n)- --- mp <- HT.lookup tbl key- case mp of- Nothing -> return Nothing- Just (FL () ptr rest) ->- case rest of- FL.Nil -> HT.delete tbl key >> return (Just ptr)- FL.Cons () v xs -> HT.insert tbl key (FL () v xs) >> return (Just ptr)----- Add a device pointer to the nursery.----{-# INLINE stash #-}-stash :: Int -> CUDA.Context -> NRS -> DevicePtr a -> IO ()-stash !n !ctx !ref (CUDA.castDevPtr -> !ptr) = withMVar ref $ \tbl -> do- let !key = (ctx, n)- --- mp <- HT.lookup tbl key- case mp of- Nothing -> HT.insert tbl key (FL.singleton () ptr)- Just xs -> HT.insert tbl key (FL.cons () ptr xs)----- Delete all entries from the nursery and free all associated device memory.----flush :: HashTable (CUDA.Context,Int) (FullList () (CUDA.DevicePtr ())) -> IO ()-flush !tbl =- let clean (!key@(ctx,_),!val) = do- bracket_ (CUDA.push ctx) CUDA.pop (FL.mapM_ (const CUDA.free) val)- HT.delete tbl key- in- message "flush nursery" >> HT.mapM_ clean tbl----- Debug--- -------{-# INLINE trace #-}-trace :: String -> IO a -> IO a-trace msg next = D.message D.dump_gc ("gc: " ++ msg) >> next--{-# INLINE message #-}-message :: String -> IO ()-message s = s `trace` return ()-
Data/Array/Accelerate/CUDA/Array/Prim.hs view
@@ -1,4 +1,5 @@ {-# LANGUAGE BangPatterns #-}+{-# LANGUAGE ConstraintKinds #-} {-# LANGUAGE CPP #-} {-# LANGUAGE GADTs #-} {-# LANGUAGE ScopedTypeVariables #-}@@ -21,24 +22,24 @@ DevicePtrs, HostPtrs, mallocArray, indexArray,- useArray, useArrayAsync, useDevicePtrs,+ useArray, useArrayAsync, useArraySlice, useDevicePtrs, copyArray, copyArrayAsync, copyArrayPeer, copyArrayPeerAsync, peekArray, peekArrayAsync, pokeArray, pokeArrayAsync, marshalDevicePtrs, marshalArrayData, marshalTextureData,- devicePtrsOfArrayData, advancePtrsOfArrayData+ withDevicePtrs, advancePtrsOfArrayData ) where -- libraries+import Prelude hiding ( lookup ) import Data.Int import Data.Word-import Data.Maybe-import Data.Functor import Data.Typeable import Control.Monad import Language.Haskell.TH import System.Mem.StableName+import Foreign.CUDA.Ptr ( plusDevPtr ) import Foreign.Ptr import Foreign.C.Types import Foreign.Storable@@ -46,13 +47,14 @@ import qualified Foreign.CUDA.Driver as CUDA import qualified Foreign.CUDA.Driver.Stream as CUDA import qualified Foreign.CUDA.Driver.Texture as CUDA-import Prelude hiding ( lookup ) -- friends import Data.Array.Accelerate.Error+import Data.Array.Accelerate.Lifetime ( withLifetime ) import Data.Array.Accelerate.Array.Data import Data.Array.Accelerate.CUDA.Context-import Data.Array.Accelerate.CUDA.Array.Table+import Data.Array.Accelerate.CUDA.Array.Slice ( TransferDesc(..), blocksOf )+import Data.Array.Accelerate.CUDA.Array.Remote import qualified Data.Array.Accelerate.CUDA.Debug as D @@ -155,7 +157,7 @@ -- release any inaccessible arrays and try again. -- mallocArray- :: forall e a. (ArrayElt e, DevicePtrs e ~ CUDA.DevicePtr a, Typeable e, Typeable a, Storable a)+ :: forall e a. (PrimElt e a, DevicePtrs e ~ CUDA.DevicePtr a) => Context -> MemoryTable -> ArrayData e@@ -164,18 +166,16 @@ mallocArray !ctx !mt !ad !n0 = do let !n = 1 `max` n0 !bytes = n * sizeOf (undefined :: a)- exists <- isJust <$> (lookup ctx mt ad :: IO (Maybe (CUDA.DevicePtr a)))- unless exists $ do- message $ "mallocArray: " ++ showBytes bytes- _ <- malloc ctx mt ad n :: IO (CUDA.DevicePtr a)- return ()+ message $ "mallocArray: " ++ showBytes bytes+ _ <- malloc ctx mt ad False n+ return () -- A combination of 'mallocArray' and 'pokeArray' to allocate space on the -- device and upload an existing array. This is specialised because if the host -- array is shared on the heap, we do not need to do anything. -- useArray- :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr a, Typeable e, Typeable a, Storable a)+ :: forall e a. (PrimElt e a, DevicePtrs e ~ CUDA.DevicePtr a) => Context -> MemoryTable -> ArrayData e@@ -185,13 +185,36 @@ let src = ptrsOfArrayData ad !n = 1 `max` n0 !bytes = n * sizeOf (undefined :: a)++ run dst = transfer "useArray/malloc" bytes $ CUDA.pokeArray n src dst in do- exists <- isJust <$> (lookup ctx mt ad :: IO (Maybe (CUDA.DevicePtr a)))- unless exists $ do- dst <- malloc ctx mt ad n- transfer "useArray/malloc" bytes $ CUDA.pokeArray n src dst+ alloc <- malloc ctx mt ad True n+ when alloc $ withDevicePtrs ctx mt ad Nothing run +-- A combination of 'mallocArray' and 'pokeArray' to allocate space on the+-- device and upload an existing array. This is specialised because if the host+-- array is shared on the heap, we do not need to do anything.+--+useArraySlice+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr a, Typeable e, Typeable a, Storable a)+ => Context+ -> MemoryTable+ -> ArrayData e+ -> ArrayData e+ -> TransferDesc+ -> IO ()+useArraySlice !ctx !mt !ad_host !ad_dev !tdesc =+ let src = ptrsOfArrayData ad_host+ k = sizeOf (undefined :: a)+ run dst =+ sequence_+ [ transfer "useArraySlice/malloc" (k * size) $ CUDA.pokeArray size (plusPtr src (k * src_offset)) (plusDevPtr dst (k * dst_offset))+ | (src_offset, dst_offset, size) <- blocksOf tdesc]+ in do+ alloc <- malloc ctx mt ad_dev True (k * nblocks tdesc * blocksize tdesc)+ when alloc $ withDevicePtrs ctx mt ad_dev Nothing run + useArrayAsync :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr a, Typeable e, Typeable a, Storable a) => Context@@ -204,11 +227,11 @@ let src = CUDA.HostPtr (ptrsOfArrayData ad) !n = 1 `max` n0 !bytes = n * sizeOf (undefined :: a)++ run dst = transfer "useArray/malloc" bytes $ CUDA.pokeArrayAsync n src dst ms in do- exists <- isJust <$> (lookup ctx mt ad :: IO (Maybe (CUDA.DevicePtr a)))- unless exists $ do- dst <- malloc ctx mt ad n- transfer "useArrayAsync/malloc" bytes $ CUDA.pokeArrayAsync n src dst ms+ alloc <- malloc ctx mt ad True n+ when alloc $ withDevicePtrs ctx mt ad ms run useDevicePtrs@@ -224,14 +247,14 @@ (adata, _) = runArrayData $ (,undefined) `fmap` newArrayData n in do message $ "useDevicePtrs: " ++ showBytes bytes- insertRemote ctx mt adata ptr+ insertUnmanaged ctx mt adata ptr return adata -- Read a single element from an array at the given row-major index -- indexArray- :: forall e a. (ArrayElt e, DevicePtrs e ~ CUDA.DevicePtr a, Typeable e, Typeable a, Storable a)+ :: forall e a. (PrimElt e a, DevicePtrs e ~ CUDA.DevicePtr a) => Context -> MemoryTable -> ArrayData e@@ -239,7 +262,7 @@ -> IO a indexArray !ctx !mt !ad !i = alloca $ \dst ->- devicePtrsOfArrayData ctx mt ad >>= \src -> do+ withDevicePtrs ctx mt ad Nothing $ \src -> do message $ "indexArray: " ++ showBytes (sizeOf (undefined::a)) CUDA.peekArray 1 (src `CUDA.advanceDevPtr` i) dst peek dst@@ -249,21 +272,21 @@ -- respect to the host, but will never overlap kernel execution. -- copyArray- :: forall e a b. (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr b, Typeable a, Typeable b, Typeable e, Storable b)+ :: forall e a. (PrimElt e a, DevicePtrs e ~ CUDA.DevicePtr a) => Context -> MemoryTable -> ArrayData e -- source array -> ArrayData e -- destination array -> Int -- number of array elements -> IO ()-copyArray !ctx !mt !from !to !n = do- src <- devicePtrsOfArrayData ctx mt from- dst <- devicePtrsOfArrayData ctx mt to- transfer "copyArray" (n * sizeOf (undefined :: b)) $+copyArray !ctx !mt !from !to !n =+ withDevicePtrs ctx mt from Nothing $ \src ->+ withDevicePtrs ctx mt to Nothing $ \dst -> do+ transfer "copyArray" (n * sizeOf (undefined :: a)) $ 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)+ :: forall e a. (PrimElt e a, DevicePtrs e ~ CUDA.DevicePtr a) => Context -> MemoryTable -> ArrayData e -- source array@@ -271,41 +294,45 @@ -> 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+copyArrayAsync !ctx !mt !from !to !n !mst =+ withDevicePtrs ctx mt from Nothing$ \src ->+ withDevicePtrs ctx mt to Nothing $ \dst -> do+ transfer "copyArrayAsync" (n * sizeOf (undefined :: a)) $+ CUDA.copyArrayAsync n src dst mst -- Copy data between two device arrays that exist in different contexts and/or -- devices. -- copyArrayPeer- :: forall e a b. (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr b, Typeable a, Typeable b, Typeable e, Storable b)+ :: forall e a. (PrimElt e a, DevicePtrs e ~ CUDA.DevicePtr a) => MemoryTable -> ArrayData e -> Context -- source array and context -> ArrayData e -> Context -- destination array and context -> Int -- number of array elements -> IO ()-copyArrayPeer !mt !from !ctxSrc !to !ctxDst !n = do- src <- devicePtrsOfArrayData ctxSrc mt from- dst <- devicePtrsOfArrayData ctxDst mt to- transfer "copyArrayPeer" (n * sizeOf (undefined :: b)) $- CUDA.copyArrayPeer n src (deviceContext ctxSrc) dst (deviceContext ctxDst)+copyArrayPeer !mt !from !ctxSrc !to !ctxDst !n =+ withDevicePtrs ctxSrc mt from Nothing $ \src ->+ withDevicePtrs ctxDst mt to Nothing $ \dst ->+ withLifetime (deviceContext ctxSrc) $ \dctxSrc ->+ withLifetime (deviceContext ctxDst) $ \dctxDst -> do+ transfer "copyArrayPeer" (n * sizeOf (undefined :: a)) $+ CUDA.copyArrayPeer n src dctxSrc dst dctxDst copyArrayPeerAsync- :: forall e a b. (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr b, Typeable a, Typeable b, Typeable e, Storable b)+ :: forall e a. (PrimElt e a, DevicePtrs e ~ CUDA.DevicePtr a) => MemoryTable -> ArrayData e -> Context -- source array and context -> ArrayData e -> Context -- destination array and context -> Int -- number of array elements -> Maybe CUDA.Stream -> IO ()-copyArrayPeerAsync !mt !from !ctxSrc !to !ctxDst !n !st = do- src <- devicePtrsOfArrayData ctxSrc mt from- dst <- devicePtrsOfArrayData ctxDst mt to- transfer "copyArrayPeerAsync" (n * sizeOf (undefined :: b)) $- CUDA.copyArrayPeerAsync n src (deviceContext ctxSrc) dst (deviceContext ctxDst) st+copyArrayPeerAsync !mt !from !ctxSrc !to !ctxDst !n !st =+ withDevicePtrs ctxSrc mt from st $ \src ->+ withDevicePtrs ctxDst mt to st $ \dst ->+ withLifetime (deviceContext ctxSrc) $ \dctxSrc ->+ withLifetime (deviceContext ctxDst) $ \dctxDst ->+ transfer "copyArrayPeerAsync" (n * sizeOf (undefined :: a)) $+ CUDA.copyArrayPeerAsync n src dctxSrc dst dctxDst st -- Copy data from the device into the associated Accelerate host-side array@@ -318,7 +345,7 @@ -> Int -> IO () peekArray !ctx !mt !ad !n =- devicePtrsOfArrayData ctx mt ad >>= \src ->+ withDevicePtrs ctx mt ad Nothing $ \src -> transfer "peekArray" (n * sizeOf (undefined :: a)) $ CUDA.peekArray n src (ptrsOfArrayData ad) @@ -331,7 +358,7 @@ -> Maybe CUDA.Stream -> IO () peekArrayAsync !ctx !mt !ad !n !st =- devicePtrsOfArrayData ctx mt ad >>= \src ->+ withDevicePtrs ctx mt ad st $ \src -> transfer "peekArrayAsync" (n * sizeOf (undefined :: a)) $ CUDA.peekArrayAsync n src (CUDA.HostPtr $ ptrsOfArrayData ad) st @@ -346,7 +373,7 @@ -> Int -> IO () pokeArray !ctx !mt !ad !n =- devicePtrsOfArrayData ctx mt ad >>= \dst ->+ withDevicePtrs ctx mt ad Nothing $ \dst -> transfer "pokeArray: " (n * sizeOf (undefined :: a)) $ CUDA.pokeArray n (ptrsOfArrayData ad) dst @@ -359,7 +386,7 @@ -> Maybe CUDA.Stream -> IO () pokeArrayAsync !ctx !mt !ad !n !st =- devicePtrsOfArrayData ctx mt ad >>= \dst ->+ withDevicePtrs ctx mt ad st $ \dst -> transfer "pokeArrayAsync: " (n * sizeOf (undefined :: a)) $ CUDA.pokeArrayAsync n (CUDA.HostPtr $ ptrsOfArrayData ad) dst st @@ -375,49 +402,63 @@ -- Wrap a device pointer corresponding corresponding to a host-side array into--- arguments that can be passed to a kernel upon invocation+-- arguments that can be passed to a kernel upon invocation and call the+-- supplied continuation. Any asynchronous CUDA functions called by the+-- continuation must be in the same stream as given by the 4th argument. -- marshalArrayData- :: (ArrayElt e, DevicePtrs e ~ CUDA.DevicePtr b, Typeable b, Typeable e)+ :: (PrimElt e a, DevicePtrs e ~ CUDA.DevicePtr a) => Context -> MemoryTable -> ArrayData e- -> IO CUDA.FunParam-marshalArrayData !ctx !mt !ad = marshalDevicePtrs ad <$> devicePtrsOfArrayData ctx mt ad+ -> Maybe CUDA.Stream+ -> (CUDA.FunParam -> IO b)+ -> IO b+marshalArrayData !ctx !mt !ad ms run = withDevicePtrs ctx mt ad ms (run . marshalDevicePtrs ad) --- Bind device memory to the given texture reference, setting appropriate type+-- Bind device memory to the given texture reference, setting appropriate type,+-- and call the supplied continuation. The texture should only be considered+-- bound during the execution of the continuation. Any asynchronous CUDA+-- functions called by the continuation must be in the same stream as given by+-- the 6th argument. -- marshalTextureData- :: forall a e. (ArrayElt e, DevicePtrs e ~ CUDA.DevicePtr a, Typeable a, Typeable e, Storable a, TextureData a)+ :: forall a e b. (PrimElt e a, DevicePtrs e ~ CUDA.DevicePtr a, TextureData a) => Context -> MemoryTable -> ArrayData e -- host array -> Int -- number of elements -> CUDA.Texture -- texture reference to bind array to- -> IO ()-marshalTextureData !ctx !mt !ad !n !tex =+ -> Maybe CUDA.Stream+ -> (CUDA.Texture -> IO b)+ -> IO b+marshalTextureData !ctx !mt !ad !n !tex ms run = let (fmt, c) = format (undefined :: a)- in devicePtrsOfArrayData ctx mt ad >>= \ptr -> do+ in withDevicePtrs ctx mt ad ms $ \ptr -> do CUDA.setFormat tex fmt c CUDA.bind tex ptr (fromIntegral $ n * sizeOf (undefined :: a))-+ run tex --- Lookup the device memory associated with our host array+-- Perform an operation using the device pointer of the given array. Any+-- asynchronous CUDA functions called by the supplied continuation must be in+-- the same stream as given by the 4th argument. ---devicePtrsOfArrayData- :: (ArrayElt e, DevicePtrs e ~ CUDA.DevicePtr b, Typeable e, Typeable b)+withDevicePtrs+ :: (PrimElt e a, DevicePtrs e ~ CUDA.DevicePtr a) => Context -> MemoryTable -> ArrayData e- -> IO (DevicePtrs e)-devicePtrsOfArrayData !ctx !mt !ad = do- mv <- lookup ctx mt ad- case mv of- Just v -> return v+ -> Maybe (CUDA.Stream)+ -> (DevicePtrs e -> IO b)+ -> IO b+withDevicePtrs !ctx !mt !ad ms run = do+ mb <- withRemote ctx mt ad run ms+ case mb of+ Just b -> return b Nothing -> do sn <- makeStableName ad- $internalError "devicePtrsOfArrayData" $ "lost device memory #" ++ show (hashStableName sn)+ $internalError "withDevicePtrs" $ "lost device memory #" ++ show (hashStableName sn) -- Advance device pointers by a given number of elements@@ -438,13 +479,9 @@ showBytes :: Int -> String showBytes x = D.showFFloatSIBase (Just 0) 1024 (fromIntegral x :: Double) "B" -{-# INLINE trace #-}-trace :: String -> IO a -> IO a-trace msg next = D.message D.dump_gc ("gc: " ++ msg) >> next- {-# INLINE message #-} message :: String -> IO ()-message s = s `trace` return ()+message msg = D.traceIO D.dump_gc ("gc: " ++ msg) {-# INLINE transfer #-} transfer :: String -> Int -> IO () -> IO ()
+ Data/Array/Accelerate/CUDA/Array/Remote.hs view
@@ -0,0 +1,249 @@+{-# LANGUAGE BangPatterns #-}+{-# LANGUAGE ConstraintKinds #-}+{-# LANGUAGE FlexibleInstances #-}+{-# LANGUAGE GADTs #-}+{-# LANGUAGE InstanceSigs #-}+{-# LANGUAGE PatternGuards #-}+{-# LANGUAGE ScopedTypeVariables #-}+{-# LANGUAGE TemplateHaskell #-}+{-# LANGUAGE TypeFamilies #-}+{-# OPTIONS_GHC -fno-warn-orphans #-}+-- |+-- Module : Data.Array.Accelerate.CUDA.Array.Remote+-- Copyright : [2015..2016] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell, Robert Clifton-Everest+-- License : BSD3+--+-- Maintainer : Robert Clifton-Everest <robertce@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.CUDA.Array.Remote (++ -- Tables for host/device memory associations+ MemoryTable, R.PrimElt,+ new, malloc, withRemote, free, insertUnmanaged, reclaim++) where++import Control.Concurrent.MVar ( MVar, newMVar, withMVar, modifyMVar, readMVar )+import Control.Exception+import Control.Monad.IO.Class ( MonadIO, liftIO )+import Control.Monad.Trans.Reader+import Data.Functor+import Data.IntMap.Strict ( IntMap )+import Data.Proxy+import Data.Typeable ( Typeable )+import Foreign.Ptr ( ptrToIntPtr )+import Foreign.Storable+import Prelude hiding ( lookup )+import qualified Data.IntMap.Strict as IM++import Foreign.CUDA.Driver.Error+import qualified Foreign.CUDA.Driver as CUDA++import Data.Array.Accelerate.Array.Data+import Data.Array.Accelerate.Error+import Data.Array.Accelerate.Lifetime ( unsafeGetValue )+import qualified Data.Array.Accelerate.Array.Remote as R++import Data.Array.Accelerate.CUDA.Context ( Context(..), push, pop )+import Data.Array.Accelerate.CUDA.Execute.Event ( Event, EventTable, waypoint, query )+import Data.Array.Accelerate.CUDA.Execute.Stream ( Stream )+import qualified Data.Array.Accelerate.CUDA.Debug as Debug+++-- Instance for basic remote memory management functionality.+--+type CRM = ReaderT (Maybe Stream) IO++instance R.RemoteMemory CRM where+ type RemotePtr CRM = CUDA.DevicePtr+ --+ mallocRemote n = ReaderT $ \_ ->+ fmap Just (CUDA.mallocArray n)+ `catch` \e -> case e of+ ExitCode OutOfMemory -> return Nothing+ _ -> trace ("malloc failed with error: " ++ show e) (throwIO e)++ pokeRemote n dst ad = ReaderT $ \mst ->+ transfer "poke" (n * sizeOfPtr dst) $+ CUDA.pokeArrayAsync n (CUDA.HostPtr (ptrsOfArrayData ad)) dst mst++ peekRemote n src ad = ReaderT $ \mst ->+ transfer "peek" (n * sizeOfPtr src) $+ CUDA.peekArrayAsync n src (CUDA.HostPtr (ptrsOfArrayData ad)) mst++ castRemotePtr _ = CUDA.castDevPtr+ totalRemoteMem = ReaderT $ \_ -> snd <$> CUDA.getMemInfo+ availableRemoteMem = ReaderT $ \_ -> fst <$> CUDA.getMemInfo+ remoteAllocationSize = return 1024+++-- We leverage the memory cache from the accelerate base package.+--+-- However, we actually need multiple caches because every pointer has an+-- associated CUDA context. We could pair every DevicePtr with its context and+-- just have a single table, but the LRU implementation in the base package+-- assumes that remote pointers can be re-used, something that would not be true+-- for pointers allocated under different contexts.+--+type MT = IntMap (R.MemoryTable CUDA.DevicePtr Task)+data MemoryTable = MemoryTable {-# UNPACK #-} !EventTable+ {-# UNPACK #-} !(MVar MT)++type Task = Maybe Event++instance R.Task Task where+ completed Nothing = return True+ completed (Just e) = query e+++-- Create a MemoryTable.+--+new :: EventTable -> IO MemoryTable+new et = do+ message "initialise CUDA memory table"+ MemoryTable et <$> newMVar IM.empty+++-- Perform action on the device ptr that matches the given host-side array. Any+-- operations+--+withRemote+ :: forall e a b. R.PrimElt e a+ => Context+ -> MemoryTable+ -> ArrayData e+ -> (CUDA.DevicePtr a -> IO b)+ -> Maybe Stream+ -> IO (Maybe b)+withRemote ctx (MemoryTable et ref) ad run ms = do+ ct <- readMVar ref+ case IM.lookup (contextId ctx) ct of+ Nothing -> $internalError "withRemote" "context not found"+ Just mc -> streaming ms $ R.withRemote mc ad run'+ where+ run' :: R.RemotePtr CRM a -> CRM (Task, b)+ run' p = liftIO $ do+ c <- run p+ case ms of+ Nothing -> return (Nothing, c)+ Just s -> do+ e <- waypoint ctx et s+ return (Just e, c)+++-- Allocate a new device array to be associated with the given host-side array.+-- Has the same properties as `Data.Array.Accelerate.Array.Remote.LRU.malloc`+malloc :: forall a b. (Typeable a, R.PrimElt a b)+ => Context+ -> MemoryTable+ -> ArrayData a+ -> Bool+ -> Int+ -> IO Bool+malloc !ctx (MemoryTable _ !ref) !ad !frozen !n = do+ mt <- modifyMVar ref $ \ct -> blocking $ do+ case IM.lookup (contextId ctx) ct of+ Nothing -> trace "malloc/context not found" $ insertContext ctx ct+ Just mt -> return (ct, mt)+ blocking $ R.malloc mt ad frozen n+++-- Explicitly free an array in the LRU table. Has the same properties as+-- `Data.Array.Accelerate.Array.Remote.LRU.free`+--+free :: R.PrimElt a b+ => Context+ -> MemoryTable+ -> ArrayData a+ -> IO ()+free !ctx (MemoryTable _ !ref) !arr = withMVar ref $ \ct ->+ case IM.lookup (contextId ctx) ct of+ Nothing -> message "free/context not found"+ Just mt -> R.free (Proxy :: Proxy CRM) mt arr+++-- Record an association between a host-side array and a device memory area that was+-- not allocated by accelerate. The device memory will NOT be freed by the memory+-- manager.+--+insertUnmanaged+ :: R.PrimElt a b+ => Context+ -> MemoryTable+ -> ArrayData a+ -> CUDA.DevicePtr b+ -> IO ()+insertUnmanaged !ctx (MemoryTable _ !ref) !arr !ptr = do+ mt <- modifyMVar ref $ \ct -> blocking $ do+ case IM.lookup (contextId ctx) ct of+ Nothing -> trace "insertUnmanaged/context not found" $ insertContext ctx ct+ Just mt -> return (ct, mt)+ blocking $ R.insertUnmanaged mt arr ptr++insertContext+ :: Context+ -> MT+ -> CRM ( MT, R.MemoryTable CUDA.DevicePtr Task )+insertContext ctx ct = liftIO $ do+ mt <- R.new (\p -> bracket_ (push ctx) pop (CUDA.free p))+ return (IM.insert (contextId ctx) mt ct, mt)+++-- Removing entries+-- ----------------++-- Initiate garbage collection and finalise any arrays that have been marked as+-- unreachable.+--+reclaim :: MemoryTable -> IO ()+reclaim (MemoryTable _ ref) = withMVar ref (blocking . mapM_ R.reclaim . IM.elems)++-- Miscellaneous+-- -------------++{-# INLINE contextId #-}+contextId :: Context -> Int+contextId !ctx =+ let CUDA.Context !p = unsafeGetValue (deviceContext ctx)+ in fromIntegral (ptrToIntPtr p)++{-# INLINE blocking #-}+blocking :: CRM a -> IO a+blocking = flip runReaderT Nothing++{-# INLINE streaming #-}+streaming :: Maybe Stream -> CRM a -> IO a+streaming = flip runReaderT++{-# INLINE sizeOfPtr #-}+sizeOfPtr :: forall a. Storable a => CUDA.DevicePtr a -> Int+sizeOfPtr _ = sizeOf (undefined :: a)++-- Debug+-- -----++{-# INLINE trace #-}+trace :: MonadIO m => String -> m a -> m a+trace msg next = message msg >> next++{-# INLINE message #-}+message :: MonadIO m => String -> m ()+message msg = liftIO $ Debug.traceIO Debug.dump_gc ("gc/lru: " ++ msg)++{-# INLINE showBytes #-}+showBytes :: Int -> String+showBytes x = Debug.showFFloatSIBase (Just 0) 1024 (fromIntegral x :: Double) "B"++{-# INLINE transfer #-}+transfer :: String -> Int -> IO () -> IO ()+transfer name bytes action+ = let showRate x t = Debug.showFFloatSIBase (Just 3) 1024 (fromIntegral x / t) "B/s"+ msg gpuTime cpuTime = "gc/lru: " ++ name ++ ": "+ ++ showBytes bytes ++ " @ " ++ showRate bytes gpuTime ++ ", "+ ++ Debug.elapsed gpuTime cpuTime+ in+ Debug.timed Debug.dump_gc msg Nothing action+
+ Data/Array/Accelerate/CUDA/Array/Slice.hs view
@@ -0,0 +1,212 @@+{-# LANGUAGE GADTs #-}+-- |+-- Module : Data.Array.Accelerate.Array.Slice+--+-- Maintainer : Frederik Meisner Madsen <fmma@diku.dk>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.CUDA.Array.Slice (++ transferDesc, blocksOf,++ TransferDesc(..)++) where++import Control.Arrow ( first )+import Data.List ( groupBy, elemIndex )+import Data.Array.Accelerate.Array.Representation ( SliceIndex(..) )+++-- Slicing algorithm.++-- Figure out how to transfer a multi-dimensional slice from a+-- multi-dimensional array of the same or higher dimensionality to a+-- contiguous memory region. The algorithm uses as few data transfers+-- as possible, where a linear memory copy with a stride consitutes+-- one transfer (cudaMemcpy2d). Each transfer is described by four+-- numbers: A start, a stride, a number of basic blocks and block+-- size.+--++-- Memory transfer(s) description suitable for (multiple calls to)+-- cudaMemcpu2D. The number of calls necessary is given by the length+-- of the starting offsets. The unit is basic elements of the target+-- array.+data TransferDesc =+ TransferDesc { starts :: [Int] -- Starting offsets.+ , stride :: Int -- Stride.+ , nblocks :: Int -- Number of blocks.+ , blocksize :: Int } -- Elements per block.+ deriving Show++-- Get the basic blocks of the given transfer descriptions, described+-- as offset and length. This function can be used to do a memory copy+-- with only contiguous data transfers, if a strided transfer is not+-- available.+blocksOf :: TransferDesc -> [(Int, Int, Int)]+blocksOf tdesc =+ [ ( start + i * stride tdesc+ , i * blocksize tdesc+ , blocksize tdesc)+ | start <- starts tdesc+ , i <- [0..nblocks tdesc-1]]++-- A slice index in one dimension, annotated with information used to+-- catogorize how the dimension is transferred.+type SliceIR = ( TransferType -- Transfer type annotation.+ , Int ) -- Size of this dimension.++data TransferType =+ Strided -- Transfer the entire dimension. Each element+ -- will be transfered in a different block, but in+ -- a single transfer.++ | Contiguous -- Transfer the entire dimension in one+ -- contiguous data transfer. Each element will+ -- be transfered in the same block in the same+ -- transfer.++ | Fixed Int -- Transfer a specific element in this+ -- dimension. The argument is the element index.++ | FixedAll -- Transfer the entire dimension. Each element+ -- will be transfered in a seperate data+ -- transfer.++-- Convert a slice index to internal representation used by this+-- algorithm. Initially, all dimensions are represented as strided+-- (transfer entire dimension) or fixed (transfer specific element).++toIR :: SliceIndex slix sl co dim+ -> slix+ -> dim+ -> [SliceIR]+toIR SliceNil () () = []+toIR (SliceAll si) (sl, ()) (sh, n) = (Strided, n):toIR si sl sh+toIR (SliceFixed si) (sl, i ) (sh, n) = (Fixed i, n):toIR si sl sh+{-+toIR :: SliceIndex slix sl co dim+ -> slix+ -> dim+ -> [SliceIR]+toIR slix sl dim =+ f slix sl dim []+ where+ f :: SliceIndex slix' sl' co' dim'+ -> slix'+ -> dim'+ -> [SliceIR]+ -> [SliceIR]+ f SliceNil () () res = res+ f (SliceAll si) (sl, ()) (sh, n) res = f si sl sh ((Strided, n):res)+ f (SliceFixed si) (sl, i ) (sh, n) res = f si sl sh ((Fixed i, n):res)+-}++-- Promote strided slice indices to contiguous slice indices when+-- possible (= all strided innermost dimensions).+strideToContiguous :: [SliceIR] -> [SliceIR]+strideToContiguous slirs =+ let (strided, rest) = span (isStrided . fst) slirs+ conti = map (first promote) strided+ in conti ++ rest+ where+ isStrided Strided = True+ isStrided _ = False++ promote Strided = Contiguous+ promote x = x++-- Find largest group (= next to each other) of strided slice+-- indices. Promote all other strided groups to fixed groups, which+-- entails multiple data transfers in these dimensions. This is+-- necessary step, since it is impossible to describe the memory+-- transfer with the given transfer description type in some+-- situations (namely when there is a gap between strided+-- dimensions). This step chooses the split that results in fewest+-- possible data transfers by selecting the largest stride pivot.+selectStridePivot :: [SliceIR] -> [SliceIR]+selectStridePivot slirs =+ let -- Group the slice dimensions in groups of same transfer types.+ groups = groupBy sameTransferType slirs++ -- Calculate the sizes of the strided groups, ignore the rest.+ sizes = map (product . map stridedSize) groups++ -- Find the largest strided group.+ Just imax = elemIndex (maximum sizes) sizes++ -- Promote the remaining strided groups to fixed groups using+ -- multiple data transfers.+ (groups1, pivotGroup:groups2) = splitAt imax groups++ groups' = map (map (first promote)) groups1 ++ pivotGroup:map (map (first promote)) groups2++ in concat groups'+ where+ sameTransferType :: SliceIR -> SliceIR -> Bool+ sameTransferType x y = fst x ~= fst y+ where+ (~=) :: TransferType -> TransferType -> Bool+ Contiguous ~= Contiguous = True+ Strided ~= Strided = True+ Fixed _ ~= Fixed _ = True+ FixedAll ~= FixedAll = True+ _ ~= _ = False++ promote :: TransferType -> TransferType+ promote Strided = FixedAll+ promote x = x++ stridedSize :: SliceIR -> Int+ stridedSize (Strided, n) = n+ stridedSize _ = 0++-- Compute a description of the transfers necessary to copy the+-- specified slice.+transferDesc :: SliceIndex slix sl co dim+ -> slix -- Slice index.+ -> dim -- Full shape.+ -> TransferDesc+transferDesc slix sl dim =+ let slirs = selectStridePivot $ strideToContiguous $ toIR slix sl dim+ tdesc0 = TransferDesc [0] 1 1 1+ size0 = 1+ in f slirs size0 tdesc0+ where+ f :: [SliceIR] -> Int -> TransferDesc -> TransferDesc+ f slirs m tdesc =+ case slirs of+ [] -> tdesc++ (ttyp, n):slirs' -> f slirs' (n * m) $+ case ttyp of+ Strided ->+ TransferDesc+ { starts = starts tdesc+ , stride = if stride tdesc == 1 then m else stride tdesc+ , nblocks = n * nblocks tdesc+ , blocksize = blocksize tdesc }++ Contiguous ->+ TransferDesc+ { starts = starts tdesc+ , stride = stride tdesc+ , nblocks = nblocks tdesc+ , blocksize = n * blocksize tdesc }++ Fixed i ->+ TransferDesc+ { starts = [start + i * m | start <- starts tdesc]+ , stride = stride tdesc+ , nblocks = nblocks tdesc+ , blocksize = blocksize tdesc }++ FixedAll ->+ TransferDesc+ { starts = [start + i * m | start <- starts tdesc, i <- [0..n-1]]+ , stride = stride tdesc+ , nblocks = nblocks tdesc+ , blocksize = blocksize tdesc }
Data/Array/Accelerate/CUDA/Array/Sugar.hs view
@@ -12,7 +12,7 @@ module Data.Array.Accelerate.CUDA.Array.Sugar ( module Data.Array.Accelerate.Array.Sugar,- newArray, allocateArray, useArray, useArrayAsync,+ fromFunction, allocateArray, useArray, useArrayAsync, ) where @@ -20,16 +20,16 @@ import Data.Array.Accelerate.CUDA.State import Data.Array.Accelerate.CUDA.Array.Data-import Data.Array.Accelerate.Array.Sugar hiding (newArray, allocateArray)+import Data.Array.Accelerate.Array.Sugar hiding ( fromFunction, allocateArray ) import qualified Data.Array.Accelerate.Array.Sugar as Sugar -- Create an array from its representation function, uploading the result to the -- device ---newArray :: (Shape sh, Elt e) => sh -> (sh -> e) -> CIO (Array sh e)-newArray sh f =- let arr = Sugar.newArray sh f+fromFunction :: (Shape sh, Elt e) => sh -> (sh -> e) -> CIO (Array sh e)+fromFunction sh f =+ let arr = Sugar.fromFunction sh f in do useArray arr return arr
− Data/Array/Accelerate/CUDA/Array/Table.hs
@@ -1,303 +0,0 @@-{-# LANGUAGE BangPatterns #-}-{-# LANGUAGE GADTs #-}-{-# LANGUAGE PatternGuards #-}-{-# LANGUAGE ScopedTypeVariables #-}-{-# LANGUAGE TemplateHaskell #-}--- |--- Module : Data.Array.Accelerate.CUDA.Array.Table--- 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>--- Stability : experimental--- Portability : non-portable (GHC extensions)-----module Data.Array.Accelerate.CUDA.Array.Table (-- -- Tables for host/device memory associations- MemoryTable, new, lookup, malloc, free, insert, insertRemote, reclaim--) where--import Control.Applicative-import Control.Concurrent ( yield )-import Control.Concurrent.MVar ( MVar, newMVar, withMVar, mkWeakMVar )-import Control.Exception ( bracket_, catch, throwIO )-import Control.Monad ( unless )-import Data.Hashable ( Hashable(..) )-import Data.Maybe ( isJust )-import Data.Typeable ( Typeable, gcast )-import System.Mem ( performGC )-import System.Mem.StableName ( StableName, makeStableName, hashStableName )-import System.Mem.Weak ( Weak, mkWeak, deRefWeak, finalize )-import Prelude hiding ( lookup )--import Foreign.Ptr ( ptrToIntPtr )-import Foreign.Storable ( Storable, sizeOf )-import Foreign.CUDA.Ptr ( DevicePtr )-import Foreign.CUDA.Driver.Error-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----- 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.------ It is important that we can garbage collect old entries from the table when--- the key is no longer reachable in the heap. Hence the value part of each--- table entry is a (Weak val), where the stable name 'key' is the key for the--- memo table, and the 'val' is the value of this table entry. When the key--- becomes unreachable, a finaliser will fire and remove this entry from the--- hash buckets, and further attempts to dereference the weak pointer will--- return Nothing. References from 'val' to the key are ignored (see the--- semantics of weak pointers in the documentation).----type HashTable key val = HT.BasicHashTable key val-type MT = MVar ( HashTable HostArray DeviceArray )-data MemoryTable = MemoryTable {-# UNPACK #-} !MT- {-# UNPACK #-} !(Weak MT)- {-# UNPACK #-} !Nursery---- Arrays on the host and device----type ContextId = Int--data HostArray where- HostArray :: Typeable e- => {-# UNPACK #-} !ContextId -- unique ID relating to the parent context- -> {-# UNPACK #-} !(StableName (ArrayData e))- -> HostArray--data DeviceArray where- DeviceArray :: Typeable e- => {-# UNPACK #-} !(Weak (DevicePtr e))- -> DeviceArray--instance Eq HostArray where- HostArray _ a1 == HostArray _ a2- = maybe False (== a2) (gcast a1)--instance Hashable HostArray where- {-# INLINE hashWithSalt #-}- hashWithSalt salt (HostArray cid sn)- = salt `hashWithSalt` cid `hashWithSalt` sn--instance Show HostArray where- show (HostArray _ sn) = "Array #" ++ show (hashStableName sn)----- Referencing arrays--- ---------------------- Create a new hash table from host to device arrays. When the structure is--- collected it will finalise all entries in the table.----new :: IO MemoryTable-new = do- message "initialise memory table"- tbl <- HT.new- ref <- newMVar tbl- nrs <- N.new- weak <- mkWeakMVar ref (table_finalizer tbl)- return $! MemoryTable ref weak nrs----- Look for the device memory corresponding to a given host-side array.----lookup :: (Typeable a, Typeable b) => Context -> MemoryTable -> ArrayData a -> IO (Maybe (DevicePtr b))-lookup ctx (MemoryTable !ref _ _) !arr = do- sa <- makeStableArray ctx arr- mw <- withMVar ref (`HT.lookup` sa)- case mw of- Nothing -> trace ("lookup/not found: " ++ show sa) $ return Nothing- Just (DeviceArray w) -> do- mv <- deRefWeak w- case mv of- Just v | Just p <- gcast v -> trace ("lookup/found: " ++ show sa) $ return (Just p)- | otherwise -> $internalError "lookup" $ "type mismatch"-- -- Note: [Weak pointer weirdness]- --- -- After the lookup is successful, there might conceivably be no further- -- references to 'arr'. If that is so, and a garbage collection- -- intervenes, the weak pointer might get tombstoned before 'deRefWeak'- -- gets to it. In that case we throw an error (below). However, because- -- we have used 'arr' in the continuation, this ensures that 'arr' is- -- reachable in the continuation of 'deRefWeak' and thus 'deRefWeak'- -- always succeeds. This sort of weirdness, typical of the world of weak- -- pointers, is why we can not reuse the stable name 'sa' computed- -- above in the error message.- --- Nothing ->- 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.--- This will attempt to use an old array from the nursery, but will otherwise--- allocate fresh data.------ Instead of allocating the exact number of elements requested, we round up to--- a fixed chunk size; currently set at 128 elements. This means there is a--- greater chance the nursery will get a hit, and moreover that we can search--- the nursery for an exact size. TLM: I believe the CUDA API allocates in--- chunks, of size 4MB.----malloc :: forall a b. (Typeable a, Typeable b, Storable b) => Context -> MemoryTable -> ArrayData a -> Int -> IO (DevicePtr b)-malloc !ctx mt@(MemoryTable _ _ !nursery) !ad !n = do- let -- next highest multiple of f from x- multiple x f = floor ((x + (f-1)) / f :: Double)- chunk = 1024-- !n' = chunk * multiple (fromIntegral n) (fromIntegral chunk)- !bytes = n' * sizeOf (undefined :: b)- --- mp <- N.malloc bytes (deviceContext ctx) nursery- ptr <- case mp of- Just p -> trace "malloc/nursery" $ return (CUDA.castDevPtr p)- Nothing -> trace "malloc/new" $- CUDA.mallocArray n' `catch` \(e :: CUDAException) ->- case e of- ExitCode OutOfMemory -> reclaim mt >> CUDA.mallocArray n'- _ -> throwIO e- insert ctx mt ad ptr bytes- return ptr----- Deallocate the device array associated with the given host-side array. This--- calls the finaliser for that array immediately, regardless of the current (or--- future) use status of that array.----free :: Typeable a => Context -> MemoryTable -> ArrayData a -> IO ()-free !ctx (MemoryTable !ref _ _) !arr = do- sa <- makeStableArray ctx arr- mw <- withMVar ref (`HT.lookup` sa)- case mw of- Nothing -> message ("free/not found: " ++ show sa)- Just (DeviceArray w) -> trace ("free/evict: " ++ show sa) $ finalize w----- Record an association between a host-side array and a new device memory area.--- The device memory will be freed when the host array is garbage collected.----insert :: (Typeable a, Typeable b) => Context -> MemoryTable -> ArrayData a -> DevicePtr b -> Int -> IO ()-insert !ctx (MemoryTable !ref !weak_ref (Nursery _ !weak_nrs)) !arr !ptr !bytes = do- key <- makeStableArray ctx arr- dev <- DeviceArray `fmap` mkWeak arr ptr (Just $ finalizer (weakContext ctx) weak_ref weak_nrs key ptr bytes)- message $ "insert: " ++ show key- withMVar ref $ \tbl -> HT.insert tbl key dev----- Record an association between a host-side array and a device memory area that was--- not allocated by accelerate. The device memory will NOT be freed when the host--- array is garbage collected.----insertRemote :: (Typeable a, Typeable b) => Context -> MemoryTable -> ArrayData a -> DevicePtr b -> IO ()-insertRemote !ctx (MemoryTable !ref !weak_ref _) !arr !ptr = do- key <- makeStableArray ctx arr- dev <- DeviceArray `fmap` mkWeak arr ptr (Just $ remoteFinalizer weak_ref key)- message $ "insert/remote: " ++ show key- withMVar ref $ \tbl -> HT.insert tbl key dev----- Removing entries--- -------------------- Initiate garbage collection and finalise any arrays that have been marked as--- unreachable.----reclaim :: MemoryTable -> IO ()-reclaim (MemoryTable _ weak_ref (Nursery nrs _)) = do- (before, total) <- CUDA.getMemInfo- performGC- yield- withMVar nrs N.flush- mr <- deRefWeak weak_ref- case mr of- Nothing -> return ()- Just ref -> withMVar ref $ \tbl ->- flip HT.mapM_ tbl $ \(_,DeviceArray w) -> do- alive <- isJust `fmap` deRefWeak w- unless alive $ finalize w- --- D.when D.dump_gc $ do- (after, _) <- CUDA.getMemInfo- message $ "reclaim: freed " ++ showBytes (fromIntegral (before - after))- ++ ", " ++ showBytes (fromIntegral after)- ++ " of " ++ showBytes (fromIntegral total) ++ " remaining"---- Because a finaliser might run at any time, we must reinstate the context in--- which the array was allocated before attempting to release it.------ Note also that finaliser threads will silently terminate if an exception is--- raised. If the context, and thereby all allocated memory, was destroyed--- externally before the thread had a chance to run, all we need do is update--- the hash tables --- but we must do this first before failing to use a dead--- context.----finalizer :: Weak CUDA.Context -> Weak MT -> Weak NRS -> HostArray -> DevicePtr b -> Int -> IO ()-finalizer !weak_ctx !weak_ref !weak_nrs !key !ptr !bytes = do- mr <- deRefWeak weak_ref- case mr of- Nothing -> message ("finalise/dead table: " ++ show key)- Just ref -> withMVar ref (`HT.delete` key)- --- mc <- deRefWeak weak_ctx- case mc of- Nothing -> message ("finalise/dead context: " ++ show key)- Just ctx -> do- --- mn <- deRefWeak weak_nrs- case mn of- Nothing -> trace ("finalise/free: " ++ show key) $ bracket_ (CUDA.push ctx) CUDA.pop (CUDA.free ptr)- Just nrs -> trace ("finalise/nursery: " ++ show key) $ N.stash bytes ctx nrs ptr--remoteFinalizer :: Weak MT -> HostArray -> IO ()-remoteFinalizer !weak_ref !key = do- mr <- deRefWeak weak_ref- case mr of- Nothing -> message ("finalise/dead table: " ++ show key)- Just ref -> trace ("finalise: " ++ show key) $ withMVar ref (`HT.delete` key)--table_finalizer :: HashTable HostArray DeviceArray -> IO ()-table_finalizer !tbl- = trace "table finaliser"- $ HT.mapM_ (\(_,DeviceArray w) -> finalize w) tbl----- Miscellaneous--- ---------------{-# INLINE makeStableArray #-}-makeStableArray :: Typeable a => Context -> ArrayData a -> IO HostArray-makeStableArray !ctx !arr =- let CUDA.Context !p = deviceContext ctx- !cid = fromIntegral (ptrToIntPtr p)- in- HostArray cid <$> makeStableName arr----- Debug--- -------{-# INLINE showBytes #-}-showBytes :: Int -> String-showBytes x = D.showFFloatSIBase (Just 0) 1024 (fromIntegral x :: Double) "B"--{-# INLINE trace #-}-trace :: String -> IO a -> IO a-trace msg next = D.message D.dump_gc ("gc: " ++ msg) >> next--{-# INLINE message #-}-message :: String -> IO ()-message s = s `trace` return ()-
− Data/Array/Accelerate/CUDA/Async.hs
@@ -1,56 +0,0 @@-{-# LANGUAGE CPP #-}--- |--- Module : Data.Array.Accelerate.CUDA.Async--- Copyright : [2009..2014] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell--- License : BSD3------ Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>--- Stability : experimental--- Portability : non-portable (GHC extensions)-----module Data.Array.Accelerate.CUDA.Async- where--import Control.Exception-import Control.Concurrent----- We need to execute the main thread asynchronously to give finalisers a chance--- to run. Make sure to catch exceptions to avoid "blocked indefinitely on MVar"--- errors.----data Async a = Async {-# UNPACK #-} !ThreadId- {-# UNPACK #-} !(MVar (Either SomeException a))---- | Fork an action to execute asynchronously.----async :: IO a -> IO (Async a)-async action = do- var <- newEmptyMVar- tid <- forkIO $ (putMVar var . Right =<< action)- `catch`- \e -> putMVar var (Left e)- return (Async tid var)---- | Block the calling thread until the computation completes, then return the--- result.----{-# INLINE wait #-}-wait :: Async a -> IO a-wait (Async _ var) = either throwIO return =<< readMVar var---- | Test whether the asynchronous computation has already completed. If so,--- return the result, else 'Nothing'.----{-# INLINE poll #-}-poll :: Async a -> IO (Maybe a)-poll (Async _ var) =- maybe (return Nothing) (either throwIO (return . Just)) =<< tryTakeMVar var---- | Cancel a running asynchronous computation.----{-# INLINE cancel #-}-cancel :: Async a -> IO ()-cancel (Async tid _) = throwTo tid ThreadKilled-
Data/Array/Accelerate/CUDA/CodeGen.hs view
@@ -4,6 +4,7 @@ {-# LANGUAGE RecordWildCards #-} {-# LANGUAGE ScopedTypeVariables #-} {-# LANGUAGE TemplateHaskell #-}+{-# LANGUAGE ViewPatterns #-} {-# OPTIONS -fno-warn-name-shadowing #-} -- | -- Module : Data.Array.Accelerate.CUDA.CodeGen@@ -18,13 +19,11 @@ module Data.Array.Accelerate.CUDA.CodeGen ( - CUTranslSkel, codegenAcc,+ CUTranslSkel, codegenAcc, codegenToSeq ) where -- libraries-import Data.Loc-import Data.Char import Data.HashSet ( HashSet ) import Control.Monad.State.Strict import Foreign.CUDA.Analysis@@ -37,16 +36,18 @@ -- friends import Data.Array.Accelerate.Error import Data.Array.Accelerate.Type-import Data.Array.Accelerate.Tuple+import Data.Array.Accelerate.Product import Data.Array.Accelerate.Trafo import Data.Array.Accelerate.Pretty () import Data.Array.Accelerate.Analysis.Shape-import Data.Array.Accelerate.Array.Sugar ( Array, Shape, Elt, EltRepr )+import Data.Array.Accelerate.Array.Sugar ( Array, Shape, Elt, EltRepr+ , Tuple(..), TupleRepr ) import Data.Array.Accelerate.Array.Representation ( SliceIndex(..) ) import qualified Data.Array.Accelerate.Array.Sugar as Sugar import qualified Data.Array.Accelerate.Analysis.Type as Sugar import Data.Array.Accelerate.CUDA.AST hiding ( Val(..), prj )+import Data.Array.Accelerate.CUDA.CodeGen.Constant import Data.Array.Accelerate.CUDA.CodeGen.Base import Data.Array.Accelerate.CUDA.CodeGen.Type import Data.Array.Accelerate.CUDA.CodeGen.Monad@@ -55,7 +56,9 @@ import Data.Array.Accelerate.CUDA.CodeGen.PrefixSum import Data.Array.Accelerate.CUDA.CodeGen.Reduction import Data.Array.Accelerate.CUDA.CodeGen.Stencil+import Data.Array.Accelerate.CUDA.CodeGen.Streaming import Data.Array.Accelerate.CUDA.Foreign.Import ( canExecuteExp )+import qualified Data.Array.Accelerate.CUDA.CodeGen.Arithmetic as A -- Local environments@@ -112,6 +115,9 @@ Stencil f b a -> mkStencil dev aenv <$> travF1 f <*> travB a b Stencil2 f b1 a1 b2 a2 -> mkStencil2 dev aenv <$> travF2 f <*> travB a1 b1 <*> travB a2 b2 + -- Sequence collection+ -- Collect _ -> unexpectedError+ -- Non-computation forms -> sadness Alet{} -> unexpectedError Avar{} -> unexpectedError@@ -163,7 +169,7 @@ travB _ Clamp = return Clamp travB _ Mirror = return Mirror travB _ Wrap = return Wrap- travB _ (Constant c) = return . Constant $ CUExp ([], codegenConst (Sugar.eltType (undefined::e)) c)+ travB _ (Constant c) = return . Constant $ CUExp ([], constant (Sugar.eltType (undefined::e)) c) -- caffeine and misery prim :: String@@ -171,7 +177,39 @@ unexpectedError = $internalError "codegenAcc" $ "unexpected array primitive: " ++ prim fusionError = $internalError "codegenAcc" $ "unexpected fusible material: " ++ prim +codegenToSeq :: forall aenv slix sl co sh e. (Shape sl, Shape sh, Elt e)+ => SliceIndex slix+ (EltRepr sl)+ co+ (EltRepr sh)+ -> DeviceProperties+ -> DelayedOpenAcc aenv (Array sh e)+ -> Gamma aenv+ -> CUTranslSkel aenv (Array sl e)+codegenToSeq slix dev acc aenv = codegen $ (mkToSeq slix dev aenv <$> travD acc)+ where+ codegen :: CUDA (CUTranslSkel aenv (Array sl e)) -> CUTranslSkel aenv (Array sl e)+ codegen cuda =+ let (skeleton, st) = runCUDA cuda+ addTo (CUTranslSkel name code) =+ CUTranslSkel name (Set.foldr (\h c -> [cedecl| $esc:("#include \"" ++ h ++ "\"") |] : c) code (headers st))+ in+ addTo skeleton + -- code generation for delayed arrays+ travD :: (Shape sh, Elt e) => DelayedOpenAcc aenv (Array sh e) -> CUDA (CUDelayedAcc aenv sh e )+ travD Manifest{} = $internalError "codegenAcc" "expected delayed array"+ travD Delayed{..} = CUDelayed <$> travE extentD+ <*> travF1 indexD+ <*> travF1 linearIndexD++ travE :: forall t. DelayedExp aenv t -> CUDA (CUExp aenv t)+ travE = codegenExp dev aenv++ travF1 :: forall a b. DelayedFun aenv (a -> b) -> CUDA (CUFun1 aenv (a -> b))+ travF1 = codegenFun1 dev aenv++ -- Scalar function abstraction -- --------------------------- @@ -313,9 +351,9 @@ case exp of Let bnd body -> elet bnd body env Var ix -> return $ prj ix env- PrimConst c -> return $ [codegenPrimConst c]- Const c -> return $ codegenConst (Sugar.eltType (undefined::t)) c- PrimApp f x -> return <$> primApp f x env+ PrimConst c -> return $ [primConst c]+ Const c -> return $ constant (Sugar.eltType (undefined::t)) c+ PrimApp f x -> 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@@ -338,6 +376,7 @@ Shape acc -> shape acc env ShapeSize sh -> shapeSize sh env Intersect sh1 sh2 -> intersect sh1 sh2 env+ Union sh1 sh2 -> union sh1 sh2 env --Foreign function Foreign ff _ e -> foreignE ff e env@@ -364,13 +403,77 @@ -- 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+ primApp :: PrimFun (a -> b) -> DelayedOpenExp env aenv a -> Val env -> Gen [C.Exp]+ primApp f x env =+ case f of+ -- operators from Num+ PrimAdd{} -> binary A.add x env+ PrimSub{} -> binary A.sub x env+ PrimMul{} -> binary A.mul x env+ PrimNeg{} -> unary A.negate x env+ PrimAbs ty -> unary (A.abs ty) x env+ PrimSig ty -> unaryM (A.signum ty) x env+ -- operators from Integral & Bits+ PrimQuot{} -> binary A.quot x env+ PrimRem{} -> binary A.rem x env+ PrimQuotRem ty -> binaryM2 (A.quotRem ty) x env+ PrimIDiv ty -> binaryM (A.idiv ty) x env+ PrimMod ty -> binaryM (A.mod ty) x env+ PrimDivMod ty -> binaryM2 (A.divMod ty) x env+ PrimBAnd{} -> binary A.band x env+ PrimBOr{} -> binary A.bor x env+ PrimBXor{} -> binary A.xor x env+ PrimBNot{} -> unary A.bnot x env+ PrimBShiftL{} -> binary A.shiftL x env+ PrimBShiftR{} -> binary A.shiftRA x env+ PrimBRotateL ty -> binaryM (A.rotateL ty) x env+ PrimBRotateR ty -> binaryM (A.rotateR ty) x env+ -- operators from Fractional and Floating+ PrimFDiv{} -> binary A.fdiv x env+ PrimRecip ty -> unary (A.recip ty) x env+ PrimSin ty -> unary (A.sin ty) x env+ PrimCos ty -> unary (A.cos ty) x env+ PrimTan ty -> unary (A.tan ty) x env+ PrimAsin ty -> unary (A.asin ty) x env+ PrimAcos ty -> unary (A.acos ty) x env+ PrimAtan ty -> unary (A.atan ty) x env+ PrimSinh ty -> unary (A.sinh ty) x env+ PrimCosh ty -> unary (A.cosh ty) x env+ PrimTanh ty -> unary (A.tanh ty) x env+ PrimAsinh ty -> unary (A.asinh ty) x env+ PrimAcosh ty -> unary (A.acosh ty) x env+ PrimAtanh ty -> unary (A.atanh ty) x env+ PrimExpFloating ty -> unary (A.exp ty) x env+ PrimSqrt ty -> unary (A.sqrt ty) x env+ PrimLog ty -> unary (A.log ty) x env+ PrimFPow ty -> binary (A.pow ty) x env+ PrimLogBase ty -> binary (A.logBase ty) x env+ -- operators from RealFrac+ PrimTruncate ta tb -> unary (A.trunc ta tb) x env+ PrimRound ta tb -> unary (A.round ta tb) x env+ PrimFloor ta tb -> unary (A.floor ta tb) x env+ PrimCeiling ta tb -> unary (A.ceiling ta tb) x env+ -- operators from RealFloat+ PrimAtan2 ty -> binary (A.atan2 ty) x env+ PrimIsNaN{} -> unary A.isNaN x env+ -- relational and equality operators+ PrimLt{} -> binary A.lt x env+ PrimGt{} -> binary A.gt x env+ PrimLtEq{} -> binary A.leq x env+ PrimGtEq{} -> binary A.geq x env+ PrimEq{} -> binary A.eq x env+ PrimNEq{} -> binary A.neq x env+ PrimMax ty -> binary (A.max ty) x env+ PrimMin ty -> binary (A.min ty) x env+ -- logical operators+ PrimLAnd -> binary A.land x env+ PrimLOr -> binary A.lor x env+ PrimLNot -> unary A.lnot x env+ -- type conversions+ PrimOrd -> unary A.ord x env+ PrimChr -> unary A.chr x env+ PrimBoolToInt -> unary A.boolToInt x env+ PrimFromIntegral ta tb -> unary (A.fromIntegral ta tb) x env where cvtE' :: DelayedOpenExp env aenv a -> Val env -> Gen C.Exp cvtE' e env = do@@ -379,6 +482,36 @@ then return r else return [cexp| ({ $items:b; $exp:r; }) |] + -- TLM: This is a bit ugly. Consider making all primitive functions from+ -- Arithmetic.hs evaluate in the Gen monad.+ --+ unary :: (C.Exp -> C.Exp) -> DelayedOpenExp env aenv a -> Val env -> Gen [C.Exp]+ unary f = unaryM (return . f)++ unaryM :: (C.Exp -> Gen C.Exp) -> DelayedOpenExp env aenv a -> Val env -> Gen [C.Exp]+ unaryM f a env = do+ a' <- cvtE' a env+ r <- f a'+ return [r]++ binary :: (C.Exp -> C.Exp -> C.Exp) -> DelayedOpenExp env aenv (a,b) -> Val env -> Gen [C.Exp]+ binary f = binaryM (\a b -> return (f a b))++ binaryM :: (C.Exp -> C.Exp -> Gen C.Exp) -> DelayedOpenExp env aenv (a,b) -> Val env -> Gen [C.Exp]+ binaryM f x env = do+ x' <- cvtE x env+ case x' of+ [a,b] -> return <$> f a b+ _ -> $internalError "primApp" "unexpected argument to binary function"++ binaryM2 :: (C.Exp -> C.Exp -> Gen (C.Exp, C.Exp)) -> DelayedOpenExp env aenv (a,b) -> Val env -> Gen [C.Exp]+ binaryM2 f x env = do+ x' <- cvtE x env+ case x' of+ [a,b] -> do (r,s) <- f a b+ return [r,s]+ _ -> $internalError "primApp" "unexpected argument to binary function"+ -- 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.@@ -429,7 +562,8 @@ => DelayedOpenExp env aenv Bool -> DelayedOpenExp env aenv t -> DelayedOpenExp env aenv t- -> Val env -> Gen [C.Exp]+ -> Val env+ -> Gen [C.Exp] cond p t f env = do p' <- cvtE p env ok <- single "Cond" <$> pushEnv p p'@@ -641,24 +775,34 @@ intersect sh1 sh2 env = zipWith (\a b -> ccall "min" [a,b]) <$> cvtE sh1 env <*> cvtE sh2 env + -- Union of two shapes, taken as the maximum in each dimension.+ --+ union :: forall env sh. Elt sh+ => DelayedOpenExp env aenv sh+ -> DelayedOpenExp env aenv sh+ -> Val env -> Gen [C.Exp]+ union sh1 sh2 env =+ zipWith (\a b -> ccall "max" [a,b]) <$> cvtE sh1 env <*> cvtE sh2 env+ -- Foreign scalar functions. We need to extract any header files that might -- be required so they can be added to the top level definitions. -- -- Additionally, we insert an explicit type cast from the foreign function -- result back into Accelerate types (c.f. Int vs int). --- foreignE :: forall f a b env. (Sugar.Foreign f, Elt a, Elt b)- => f a b+ foreignE :: forall asm a b env. (Sugar.Foreign asm, Elt a, Elt b)+ => asm (a -> b) -> DelayedOpenExp env aenv a -> Val env -> Gen [C.Exp]- foreignE ff x env = case canExecuteExp ff of- 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)]+ foreignE ff x env =+ case canExecuteExp ff of+ Nothing -> $internalError "codegenOpenExp" "failed to recover foreign function a second time"+ 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.@@ -677,264 +821,21 @@ single loc _ = $internalError loc "expected single expression" --- Scalar Primitives--- -------------------codegenPrimConst :: PrimConst a -> C.Exp-codegenPrimConst (PrimMinBound ty) = codegenMinBound ty-codegenPrimConst (PrimMaxBound ty) = codegenMaxBound ty-codegenPrimConst (PrimPi ty) = codegenPi ty---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"--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----codegenConst :: TupleType a -> a -> [C.Exp]-codegenConst UnitTuple _ = []-codegenConst (SingleTuple ty) c = [codegenScalar ty c]-codegenConst (PairTuple ty1 ty0) (cs,c) = codegenConst ty1 cs ++ codegenConst ty0 c----- Scalar constants----codegenScalar :: ScalarType a -> a -> C.Exp-codegenScalar (NumScalarType ty) = codegenNumScalar ty-codegenScalar (NonNumScalarType ty) = codegenNonNumScalar ty--codegenNumScalar :: NumType a -> a -> C.Exp-codegenNumScalar (IntegralNumType ty) = codegenIntegralScalar ty-codegenNumScalar (FloatingNumType ty) = codegenFloatingScalar ty--codegenIntegralScalar :: IntegralType a -> a -> C.Exp-codegenIntegralScalar ty x | IntegralDict <- integralDict ty = [cexp| ( $ty:(codegenIntegralType ty) ) $exp:(cintegral x) |]--codegenFloatingScalar :: FloatingType a -> a -> C.Exp-codegenFloatingScalar (TypeFloat _) x = C.Const (C.FloatConst (shows x "f") (toRational x) noLoc) noLoc-codegenFloatingScalar (TypeCFloat _) x = C.Const (C.FloatConst (shows x "f") (toRational x) noLoc) noLoc-codegenFloatingScalar (TypeDouble _) x = C.Const (C.DoubleConst (show x) (toRational x) noLoc) noLoc-codegenFloatingScalar (TypeCDouble _) x = C.Const (C.DoubleConst (show x) (toRational x) noLoc) noLoc--codegenNonNumScalar :: NonNumType a -> a -> C.Exp-codegenNonNumScalar (TypeBool _) x = cbool x-codegenNonNumScalar (TypeChar _) x = [cexp|$char:x|]-codegenNonNumScalar (TypeCChar _) x = [cexp|$char:(chr (fromIntegral x))|]-codegenNonNumScalar (TypeCUChar _) x = [cexp|$char:(chr (fromIntegral x))|]-codegenNonNumScalar (TypeCSChar _) x = [cexp|$char:(chr (fromIntegral x))|]----- Constant methods of floating----codegenPi :: FloatingType a -> C.Exp-codegenPi ty | FloatingDict <- floatingDict ty = codegenFloatingScalar ty pi----- Constant methods of bounded----codegenMinBound :: BoundedType a -> C.Exp-codegenMinBound (IntegralBoundedType ty) | IntegralDict <- integralDict ty = codegenIntegralScalar ty minBound-codegenMinBound (NonNumBoundedType ty) | NonNumDict <- nonNumDict ty = codegenNonNumScalar ty minBound---codegenMaxBound :: BoundedType a -> C.Exp-codegenMaxBound (IntegralBoundedType ty) | IntegralDict <- integralDict ty = codegenIntegralScalar ty maxBound-codegenMaxBound (NonNumBoundedType ty) | NonNumDict <- nonNumDict ty = codegenNonNumScalar ty maxBound----- Methods from Num, Floating, Fractional and RealFrac----codegenAbs :: NumType a -> C.Exp -> C.Exp-codegenAbs (FloatingNumType ty) x = ccall (FloatingNumType ty `postfix` "fabs") [x]-codegenAbs (IntegralNumType ty) x =- case ty of- TypeWord _ -> x- TypeWord8 _ -> x- TypeWord16 _ -> x- TypeWord32 _ -> x- TypeWord64 _ -> x- TypeCUShort _ -> x- TypeCUInt _ -> x- TypeCULong _ -> x- TypeCULLong _ -> x- _ -> ccall "abs" [x]---codegenSig :: NumType a -> C.Exp -> C.Exp-codegenSig (IntegralNumType ty) = codegenIntegralSig ty-codegenSig (FloatingNumType ty) = codegenFloatingSig ty--codegenIntegralSig :: IntegralType a -> C.Exp -> C.Exp-codegenIntegralSig ty x =- case ty of- TypeWord _ -> unsigned- TypeWord8 _ -> unsigned- TypeWord16 _ -> unsigned- TypeWord32 _ -> unsigned- TypeWord64 _ -> unsigned- TypeCUShort _ -> unsigned- TypeCUInt _ -> unsigned- TypeCULong _ -> unsigned- TypeCULLong _ -> unsigned- _ -> signed- where- unsigned = [cexp| $exp:x > $exp:zero |]- signed = [cexp| ($exp:x > $exp:zero) - ($exp:x < $exp:zero) |]- zero | IntegralDict <- integralDict ty- = codegenIntegralScalar ty 0--codegenFloatingSig :: FloatingType a -> C.Exp -> C.Exp-codegenFloatingSig ty x =- [cexp|$exp:x == $exp:zero- ? $exp:zero- : $exp:(ccall (FloatingNumType ty `postfix` "copysign") [one,x]) |]- where- zero | FloatingDict <- floatingDict ty = codegenFloatingScalar ty 0- one | FloatingDict <- floatingDict ty = codegenFloatingScalar ty 1---codegenRecip :: FloatingType a -> C.Exp -> C.Exp-codegenRecip ty x | FloatingDict <- floatingDict ty = [cexp|$exp:(codegenFloatingScalar ty 1) / $exp:x|]---codegenLogBase :: FloatingType a -> C.Exp -> C.Exp -> C.Exp-codegenLogBase ty x y = let a = ccall (FloatingNumType ty `postfix` "log") [x]- b = ccall (FloatingNumType ty `postfix` "log") [y]- in- [cexp|$exp:b / $exp:a|]---codegenMin :: ScalarType a -> C.Exp -> C.Exp -> C.Exp-codegenMin (NumScalarType ty@(IntegralNumType _)) a b = ccall (ty `postfix` "min") [a,b]-codegenMin (NumScalarType ty@(FloatingNumType _)) a b = ccall (ty `postfix` "fmin") [a,b]-codegenMin (NonNumScalarType _) a b =- let ty = scalarType :: ScalarType Int32- in codegenMin ty (ccast ty a) (ccast ty b)---codegenMax :: ScalarType a -> C.Exp -> C.Exp -> C.Exp-codegenMax (NumScalarType ty@(IntegralNumType _)) a b = ccall (ty `postfix` "max") [a,b]-codegenMax (NumScalarType ty@(FloatingNumType _)) a b = ccall (ty `postfix` "fmax") [a,b]-codegenMax (NonNumScalarType _) a b =- let ty = scalarType :: ScalarType Int32- in codegenMax ty (ccast ty a) (ccast ty b)----- Type coercions----codegenOrd :: C.Exp -> C.Exp-codegenOrd = ccast (scalarType :: ScalarType Int)--codegenChr :: C.Exp -> C.Exp-codegenChr = ccast (scalarType :: ScalarType Char)--codegenBoolToInt :: C.Exp -> C.Exp-codegenBoolToInt = ccast (scalarType :: ScalarType Int)--codegenFromIntegral :: IntegralType a -> NumType b -> C.Exp -> C.Exp-codegenFromIntegral _ ty = ccast (NumScalarType ty)--codegenTruncate :: FloatingType a -> IntegralType b -> C.Exp -> C.Exp-codegenTruncate ta tb x- = ccast (NumScalarType (IntegralNumType tb))- $ ccall (FloatingNumType ta `postfix` "trunc") [x]--codegenRound :: FloatingType a -> IntegralType b -> C.Exp -> C.Exp-codegenRound ta tb x- = ccast (NumScalarType (IntegralNumType tb))- $ ccall (FloatingNumType ta `postfix` "round") [x]--codegenFloor :: FloatingType a -> IntegralType b -> C.Exp -> C.Exp-codegenFloor ta tb x- = ccast (NumScalarType (IntegralNumType tb))- $ ccall (FloatingNumType ta `postfix` "floor") [x]--codegenCeiling :: FloatingType a -> IntegralType b -> C.Exp -> C.Exp-codegenCeiling ta tb x- = ccast (NumScalarType (IntegralNumType tb))- $ ccall (FloatingNumType ta `postfix` "ceil") [x]-- -- Auxiliary Functions -- ------------------- ccast :: ScalarType a -> C.Exp -> C.Exp-ccast ty x = [cexp|($ty:(codegenScalarType ty)) $exp:x|]+ccast ty x = [cexp| ($ty:(typeOf ty)) $exp:x |] ccastTup :: TupleType e -> [C.Exp] -> [C.Exp] ccastTup ty = fst . travTup ty where- travTup :: TupleType e -> [C.Exp] -> ([C.Exp],[C.Exp])+ travTup :: TupleType e -> [C.Exp] -> ([C.Exp], [C.Exp]) travTup UnitTuple xs = ([], xs) travTup (SingleTuple ty') (x:xs) = ([ccast ty' x], xs)- travTup (PairTuple l r) xs = let- (ls, xs' ) = travTup l xs- (rs, xs'') = travTup r xs'- in (ls ++ rs, xs'')- travTup _ _ = $internalError "ccastTup" "not enough expressions to match type"---postfix :: NumType a -> String -> String-postfix (FloatingNumType (TypeFloat _)) x = x ++ "f"-postfix (FloatingNumType (TypeCFloat _)) x = x ++ "f"-postfix _ x = x+ travTup (PairTuple l r) xs =+ let (ls, xs' ) = travTup l xs+ (rs, xs'') = travTup r xs'+ in (ls ++ rs, xs'')+ travTup _ _ = $internalError "ccastTup" "not enough expressions to match type"
+ Data/Array/Accelerate/CUDA/CodeGen/Arithmetic.hs view
@@ -0,0 +1,393 @@+{-# LANGUAGE GADTs #-}+{-# LANGUAGE NoImplicitPrelude #-}+{-# LANGUAGE PatternGuards #-}+{-# LANGUAGE QuasiQuotes #-}+{-# LANGUAGE ScopedTypeVariables #-}+{-# LANGUAGE TemplateHaskell #-}+{-# LANGUAGE ViewPatterns #-}+-- |+-- Module : Data.Array.Accelerate.CUDA.CodeGen.Arithmetic+-- 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>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.CUDA.CodeGen.Arithmetic+ where++-- friends+import Data.Array.Accelerate.Type+import Data.Array.Accelerate.Error+import Data.Array.Accelerate.CUDA.CodeGen.Base+import Data.Array.Accelerate.CUDA.CodeGen.Constant+import Data.Array.Accelerate.CUDA.CodeGen.Monad+import Data.Array.Accelerate.CUDA.CodeGen.Type++-- libraries+import Prelude ( String, ($), (++), (-), undefined, otherwise )+import Data.Bits ( finiteBitSize )+import Control.Monad.State.Strict+import Language.C+import Language.C.Quote.CUDA+import Foreign.Storable ( sizeOf )+++-- Operations from Num+-- ===================++add :: Exp -> Exp -> Exp+add x y = [cexp| $exp:x + $exp:y |]++sub :: Exp -> Exp -> Exp+sub x y = [cexp| $exp:x - $exp:y |]++mul :: Exp -> Exp -> Exp+mul x y = [cexp| $exp:x * $exp:y |]++negate :: Exp -> Exp+negate x = [cexp| - $exp:x |]++abs :: forall a. NumType a -> Exp -> Exp+abs (FloatingNumType t) x+ = mathf t "fabs" [x]++abs (IntegralNumType t) x+ | signedIntegralNum t+ , IntegralDict <- integralDict t+ = case sizeOf (undefined::a) of+ 8 -> ccall "llabs" [x]+ _ -> ccall "abs" [x]++ | otherwise+ = x++signum :: NumType a -> Exp -> Gen Exp+signum (IntegralNumType t) x+ | IntegralDict <- integralDict t+ , unsignedIntegralNum t+ = return [cexp| $exp:x > $exp:(integral t 0) |]++ | IntegralDict <- integralDict t+ = do x' <- bind (typeOf t) x+ return [cexp| ($exp:x' > $exp:(integral t 0)) - ($exp:x' < $exp:(integral t 0)) |]++signum (FloatingNumType t) x+ | FloatingDict <- floatingDict t+ = do x' <- bind (typeOf t) x+ return [cexp| $exp:x' == $exp:(floating t 0)+ ? $exp:(floating t 0)+ : $exp:(mathf t "copysign" [floating t 1, x']) |]+++-- Operators from Integral & Bits+-- ==============================++quot :: Exp -> Exp -> Exp+quot x y = [cexp| $exp:x / $exp:y |]++rem :: Exp -> Exp -> Exp+rem x y = [cexp| $exp:x % $exp:y |]++quotRem :: IntegralType a -> Exp -> Exp -> Gen (Exp,Exp)+quotRem (typeOf -> t') x y = do+ x' <- bind t' x+ y' <- bind t' y+ q <- bind t' (x' `quot` y')+ r <- bind t' (x' `sub` (y' `mul` q))+ return (q, r)++idiv :: IntegralType a -> Exp -> Exp -> Gen Exp+idiv t x y+ | unsignedIntegralNum t+ = return (x `quot` y)++ | IntegralDict <- integralDict t+ , zero <- integral t 0+ , one <- integral t 1+ = do+ x' <- bind (typeOf t) x+ y' <- bind (typeOf t) y+ return $+ cases [ ((x' `gt` zero) `land` (y' `lt` zero), ((x' `sub` one) `quot` y') `sub` one)+ , ((x' `lt` zero) `land` (y' `gt` zero), ((x' `add` one) `quot` y') `sub` one)+ ]+ (x' `quot` y')++mod :: IntegralType a -> Exp -> Exp -> Gen Exp+mod t x y+ | unsignedIntegralNum t+ = return (x `rem` y)++ | IntegralDict <- integralDict t+ , zero <- integral t 0+ = do+ x' <- bind (typeOf t) x+ y' <- bind (typeOf t) y+ r <- bind (typeOf t) (x' `rem` y')+ return $+ ((((x' `gt` zero) `land` (y' `lt` zero)) `lor` ((x' `lt` zero) `land` (y' `gt` zero)))+ ?: ( r `neq` zero ?: ( r `add` y', zero )+ , r ))++divMod :: IntegralType a -> Exp -> Exp -> Gen (Exp, Exp)+divMod t x y | IntegralDict <- integralDict t = do+ x' <- bind (typeOf t) x+ y' <- bind (typeOf t) y+ (q,r) <- quotRem t x' y'++ sr <- signum (IntegralNumType t) r+ sy' <- signum (IntegralNumType t) y'++ -- Somewhat awful way to inject an ifThenElse statement+ vd <- lift fresh+ vm <- lift fresh+ modify (\st -> st { localBindings = [citem| $ty:(typeOf t) $id:vd, $id:vm; |] : localBindings st })+ modify (\st -> st { localBindings = [citem| if ( $exp:(sr `eq` negate sy') ) {+ $id:vd = $exp:(q `sub` integral t 1);+ $id:vm = $exp:(r `add` y') ;+ } else {+ $id:vd = $exp:q;+ $id:vm = $exp:r;+ } |] : localBindings st })+ return ( cvar vd, cvar vm )+++band :: Exp -> Exp -> Exp+band x y = [cexp| $exp:x & $exp:y |]++bor :: Exp -> Exp -> Exp+bor x y = [cexp| $exp:x | $exp:y |]++xor :: Exp -> Exp -> Exp+xor x y = [cexp| $exp:x ^ $exp:y |]++bnot :: Exp -> Exp+bnot x = [cexp| ~ $exp:x |]++shiftL :: Exp -> Exp -> Exp+shiftL x i = [cexp| $exp:x << $exp:i |]++-- Arithmetic right shift (unchecked)+--+shiftRA :: Exp -> Exp -> Exp+shiftRA x i = [cexp| $exp:x >> $exp:i |]++-- Logical right shift (unchecked)+--+shiftRL :: IntegralType a -> Exp -> Exp -> Exp+shiftRL ty x i =+ let int = typeOf (integralType :: IntegralType Int)+ word = typeOf (integralType :: IntegralType Word)+ in+ case ty of+ TypeInt{} -> [cexp| ($ty:int) (($ty:word) $exp:x >> $exp:i) |]+ TypeInt8{} -> [cexp| (typename Int8) ((typename Word8) $exp:x >> $exp:i) |]+ TypeInt16{} -> [cexp| (typename Int16) ((typename Word16) $exp:x >> $exp:i) |]+ TypeInt32{} -> [cexp| (typename Int32) ((typename Word32) $exp:x >> $exp:i) |]+ TypeInt64{} -> [cexp| (typename Int64) ((typename Word64) $exp:x >> $exp:i) |]+ TypeCShort{} -> [cexp| (short) ((unsigned short) $exp:x >> $exp:i) |]+ TypeCInt{} -> [cexp| (int) ((unsigned int) $exp:x >> $exp:i) |]+ TypeCLong{} -> [cexp| (long) ((unsigned long) $exp:x >> $exp:i) |]+ TypeCLLong{} -> [cexp| (long long) ((unsigned long long) $exp:x >> $exp:i) |]++ -- unsigned types use arithmetic shift+ _ -> $internalCheck "shiftRL" "unhandled signed type" (unsignedIntegralNum ty) (shiftRA x i)+++rotateL :: forall a. IntegralType a -> Exp -> Exp -> Gen Exp+rotateL t x i | IntegralDict <- integralDict t = do+ let int = integralType :: IntegralType Int+ wsib = finiteBitSize (undefined::a)+ --+ x' <- bind (typeOf t) x+ i' <- bind (typeOf int) (i `band` integral int (wsib - 1))+ return $ (x' `shiftL` i') `bor` (shiftRL t x' (integral int wsib `sub` i'))++rotateR :: IntegralType a -> Exp -> Exp -> Gen Exp+rotateR t x i = rotateL t x (negate i)+++-- Operators from Fractional & Floating+-- ====================================++fdiv :: Exp -> Exp -> Exp+fdiv x y = [cexp| $exp:x / $exp:y |]++recip :: FloatingType a -> Exp -> Exp+recip t x | FloatingDict <- floatingDict t = fdiv (floating t 1) x++sin :: FloatingType a -> Exp -> Exp+sin t x = mathf t "sin" [x]++cos :: FloatingType a -> Exp -> Exp+cos t x = mathf t "cos" [x]++tan :: FloatingType a -> Exp -> Exp+tan t x = mathf t "tan" [x]++asin :: FloatingType a -> Exp -> Exp+asin t x = mathf t "asin" [x]++acos :: FloatingType a -> Exp -> Exp+acos t x = mathf t "acos" [x]++atan :: FloatingType a -> Exp -> Exp+atan t x = mathf t "atan" [x]++sinh :: FloatingType a -> Exp -> Exp+sinh t x = mathf t "sinh" [x]++cosh :: FloatingType a -> Exp -> Exp+cosh t x = mathf t "cosh" [x]++tanh :: FloatingType a -> Exp -> Exp+tanh t x = mathf t "tanh" [x]++asinh :: FloatingType a -> Exp -> Exp+asinh t x = mathf t "asinh" [x]++acosh :: FloatingType a -> Exp -> Exp+acosh t x = mathf t "acosh" [x]++atanh :: FloatingType a -> Exp -> Exp+atanh t x = mathf t "atanh" [x]++exp :: FloatingType a -> Exp -> Exp+exp t x = mathf t "exp" [x]++sqrt :: FloatingType a -> Exp -> Exp+sqrt t x = mathf t "sqrt" [x]++pow :: FloatingType a -> Exp -> Exp -> Exp+pow t x y = mathf t "pow" [x,y]++log :: FloatingType a -> Exp -> Exp+log t x = mathf t "log" [x]++logBase :: FloatingType a -> Exp -> Exp -> Exp+logBase t x y = log t y `fdiv` log t x+++-- Operators from RealFrac+-- =======================++trunc :: FloatingType a -> IntegralType b -> Exp -> Exp+trunc ta tb x = cast tb $ mathf ta "trunc" [x]++round :: FloatingType a -> IntegralType b -> Exp -> Exp+round ta tb x = cast tb $ mathf ta "round" [x]++floor :: FloatingType a -> IntegralType b -> Exp -> Exp+floor ta tb x = cast tb $ mathf ta "floor" [x]++ceiling :: FloatingType a -> IntegralType b -> Exp -> Exp+ceiling ta tb x = cast tb $ mathf ta "ceil" [x]+++-- Operators from RealFloat+-- ========================++atan2 :: FloatingType a -> Exp -> Exp -> Exp+atan2 t x y = mathf t "atan2" [x, y]++isNaN :: Exp -> Exp+isNaN x = ccall "isnan" [x]+++-- Relational and equality operators+-- =================================++lt :: Exp -> Exp -> Exp+lt x y = [cexp| $exp:x < $exp:y |]++gt :: Exp -> Exp -> Exp+gt x y = [cexp| $exp:x > $exp:y |]++leq :: Exp -> Exp -> Exp+leq x y = [cexp| $exp:x <= $exp:y |]++geq :: Exp -> Exp -> Exp+geq x y = [cexp| $exp:x >= $exp:y |]++eq :: Exp -> Exp -> Exp+eq x y = [cexp| $exp:x == $exp:y |]++neq :: Exp -> Exp -> Exp+neq x y = [cexp| $exp:x != $exp:y |]++max :: ScalarType a -> Exp -> Exp -> Exp+max (NonNumScalarType _) x y =+ let t = scalarType :: ScalarType Int32+ in max t (cast t x) (cast t y)++max (NumScalarType (IntegralNumType _)) x y = ccall "max" [x,y]+max (NumScalarType (FloatingNumType t)) x y = mathf t "fmax" [x,y]++min :: ScalarType a -> Exp -> Exp -> Exp+min (NonNumScalarType _) x y =+ let t = scalarType :: ScalarType Int32+ in min t (cast t x) (cast t y)++min (NumScalarType (IntegralNumType _)) x y = ccall "min" [x,y]+min (NumScalarType (FloatingNumType t)) x y = mathf t "fmin" [x,y]+++-- Logical operators+-- =================++land :: Exp -> Exp -> Exp+land x y = [cexp| $exp:x && $exp:y |]++lor :: Exp -> Exp -> Exp+lor x y = [cexp| $exp:x || $exp:y |]++lnot :: Exp -> Exp+lnot x = [cexp| ! $exp:x |]+++-- Type Conversions+-- ================++ord :: Exp -> Exp+ord = cast (scalarType :: ScalarType Int)++chr :: Exp -> Exp+chr = cast (scalarType :: ScalarType Char)++boolToInt :: Exp -> Exp+boolToInt = cast (scalarType :: ScalarType Int)++fromIntegral :: IntegralType a -> NumType b -> Exp -> Exp+fromIntegral _ tb = cast tb+++-- Helpers+-- =======++cast :: TypeOf a => a -> Exp -> Exp+cast t x = [cexp| ($ty:(typeOf t)) $exp:x |]++mathf :: forall t. FloatingType t -> String -> [Exp] -> Exp+mathf ty f args | FloatingDict <- floatingDict ty =+ let+ fun = f ++ case sizeOf (undefined :: t) of+ 4 -> "f"+ 8 -> []+ 16 -> "l" -- long double+ _ -> $internalError "mathf" "unsupported floating point size"+ in+ ccall fun args+++infix 0 ?:+(?:) :: Exp -> (Exp, Exp) -> Exp+(?:) p (t,e) = [cexp| $exp:p ? $exp:t : $exp:e |]++cases :: [(Exp, Exp)] -> Exp -> Exp+cases [] def = def+cases ((p,b):rest) def = p ?: (b, cases rest def)+
Data/Array/Accelerate/CUDA/CodeGen/Base.hs view
@@ -1,12 +1,17 @@+{-# LANGUAGE CPP #-} {-# LANGUAGE FlexibleInstances #-} {-# LANGUAGE GADTs #-} {-# LANGUAGE ImpredicativeTypes #-} {-# LANGUAGE MultiParamTypeClasses #-}-{-# LANGUAGE OverlappingInstances #-} {-# LANGUAGE PatternGuards #-} {-# LANGUAGE QuasiQuotes #-} {-# LANGUAGE ScopedTypeVariables #-} {-# LANGUAGE TemplateHaskell #-}+#if __GLASGOW_HASKELL__ <= 708+{-# LANGUAGE OverlappingInstances #-}+{-# OPTIONS_GHC -fno-warn-unrecognised-pragmas #-}+#endif+ -- | -- Module : Data.Array.Accelerate.CUDA.CodeGen.Base -- Copyright : [2008..2014] Manuel M T Chakravarty, Gabriele Keller@@ -26,7 +31,7 @@ Name, namesOfArray, namesOfAvar, groupOfInt, -- Declaration generation- cint, cvar, ccall, cchar, cintegral, cbool, cshape, csize, cindexHead, cindexTail, ctoIndex, cfromIndex,+ cint, cvar, ccall, cchar, cintegral, cbool, cshape, cslice, csize, cindexHead, cindexTail, ctoIndex, cfromIndex, readArray, writeArray, shared, indexArray, environment, arrayAsTex, arrayAsArg, umul24, gridSize, threadIdx,@@ -50,6 +55,7 @@ -- friends import Data.Array.Accelerate.Type import Data.Array.Accelerate.Error+import Data.Array.Accelerate.Array.Representation ( SliceIndex(..) ) import Data.Array.Accelerate.Array.Sugar ( Array, Shape, Elt ) import Data.Array.Accelerate.Analysis.Shape import Data.Array.Accelerate.CUDA.CodeGen.Type@@ -136,7 +142,7 @@ -- ----------------------- cint :: C.Type-cint = codegenScalarType (scalarType :: ScalarType Int)+cint = typeOf (scalarType :: ScalarType Int) cvar :: Name -> C.Exp cvar x = [cexp|$id:x|]@@ -153,6 +159,17 @@ cbool :: Bool -> C.Exp cbool = cintegral . fromEnum +cslice :: SliceIndex slix sl co dim -> Name -> ([C.Param], [C.Exp], [(C.Type, Name)])+cslice slix sl =+ let xs = cshape' (ncodims slix) sl+ args = [ [cparam| const $ty:cint $id:x |] | x <- xs ]+ in (args, map cvar xs, zip (repeat cint) xs)+ where+ ncodims :: SliceIndex slix sl co dim -> Int+ ncodims SliceNil = 0+ ncodims (SliceAll s) = ncodims s+ ncodims (SliceFixed s) = ncodims s + 1+ -- Generate all the names of a shape given a base name and dimensionality cshape :: Int -> Name -> [C.Exp] cshape dim sh = [ cvar x | x <- cshape' dim sh ]@@ -246,7 +263,9 @@ :: forall aenv sh e. (Shape sh, Elt e) => Name -- group names -> Array sh e -- dummy to fix types- -> ( [C.Param], CUDelayedAcc aenv sh e )+ -> ( [C.Param]+ , [C.Exp]+ , CUDelayedAcc aenv sh e ) readArray grp dummy = let (sh, arrs) = namesOfArray grp (undefined :: e) args = arrayAsArg dummy grp@@ -257,7 +276,7 @@ manifest = CUDelayed (CUExp ([], sh')) ($internalError "readArray" "linear indexing only") (CUFun1 (zip (repeat True)) (\[i] -> get (rvalue i)))- in ( args, manifest )+ in ( args, sh', manifest ) -- Generate function parameters and corresponding variable names for the@@ -383,7 +402,6 @@ in unzip3 $ zipWith local elt [n-1, n-2 .. 0] - class Lvalue a where lvalue :: a -> C.Exp -> C.BlockItem @@ -438,17 +456,17 @@ instance (Lvalue l, Rvalue r) => Assign l r where assign lhs rhs = [ lvalue lhs (rvalue rhs) ] -instance Assign l r => Assign (Bool,l) r where+instance {-# OVERLAPS #-} Assign l r => Assign (Bool,l) r where assign (used,lhs) rhs | used = assign lhs rhs | otherwise = [] -instance Assign l r => Assign [l] [r] where+instance {-# OVERLAPS #-} Assign l r => Assign [l] [r] where assign [] [] = [] assign (x:xs) (y:ys) = assign x y ++ assign xs ys assign _ _ = $internalError ".=." "argument mismatch" -instance Assign l r => Assign l ([C.BlockItem], r) where+instance {-# OVERLAPS #-} Assign l r => Assign l ([C.BlockItem], r) where assign lhs (env, rhs) = env ++ assign lhs rhs
+ Data/Array/Accelerate/CUDA/CodeGen/Constant.hs view
@@ -0,0 +1,93 @@+{-# LANGUAGE GADTs #-}+{-# LANGUAGE PatternGuards #-}+{-# LANGUAGE QuasiQuotes #-}+{-# LANGUAGE TemplateHaskell #-}+-- |+-- Module : Data.Array.Accelerate.CUDA.CodeGen.Constant+-- 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>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.CUDA.CodeGen.Constant+ where++-- friends+import Data.Array.Accelerate.AST ( PrimConst(..) )+import Data.Array.Accelerate.Type+import Data.Array.Accelerate.CUDA.CodeGen.Base+import Data.Array.Accelerate.CUDA.CodeGen.Type++-- libraries+import Data.Loc+import Data.Char+import Language.C+import Language.C.Quote.CUDA+++-- | A constant value. Note that this follows the EltRepr representation of the+-- type, meaning that any nested tupling on the surface type is flattened.+--+constant :: TupleType a -> a -> [Exp]+constant UnitTuple _ = []+constant (SingleTuple ty) c = [scalar ty c]+constant (PairTuple ty1 ty0) (cs,c) = constant ty1 cs ++ constant ty0 c++-- | A constant scalar value+--+scalar :: ScalarType a -> a -> Exp+scalar (NumScalarType ty) = num ty+scalar (NonNumScalarType ty) = nonnum ty++-- | A constant numeric value+--+num :: NumType a -> a -> Exp+num (IntegralNumType ty) = integral ty+num (FloatingNumType ty) = floating ty++-- | A constant integral value+--+integral :: IntegralType a -> a -> Exp+integral ty x | IntegralDict <- integralDict ty = [cexp| ( $ty:(typeOf ty) ) $exp:(cintegral x) |]++-- | A constant floating-point value+--+floating :: FloatingType a -> a -> Exp+floating (TypeFloat _) x = Const (FloatConst (shows x "f") (toRational x) noLoc) noLoc+floating (TypeCFloat _) x = Const (FloatConst (shows x "f") (toRational x) noLoc) noLoc+floating (TypeDouble _) x = Const (DoubleConst (show x) (toRational x) noLoc) noLoc+floating (TypeCDouble _) x = Const (DoubleConst (show x) (toRational x) noLoc) noLoc+++-- | A constant non-numeric value+--+nonnum :: NonNumType a -> a -> Exp+nonnum (TypeBool _) x = cbool x+nonnum (TypeChar _) x = [cexp|$char:x|]+nonnum (TypeCChar _) x = [cexp|$char:(chr (fromIntegral x))|]+nonnum (TypeCUChar _) x = [cexp|$char:(chr (fromIntegral x))|]+nonnum (TypeCSChar _) x = [cexp|$char:(chr (fromIntegral x))|]+++-- | Primitive constants+--+primConst :: PrimConst t -> Exp+primConst (PrimMinBound t) = primMinBound t+primConst (PrimMaxBound t) = primMaxBound t+primConst (PrimPi t) = primPi t++primMinBound :: BoundedType a -> Exp+primMinBound (IntegralBoundedType ty) | IntegralDict <- integralDict ty = integral ty minBound+primMinBound (NonNumBoundedType ty) | NonNumDict <- nonNumDict ty = nonnum ty minBound++primMaxBound :: BoundedType a -> Exp+primMaxBound (IntegralBoundedType ty) | IntegralDict <- integralDict ty = integral ty maxBound+primMaxBound (NonNumBoundedType ty) | NonNumDict <- nonNumDict ty = nonnum ty maxBound++primPi :: FloatingType a -> Exp+primPi ty | FloatingDict <- floatingDict ty = floating ty pi+
Data/Array/Accelerate/CUDA/CodeGen/IndexSpace.hs view
@@ -204,7 +204,8 @@ $items:(atomically jx [ dce_y y .=. setOut jx- , setOut jx .=. combine x y ]+ , setOut jx .=. combine x y+ ] ) } }@@ -288,14 +289,15 @@ | otherwise = [ [citem| typename Int32 done = 0; |] , [citem| do {- __threadfence();+ typename Int32 *addr = &lock[ $exp:(cvar i) ]; - if ( atomicExch(&lock[ $exp:(cvar i) ], 1) == 0 ) {+ if ( atomicExch( addr, 1 ) == 0 ) { $items:body done = 1;- atomicExch(&lock[ $exp:(cvar i) ], 0);+ atomicExch( addr, 0 ); }+ __threadfence(); } while (done == 0); |] ]
Data/Array/Accelerate/CUDA/CodeGen/PrefixSum.hs view
@@ -4,6 +4,7 @@ {-# LANGUAGE PatternGuards #-} {-# LANGUAGE QuasiQuotes #-} {-# LANGUAGE ScopedTypeVariables #-}+{-# LANGUAGE TypeOperators #-} -- | -- Module : Data.Array.Accelerate.CUDA.CodeGen.PrefixSum -- Copyright : [2008..2014] Manuel M T Chakravarty, Gabriele Keller@@ -28,60 +29,96 @@ import Language.C.Quote.CUDA import qualified Language.C.Syntax as C -import Data.Array.Accelerate.Array.Sugar ( Vector, Scalar, Elt, DIM1 )+import Data.Array.Accelerate.Array.Sugar+import Data.Array.Accelerate.Analysis.Match+ import Data.Array.Accelerate.CUDA.AST+import Data.Array.Accelerate.CUDA.Analysis.Shape import Data.Array.Accelerate.CUDA.CodeGen.Base +errorMsg :: String+errorMsg+ = error+ $ unlines [ "accelerate-cuda does not support rank-polymorphic scans. Please switch to accelerate-llvm-ptx instead."+ , ""+ , "*** https://hackage.haskell.org/package/accelerate-llvm-ptx ***"+ , "*** https://github.com/AccelerateHS/accelerate-llvm ***"+ ] -- Wrappers -- -------- mkScanl, mkScanr- :: Elt e+ :: forall aenv sh e. (Shape sh, Elt e) => DeviceProperties -> Gamma aenv -> CUFun2 aenv (e -> e -> e) -> CUExp aenv e- -> CUDelayedAcc aenv DIM1 e- -> [CUTranslSkel aenv (Vector e)]-mkScanl dev aenv f z a =- [ mkScan L dev aenv f (Just z) a- , mkScanUp1 L dev aenv f a- , mkScanUp2 L dev aenv f (Just z) ]+ -> CUDelayedAcc aenv (sh:.Int) e+ -> [CUTranslSkel aenv (Array (sh:.Int) e)]+mkScanl dev aenv f z a+ | Just Refl <- matchShapeType (undefined::sh) (undefined::Z)+ = [ mkScan L dev aenv f (Just z) a+ , mkScanUp1 L dev aenv f a+ , mkScanUp2 L dev aenv f (Just z) ] -mkScanr dev aenv f z a =- [ mkScan R dev aenv f (Just z) a- , mkScanUp1 R dev aenv f a- , mkScanUp2 R dev aenv f (Just z) ]+ | otherwise+ = error errorMsg +mkScanr dev aenv f z a+ | Just Refl <- matchShapeType (undefined::sh) (undefined::Z)+ = [ mkScan R dev aenv f (Just z) a+ , mkScanUp1 R dev aenv f a+ , mkScanUp2 R dev aenv f (Just z) ]++ | otherwise+ = error errorMsg+ mkScanl1, mkScanr1- :: Elt e+ :: forall aenv sh e. (Shape sh, Elt e) => DeviceProperties -> Gamma aenv -> CUFun2 aenv (e -> e -> e)- -> CUDelayedAcc aenv DIM1 e- -> [CUTranslSkel aenv (Vector e)]-mkScanl1 dev aenv f a =- [ mkScan L dev aenv f Nothing a- , mkScanUp1 L dev aenv f a- , mkScanUp2 L dev aenv f Nothing ]+ -> CUDelayedAcc aenv (sh:.Int) e+ -> [CUTranslSkel aenv (Array (sh:.Int) e)]+mkScanl1 dev aenv f a+ | Just Refl <- matchShapeType (undefined::sh) (undefined::Z)+ = [ mkScan L dev aenv f Nothing a+ , mkScanUp1 L dev aenv f a+ , mkScanUp2 L dev aenv f Nothing ] -mkScanr1 dev aenv f a =- [ mkScan R dev aenv f Nothing a- , mkScanUp1 R dev aenv f a- , mkScanUp2 R dev aenv f Nothing ]+ | otherwise+ = error errorMsg +mkScanr1 dev aenv f a+ | Just Refl <- matchShapeType (undefined::sh) (undefined::Z)+ = [ mkScan R dev aenv f Nothing a+ , mkScanUp1 R dev aenv f a+ , mkScanUp2 R dev aenv f Nothing ]++ | otherwise+ = error errorMsg+ mkScanl', mkScanr'- :: Elt e+ :: forall aenv sh e. (Shape sh, Elt e) => DeviceProperties -> Gamma aenv -> CUFun2 aenv (e -> e -> e) -> CUExp aenv e- -> CUDelayedAcc aenv DIM1 e- -> [CUTranslSkel aenv (Vector e, Scalar e)]-mkScanl' dev aenv f z = map cast . mkScanl dev aenv f z-mkScanr' dev aenv f z = map cast . mkScanr dev aenv f z+ -> CUDelayedAcc aenv (sh:.Int) e+ -> [CUTranslSkel aenv (Array (sh:.Int) e, Array sh e)]+mkScanl' dev aenv f z+ | Just Refl <- matchShapeType (undefined::sh) (undefined::Z)+ = map cast . mkScanl dev aenv f z+ | otherwise+ = error errorMsg +mkScanr' dev aenv f z+ | Just Refl <- matchShapeType (undefined::sh) (undefined::Z)+ = map cast . mkScanr dev aenv f z+ | otherwise+ = error errorMsg+ cast :: CUTranslSkel aenv a -> CUTranslSkel aenv b cast (CUTranslSkel entry code) = CUTranslSkel entry code @@ -439,7 +476,7 @@ -> Maybe (CUExp aenv e) -> CUTranslSkel aenv (Vector e) mkScanUp2 dir dev aenv f z- = let (_, get) = readArray "Blk" (undefined :: Vector e)+ = let (_, _, get) = readArray "Blk" (undefined :: Vector e) in mkScan dir dev aenv f z get
Data/Array/Accelerate/CUDA/CodeGen/Reduction.hs view
@@ -93,7 +93,7 @@ -> CUDelayedAcc aenv (sh :. Int) e -> [ CUTranslSkel aenv (Array sh e) ] mkFoldAll dev aenv f z a- = let (_, rec) = readArray "Rec" (undefined :: Array (sh:.Int) e)+ = let (_, _, rec) = readArray "Rec" (undefined :: Array (sh:.Int) e) in [ mkFoldAll' False dev aenv f z a , mkFoldAll' True dev aenv f z rec ]@@ -188,9 +188,9 @@ foldAll = maybe "fold1All" (const "foldAll") mseed (texIn, argIn) = environment dev aenv (argOut, _, setOut) = writeArray "Out" (undefined :: Array (sh :. Int) e)- (argRec, _)+ (argRec, _, _) | recursive = readArray "Rec" (undefined :: Array (sh :. Int) e)- | otherwise = ([], undefined)+ | otherwise = ([], undefined, undefined) (_, x, declx) = locals "x" (undefined :: e) (_, y, decly) = locals "y" (undefined :: e)
+ Data/Array/Accelerate/CUDA/CodeGen/Streaming.hs view
@@ -0,0 +1,84 @@+{-# LANGUAGE GADTs #-}+{-# LANGUAGE ImpredicativeTypes #-}+{-# LANGUAGE PatternGuards #-}+{-# LANGUAGE QuasiQuotes #-}+{-# LANGUAGE ScopedTypeVariables #-}+{-# LANGUAGE TemplateHaskell #-}++-- |+-- Module : Data.Array.Accelerate.CUDA.CodeGen.Streaming+--+-- Maintainer : Frederik Meisner Madsen <fmma@diku.dk>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.CUDA.CodeGen.Streaming (++ mkToSeq++) where++import Language.C.Quote.CUDA+import Foreign.CUDA.Analysis.Device++import Data.Array.Accelerate.Error ( internalError )+import Data.Array.Accelerate.Array.Representation ( SliceIndex(..) )+import Data.Array.Accelerate.Array.Sugar ( Array, Shape, Elt, EltRepr )+import Data.Array.Accelerate.CUDA.AST+import Data.Array.Accelerate.CUDA.CodeGen.Base++mkToSeq+ :: forall slix aenv sh co sl a. (Shape sl, Shape sh, Elt a)+ => SliceIndex slix+ (EltRepr sl)+ co+ (EltRepr sh)+ -> DeviceProperties+ -> Gamma aenv+ -> CUDelayedAcc aenv sh a+ -> CUTranslSkel aenv (Array sl a)+mkToSeq slix dev aenv arr+ | CUDelayed (CUExp shIn) (CUFun1 _ get) _ <- arr+ = CUTranslSkel "tostream" [cunit|++ $esc:("#include <accelerate_cuda.h>")+ $edecls:texIn++ extern "C" __global__ void+ tostream+ (+ $params:argIn,+ $params:slIn,+ $params:argOut+ )+ {+ $items:(sh .=. shIn)++ const int shapeSize = $exp:(csize shOut);+ const int gridSize = $exp:(gridSize dev);+ int ix;++ for ( ix = $exp:(threadIdx dev)+ ; ix < shapeSize+ ; ix += gridSize )+ {+ $items:(src .=. cfromIndex shOut "ix" "tmp")+ $items:(setOut "ix" .=. get fullsrc)+ }+ }+ |]+ where+ (slIn, _, cosrc) = cslice slix "cosrc"+ (sh, _, _) = locals "shIn" (undefined :: sh)+ (src, _, _) = locals "src" (undefined :: sl)+ fullsrc = reverse (combine slix (reverse src) (reverse cosrc))+ (texIn, argIn) = environment dev aenv+ (argOut, shOut, setOut) = writeArray "Out" (undefined :: Array sl a)++combine :: SliceIndex slix sl co dim -> [a] -> [a] -> [a]+combine SliceNil [] [] = []+combine (SliceAll sl) (x:xs) ys = x:(combine sl xs ys)+combine (SliceFixed sl) xs (y:ys) = y:(combine sl xs ys)+combine _ _ _ = $internalError "mkToSeq" "Something went wrong with the slice index."+
Data/Array/Accelerate/CUDA/CodeGen/Type.hs view
@@ -21,8 +21,9 @@ accType, accTypeTex, segmentsType, expType, eltType, eltTypeTex, eltSizeOf, - -- primitive bits...- codegenIntegralType, codegenScalarType+ -- working with reified dictionaries+ TypeOf(..),+ signedIntegralNum, unsignedIntegralNum, ) where @@ -36,23 +37,44 @@ -- libraries import Data.Bits-import Data.Typeable import Language.C.Quote.CUDA+import qualified Data.Typeable as T import qualified Language.C as C -typename :: String -> C.Type-typename name = [cty| typename $id:name |]+class TypeOf a where+ typeOf :: a -> C.Type+ texTypeOf :: a -> C.Type +instance TypeOf (ScalarType a) where+ typeOf = cScalarType+ texTypeOf = cScalarTypeTex++instance TypeOf (NumType a) where+ typeOf = cNumType+ texTypeOf = cNumTypeTex++instance TypeOf (IntegralType a) where+ typeOf = cIntegralType+ texTypeOf = cIntegralTypeTex++instance TypeOf (FloatingType a) where+ typeOf = cFloatingType+ texTypeOf = cFloatingTypeTex++instance TypeOf (NonNumType a) where+ typeOf = cNonNumType+ texTypeOf = cNonNumTypeTex++ -- Surface element types -- --------------------- - accType :: DelayedOpenAcc aenv (Sugar.Array dim e) -> [C.Type]-accType = codegenTupleType . Sugar.delayedAccType+accType = cTupleType . Sugar.delayedAccType expType :: DelayedOpenExp aenv env t -> [C.Type]-expType = codegenTupleType . Sugar.preExpType Sugar.delayedAccType+expType = cTupleType . Sugar.preExpType Sugar.delayedAccType segmentsType :: DelayedOpenAcc aenv (Sugar.Segments i) -> C.Type segmentsType seg@@ -61,10 +83,10 @@ eltType :: Sugar.Elt a => a {- dummy -} -> [C.Type]-eltType = codegenTupleType . Sugar.eltType+eltType = cTupleType . Sugar.eltType eltTypeTex :: Sugar.Elt a => a {- dummy -} -> [C.Type]-eltTypeTex = codegenTupleTex . Sugar.eltType+eltTypeTex = cTupleTypeTex . Sugar.eltType eltSizeOf :: Sugar.Elt a => a {- dummy -} -> [Int] eltSizeOf = sizeOf' . Sugar.eltType@@ -75,108 +97,131 @@ sizeOf' (PairTuple a b) = sizeOf' a ++ sizeOf' b --- Implementation----codegenTupleType :: TupleType a -> [C.Type]-codegenTupleType UnitTuple = []-codegenTupleType (SingleTuple ty) = [codegenScalarType ty]-codegenTupleType (PairTuple t1 t0) = codegenTupleType t1 ++ codegenTupleType t0 -codegenScalarType :: ScalarType a -> C.Type-codegenScalarType (NumScalarType ty) = codegenNumType ty-codegenScalarType (NonNumScalarType ty) = codegenNonNumType ty+cTupleType :: TupleType a -> [C.Type]+cTupleType UnitTuple = []+cTupleType (SingleTuple ty) = [cScalarType ty]+cTupleType (PairTuple t1 t0) = cTupleType t1 ++ cTupleType t0 -codegenNumType :: NumType a -> C.Type-codegenNumType (IntegralNumType ty) = codegenIntegralType ty-codegenNumType (FloatingNumType ty) = codegenFloatingType ty+cScalarType :: ScalarType a -> C.Type+cScalarType (NumScalarType ty) = cNumType ty+cScalarType (NonNumScalarType ty) = cNonNumType ty -codegenIntegralType :: IntegralType a -> C.Type-codegenIntegralType (TypeInt8 _) = typename "Int8"-codegenIntegralType (TypeInt16 _) = typename "Int16"-codegenIntegralType (TypeInt32 _) = typename "Int32"-codegenIntegralType (TypeInt64 _) = typename "Int64"-codegenIntegralType (TypeWord8 _) = typename "Word8"-codegenIntegralType (TypeWord16 _) = typename "Word16"-codegenIntegralType (TypeWord32 _) = typename "Word32"-codegenIntegralType (TypeWord64 _) = typename "Word64"-codegenIntegralType (TypeCShort _) = [cty|short|]-codegenIntegralType (TypeCUShort _) = [cty|unsigned short|]-codegenIntegralType (TypeCInt _) = [cty|int|]-codegenIntegralType (TypeCUInt _) = [cty|unsigned int|]-codegenIntegralType (TypeCLong _) = [cty|long int|]-codegenIntegralType (TypeCULong _) = [cty|unsigned long int|]-codegenIntegralType (TypeCLLong _) = [cty|long long int|]-codegenIntegralType (TypeCULLong _) = [cty|unsigned long long int|]-codegenIntegralType (TypeInt _) = typename (showsTypeRep (typeOf (undefined::HTYPE_INT)) "")-codegenIntegralType (TypeWord _) = typename (showsTypeRep (typeOf (undefined::HTYPE_WORD)) "")+cNumType :: NumType a -> C.Type+cNumType (IntegralNumType ty) = cIntegralType ty+cNumType (FloatingNumType ty) = cFloatingType ty -codegenFloatingType :: FloatingType a -> C.Type-codegenFloatingType (TypeFloat _) = [cty|float|]-codegenFloatingType (TypeCFloat _) = [cty|float|]-codegenFloatingType (TypeDouble _) = [cty|double|]-codegenFloatingType (TypeCDouble _) = [cty|double|]+cIntegralType :: IntegralType a -> C.Type+cIntegralType (TypeInt8 _) = typename "Int8"+cIntegralType (TypeInt16 _) = typename "Int16"+cIntegralType (TypeInt32 _) = typename "Int32"+cIntegralType (TypeInt64 _) = typename "Int64"+cIntegralType (TypeWord8 _) = typename "Word8"+cIntegralType (TypeWord16 _) = typename "Word16"+cIntegralType (TypeWord32 _) = typename "Word32"+cIntegralType (TypeWord64 _) = typename "Word64"+cIntegralType (TypeCShort _) = [cty|short|]+cIntegralType (TypeCUShort _) = [cty|unsigned short|]+cIntegralType (TypeCInt _) = [cty|int|]+cIntegralType (TypeCUInt _) = [cty|unsigned int|]+cIntegralType (TypeCLong _) = [cty|long int|]+cIntegralType (TypeCULong _) = [cty|unsigned long int|]+cIntegralType (TypeCLLong _) = [cty|long long int|]+cIntegralType (TypeCULLong _) = [cty|unsigned long long int|]+cIntegralType (TypeInt _) = typename (T.showsTypeRep (T.typeOf (undefined::HTYPE_INT)) "")+cIntegralType (TypeWord _) = typename (T.showsTypeRep (T.typeOf (undefined::HTYPE_WORD)) "") -codegenNonNumType :: NonNumType a -> C.Type-codegenNonNumType (TypeBool _) = typename "Word8"-codegenNonNumType (TypeChar _) = typename "Word32"-codegenNonNumType (TypeCChar _) = [cty|char|]-codegenNonNumType (TypeCSChar _) = [cty|signed char|]-codegenNonNumType (TypeCUChar _) = [cty|unsigned char|]+cFloatingType :: FloatingType a -> C.Type+cFloatingType (TypeFloat _) = [cty|float|]+cFloatingType (TypeCFloat _) = [cty|float|]+cFloatingType (TypeDouble _) = [cty|double|]+cFloatingType (TypeCDouble _) = [cty|double|] +cNonNumType :: NonNumType a -> C.Type+cNonNumType (TypeBool _) = typename "Word8"+cNonNumType (TypeChar _) = typename "Word32"+cNonNumType (TypeCChar _) = [cty|char|]+cNonNumType (TypeCSChar _) = [cty|signed char|]+cNonNumType (TypeCUChar _) = [cty|unsigned char|] + -- Texture types -- ------------- accTypeTex :: DelayedOpenAcc aenv (Sugar.Array dim e) -> [C.Type]-accTypeTex = codegenTupleTex . Sugar.delayedAccType+accTypeTex = cTupleTypeTex . Sugar.delayedAccType -- Implementation ---codegenTupleTex :: TupleType a -> [C.Type]-codegenTupleTex UnitTuple = []-codegenTupleTex (SingleTuple t) = [codegenScalarTex t]-codegenTupleTex (PairTuple t1 t0) = codegenTupleTex t1 ++ codegenTupleTex t0+cTupleTypeTex :: TupleType a -> [C.Type]+cTupleTypeTex UnitTuple = []+cTupleTypeTex (SingleTuple t) = [cScalarTypeTex t]+cTupleTypeTex (PairTuple t1 t0) = cTupleTypeTex t1 ++ cTupleTypeTex t0 -codegenScalarTex :: ScalarType a -> C.Type-codegenScalarTex (NumScalarType ty) = codegenNumTex ty-codegenScalarTex (NonNumScalarType ty) = codegenNonNumTex ty;+cScalarTypeTex :: ScalarType a -> C.Type+cScalarTypeTex (NumScalarType ty) = cNumTypeTex ty+cScalarTypeTex (NonNumScalarType ty) = cNonNumTypeTex ty; -codegenNumTex :: NumType a -> C.Type-codegenNumTex (IntegralNumType ty) = codegenIntegralTex ty-codegenNumTex (FloatingNumType ty) = codegenFloatingTex ty+cNumTypeTex :: NumType a -> C.Type+cNumTypeTex (IntegralNumType ty) = cIntegralTypeTex ty+cNumTypeTex (FloatingNumType ty) = cFloatingTypeTex ty -codegenIntegralTex :: IntegralType a -> C.Type-codegenIntegralTex (TypeInt8 _) = typename "TexInt8"-codegenIntegralTex (TypeInt16 _) = typename "TexInt16"-codegenIntegralTex (TypeInt32 _) = typename "TexInt32"-codegenIntegralTex (TypeInt64 _) = typename "TexInt64"-codegenIntegralTex (TypeWord8 _) = typename "TexWord8"-codegenIntegralTex (TypeWord16 _) = typename "TexWord16"-codegenIntegralTex (TypeWord32 _) = typename "TexWord32"-codegenIntegralTex (TypeWord64 _) = typename "TexWord64"-codegenIntegralTex (TypeCShort _) = typename "TexCShort"-codegenIntegralTex (TypeCUShort _) = typename "TexCUShort"-codegenIntegralTex (TypeCInt _) = typename "TexCInt"-codegenIntegralTex (TypeCUInt _) = typename "TexCUInt"-codegenIntegralTex (TypeCLong _) = typename "TexCLong"-codegenIntegralTex (TypeCULong _) = typename "TexCULong"-codegenIntegralTex (TypeCLLong _) = typename "TexCLLong"-codegenIntegralTex (TypeCULLong _) = typename "TexCULLong"-codegenIntegralTex (TypeInt _) = typename ("TexInt" ++ show (finiteBitSize (undefined::Int)))-codegenIntegralTex (TypeWord _) = typename ("TexWord" ++ show (finiteBitSize (undefined::Word)))+cIntegralTypeTex :: IntegralType a -> C.Type+cIntegralTypeTex (TypeInt8 _) = typename "TexInt8"+cIntegralTypeTex (TypeInt16 _) = typename "TexInt16"+cIntegralTypeTex (TypeInt32 _) = typename "TexInt32"+cIntegralTypeTex (TypeInt64 _) = typename "TexInt64"+cIntegralTypeTex (TypeWord8 _) = typename "TexWord8"+cIntegralTypeTex (TypeWord16 _) = typename "TexWord16"+cIntegralTypeTex (TypeWord32 _) = typename "TexWord32"+cIntegralTypeTex (TypeWord64 _) = typename "TexWord64"+cIntegralTypeTex (TypeCShort _) = typename "TexCShort"+cIntegralTypeTex (TypeCUShort _) = typename "TexCUShort"+cIntegralTypeTex (TypeCInt _) = typename "TexCInt"+cIntegralTypeTex (TypeCUInt _) = typename "TexCUInt"+cIntegralTypeTex (TypeCLong _) = typename "TexCLong"+cIntegralTypeTex (TypeCULong _) = typename "TexCULong"+cIntegralTypeTex (TypeCLLong _) = typename "TexCLLong"+cIntegralTypeTex (TypeCULLong _) = typename "TexCULLong"+cIntegralTypeTex (TypeInt _) = typename ("TexInt" ++ show (finiteBitSize (undefined::Int)))+cIntegralTypeTex (TypeWord _) = typename ("TexWord" ++ show (finiteBitSize (undefined::Word))) -codegenFloatingTex :: FloatingType a -> C.Type-codegenFloatingTex (TypeFloat _) = typename "TexFloat"-codegenFloatingTex (TypeCFloat _) = typename "TexCFloat"-codegenFloatingTex (TypeDouble _) = typename "TexDouble"-codegenFloatingTex (TypeCDouble _) = typename "TexCDouble"+cFloatingTypeTex :: FloatingType a -> C.Type+cFloatingTypeTex (TypeFloat _) = typename "TexFloat"+cFloatingTypeTex (TypeCFloat _) = typename "TexCFloat"+cFloatingTypeTex (TypeDouble _) = typename "TexDouble"+cFloatingTypeTex (TypeCDouble _) = typename "TexCDouble" -codegenNonNumTex :: NonNumType a -> C.Type-codegenNonNumTex (TypeBool _) = typename "TexWord8"-codegenNonNumTex (TypeChar _) = typename "TexWord32"-codegenNonNumTex (TypeCChar _) = typename "TexCChar"-codegenNonNumTex (TypeCSChar _) = typename "TexCSChar"-codegenNonNumTex (TypeCUChar _) = typename "TexCUChar"+cNonNumTypeTex :: NonNumType a -> C.Type+cNonNumTypeTex (TypeBool _) = typename "TexWord8"+cNonNumTypeTex (TypeChar _) = typename "TexWord32"+cNonNumTypeTex (TypeCChar _) = typename "TexCChar"+cNonNumTypeTex (TypeCSChar _) = typename "TexCSChar"+cNonNumTypeTex (TypeCUChar _) = typename "TexCUChar"+++-- Utilities+-- ---------++typename :: String -> C.Type+typename name = [cty| typename $id:name |]++signedIntegralNum :: IntegralType a -> Bool+signedIntegralNum t =+ case t of+ TypeInt _ -> True+ TypeInt8 _ -> True+ TypeInt16 _ -> True+ TypeInt32 _ -> True+ TypeInt64 _ -> True+ TypeCShort _ -> True+ TypeCInt _ -> True+ TypeCLong _ -> True+ TypeCLLong _ -> True+ _ -> False++unsignedIntegralNum :: IntegralType a -> Bool+unsignedIntegralNum = not . signedIntegralNum
Data/Array/Accelerate/CUDA/Compile.hs view
@@ -19,13 +19,13 @@ module Data.Array.Accelerate.CUDA.Compile ( -- * generate and compile kernels to realise a computation- compileAcc, compileAfun+ compileAcc, compileAfun, ) where -- friends import Data.Array.Accelerate.Error-import Data.Array.Accelerate.Tuple+import Data.Array.Accelerate.Lifetime import Data.Array.Accelerate.Trafo import Data.Array.Accelerate.CUDA.AST import Data.Array.Accelerate.CUDA.State@@ -35,7 +35,7 @@ import Data.Array.Accelerate.CUDA.Analysis.Launch import Data.Array.Accelerate.CUDA.Foreign.Import ( canExecuteAcc, canExecuteExp ) import Data.Array.Accelerate.CUDA.Persistent as KT-import qualified Data.Array.Accelerate.CUDA.FullList as FL+import qualified Data.Array.Accelerate.FullList as FL import qualified Data.Array.Accelerate.CUDA.Debug as D -- libraries@@ -58,7 +58,6 @@ import System.IO import System.IO.Error import System.IO.Unsafe-import System.Mem.Weak import System.Process import Text.PrettyPrint.Mainland ( ppr, renderCompact, displayLazyText ) import qualified Data.ByteString as B@@ -80,7 +79,7 @@ #if defined(UNIX) import System.Posix.Process #elif defined(WIN32)-import System.Win32.Process hiding (ProcessHandle)+import System.Win32.Process hiding ( ProcessHandle ) #else #error "I don't know what operating system I am" #endif@@ -134,7 +133,7 @@ Aprj ix tup -> node =<< liftA (Aprj ix) <$> travA tup -- Foreign- Aforeign ff afun a -> node =<< foreignA ff afun a+ Aforeign ff afun a -> foreignA ff afun a -- Array injection Unit e -> node =<< liftA Unit <$> travE e@@ -168,6 +167,9 @@ Stencil2 f b1 a1 b2 a2 -> exec =<< liftA3 stencil2 <$> travF f <*> travA a1 <*> travA a2 where stencil2 f' a1' a2' = Stencil2 f' b1 a1' b2 a2' + -- Loops+ -- Collect l -> ExecSeq <$> compileOpenSeq l+ where use :: ArraysR a -> a -> CIO () use ArraysRunit () = return ()@@ -198,6 +200,10 @@ travAtup NilAtup = return (pure NilAtup) travAtup (SnocAtup t a) = liftA2 SnocAtup <$> travAtup t <*> travA a + travE :: DelayedOpenExp env aenv e+ -> CIO (Free aenv, PreOpenExp ExecOpenAcc env aenv e)+ travE = compileOpenExp+ travF :: DelayedOpenFun env aenv t -> CIO (Free aenv, PreOpenFun ExecOpenAcc env aenv t) travF (Body b) = liftA Body <$> travE b travF (Lam f) = liftA Lam <$> travF f@@ -210,109 +216,211 @@ fullOfList [x] = FL.singleton () x fullOfList (x:xs) = FL.cons () x (fullOfList xs) - -- If it is a foreign call for the CUDA backend, don't bother compiling- -- the pure version+ -- If the foreign function targets this backend, drop the remaining+ -- alternatives from the AST. Similarly, we drop the foreign node if it+ -- does not target this backend. --- foreignA :: (Arrays a, Arrays r, Foreign f)- => f a r- -> DelayedAfun (a -> r)- -> DelayedOpenAcc aenv a- -> CIO (Free aenv, PreOpenAcc ExecOpenAcc aenv r)- foreignA ff afun a = case canExecuteAcc ff of- Nothing -> liftA2 (Aforeign ff) <$> pure <$> compileAfun afun <*> travA a- Just _ -> liftA (Aforeign ff err) <$> travA a+ foreignA :: (Arrays as, Arrays bs, Foreign asm)+ => asm (as -> bs)+ -> DelayedAfun (as -> bs)+ -> DelayedOpenAcc aenv as+ -> CIO (ExecOpenAcc aenv bs)+ foreignA ff afun a =+ case canExecuteAcc ff of+ Nothing -> traverseAcc $ Manifest (Apply (weaken absurd afun) a)+ Just{} -> node =<< liftA (Aforeign ff err) <$> travA a where- err = $internalError "compile" "Executing pure version of a CUDA foreign function"+ absurd :: Idx () t -> Idx env t+ absurd = absurd+ err = $internalError "Aforeign" "failed to recover foreign function a second time" - -- Traverse a scalar expression+-- Traverse a scalar expression+--+compileOpenExp+ :: DelayedOpenExp env aenv e+ -> CIO (Free aenv, PreOpenExp ExecOpenAcc env aenv e)+compileOpenExp topExp =+ case topExp of+ Var ix -> return $ pure (Var ix)+ Const c -> return $ pure (Const c)+ PrimConst c -> return $ pure (PrimConst c)+ IndexAny -> return $ pure IndexAny+ IndexNil -> return $ pure IndexNil+ Foreign ff f x -> foreignE ff f x --+ Let a b -> liftA2 Let <$> travE a <*> travE b+ IndexCons t h -> liftA2 IndexCons <$> travE t <*> travE h+ IndexHead h -> liftA IndexHead <$> travE h+ IndexTail t -> liftA IndexTail <$> travE t+ IndexSlice slix x s -> liftA2 (IndexSlice slix) <$> travE x <*> travE s+ IndexFull slix x s -> liftA2 (IndexFull slix) <$> travE x <*> travE s+ ToIndex s i -> liftA2 ToIndex <$> travE s <*> travE i+ FromIndex s i -> liftA2 FromIndex <$> travE s <*> travE i+ Tuple t -> liftA Tuple <$> travT t+ Prj ix e -> liftA (Prj ix) <$> travE e+ Cond p t e -> liftA3 Cond <$> travE p <*> travE t <*> travE e+ While p f x -> liftA3 While <$> travF p <*> travF f <*> travE x+ PrimApp f e -> liftA (PrimApp f) <$> travE e+ Index a e -> liftA2 Index <$> travA a <*> travE e+ LinearIndex a e -> liftA2 LinearIndex <$> travA a <*> travE e+ Shape a -> liftA Shape <$> travA a+ ShapeSize e -> liftA ShapeSize <$> travE e+ Intersect x y -> liftA2 Intersect <$> travE x <*> travE y+ Union x y -> liftA2 Union <$> travE x <*> travE y++ where+ travA :: (Shape sh, Elt e)+ => DelayedOpenAcc aenv (Array sh e)+ -> CIO (Free aenv, ExecOpenAcc aenv (Array sh e))+ travA a = do+ a' <- compileOpenAcc a+ return $ (bind a', a')++ travT :: Tuple (DelayedOpenExp env aenv) t+ -> CIO (Free aenv, Tuple (PreOpenExp ExecOpenAcc env aenv) t)+ travT NilTup = return (pure NilTup)+ travT (SnocTup t e) = liftA2 SnocTup <$> travT t <*> travE e+ travE :: DelayedOpenExp env aenv e -> CIO (Free aenv, PreOpenExp ExecOpenAcc env aenv e)- travE exp =- case exp of- Var ix -> return $ pure (Var ix)- Const c -> return $ pure (Const c)- PrimConst c -> return $ pure (PrimConst c)- IndexAny -> return $ pure IndexAny- IndexNil -> return $ pure IndexNil- Foreign ff f x -> foreignE ff f x- --- Let a b -> liftA2 Let <$> travE a <*> travE b- IndexCons t h -> liftA2 IndexCons <$> travE t <*> travE h- IndexHead h -> liftA IndexHead <$> travE h- IndexTail t -> liftA IndexTail <$> travE t- IndexSlice slix x s -> liftA2 (IndexSlice slix) <$> travE x <*> travE s- IndexFull slix x s -> liftA2 (IndexFull slix) <$> travE x <*> travE s- ToIndex s i -> liftA2 ToIndex <$> travE s <*> travE i- FromIndex s i -> liftA2 FromIndex <$> travE s <*> travE i- Tuple t -> liftA Tuple <$> travT t- Prj ix e -> liftA (Prj ix) <$> travE e- Cond p t e -> liftA3 Cond <$> travE p <*> travE t <*> travE e- While p f x -> liftA3 While <$> travF p <*> travF f <*> travE x- PrimApp f e -> liftA (PrimApp f) <$> travE e- Index a e -> liftA2 Index <$> travA a <*> travE e- LinearIndex a e -> liftA2 LinearIndex <$> travA a <*> travE e- Shape a -> liftA Shape <$> travA a- ShapeSize e -> liftA ShapeSize <$> travE e- Intersect x y -> liftA2 Intersect <$> travE x <*> travE y-- where- travA :: (Shape sh, Elt e)- => DelayedOpenAcc aenv (Array sh e)- -> CIO (Free aenv, ExecOpenAcc aenv (Array sh e))- travA a = do- a' <- traverseAcc a- return $ (bind a', a')+ travE = compileOpenExp - travT :: Tuple (DelayedOpenExp env aenv) t- -> CIO (Free aenv, Tuple (PreOpenExp ExecOpenAcc env aenv) t)- travT NilTup = return (pure NilTup)- travT (SnocTup t e) = liftA2 SnocTup <$> travT t <*> travE e+ travF :: DelayedOpenFun env aenv t -> CIO (Free aenv, PreOpenFun ExecOpenAcc env aenv t)+ travF (Body b) = liftA Body <$> travE b+ travF (Lam f) = liftA Lam <$> travF f - travF :: DelayedOpenFun env aenv t -> CIO (Free aenv, PreOpenFun ExecOpenAcc env aenv t)- travF (Body b) = liftA Body <$> travE b- travF (Lam f) = liftA Lam <$> travF f+ foreignE :: (Elt a, Elt b, Foreign asm)+ => asm (a -> b)+ -> DelayedFun () (a -> b)+ -> DelayedOpenExp env aenv a+ -> CIO (Free aenv, PreOpenExp ExecOpenAcc env aenv b)+ foreignE ff f x = case canExecuteExp ff of+ -- If it's a foreign function that we can generate code from, just+ -- leave it alone. As the pure function is closed, the array+ -- environment needs to be replaced with one of the right type.+ --+ Just _ -> liftA2 (Foreign ff) <$> pure <$> snd <$> travF f <*> travE x - foreignE :: (Elt a, Elt b, Foreign f)- => f a b- -> DelayedFun () (a -> b)- -> DelayedOpenExp env aenv a- -> CIO (Free aenv, PreOpenExp ExecOpenAcc env aenv b)- foreignE ff f x = case canExecuteExp ff of- -- If it's a foreign function that we can generate code from, just- -- leave it alone. As the pure function is closed, the array- -- environment needs to be replaced with one of the right type.+ -- If the foreign function is not intended for this backend, this node+ -- needs to be replaced by a pure accelerate node giving the same+ -- result. Due to the lack of an 'apply' node in the scalar language,+ -- this is done by substitution.+ --+ Nothing -> travE (apply f x)+ where+ -- Twiddle the environment variables --- Just _ -> liftA2 (Foreign ff) <$> pure <$> snd <$> travF f <*> travE x+ apply :: DelayedFun () (a -> b) -> DelayedOpenExp env aenv a -> DelayedOpenExp env aenv b+ apply (Lam (Body b)) e = Let e $ weaken wAcc $ weakenE wExp b+ apply _ _ = error "This was a triumph." - -- If the foreign function is not intended for this backend, this node- -- needs to be replaced by a pure accelerate node giving the same- -- result. Due to the lack of an 'apply' node in the scalar language,- -- this is done by substitution.+ -- As the expression we want to weaken is closed with respect to the array+ -- environment, the index manipulation function becomes a dummy argument. --- Nothing -> travE (apply f x)- where- -- Twiddle the environment variables- --- apply :: DelayedFun () (a -> b) -> DelayedOpenExp env aenv a -> DelayedOpenExp env aenv b- apply (Lam (Body b)) e = Let e $ weakenEA rebuildAcc wAcc $ weakenE wExp b- apply _ _ = error "This was a triumph."+ wAcc :: Idx () t -> Idx aenv t+ wAcc _ = error "I'm making a note here:" - -- As the expression we want to weaken is closed with respect to the array- -- environment, the index manipulation function becomes a dummy argument.+ wExp :: Idx ((),a) t -> Idx (env,a) t+ wExp ZeroIdx = ZeroIdx+ wExp _ = error "HUGE SUCCESS"++ bind :: (Shape sh, Elt e) => ExecOpenAcc aenv (Array sh e) -> Free aenv+ bind (ExecAcc _ _ (Avar ix)) = freevar ix+ bind _ = $internalError "bind" "expected array variable"++{--+compileSeq :: DelayedSeq a -> CIO (ExecSeq a)+compileSeq (DelayedSeq aenv s) = ExecS <$> compileExtend aenv <*> compileOpenSeq s+ where+ compileExtend :: Extend DelayedOpenAcc aenv aenv' -> CIO (Extend ExecOpenAcc aenv aenv')+ compileExtend BaseEnv = return BaseEnv+ compileExtend (PushEnv e a) = PushEnv <$> compileExtend e <*> compileOpenAcc a++compileOpenSeq+ :: forall aenv lenv arrs'.+ PreOpenSeq DelayedOpenAcc aenv lenv arrs'+ -> CIO (ExecOpenSeq aenv lenv arrs')+compileOpenSeq l =+ case l of+ Producer p l' -> ExecP <$> compileP p <*> compileOpenSeq l'+ Consumer c -> ExecC <$> compileC c+ Reify ix -> return $ ExecR ix Nothing+ where+ compileP :: forall a. Producer DelayedOpenAcc aenv lenv a -> CIO (ExecP aenv lenv a)+ compileP p =+ case p of+ ToSeq slix (_ :: proxy slix) acc -> do+ case acc of+ -- In the case of converting an array that has not already been copied+ -- to device memory, we are smart and treat it specially.+ Manifest (Use a) -> return $ ExecUseLazy slix (toArr a) ([] :: [slix])+ _ -> do+ (free1, acc') <- travA acc+ let gamma = makeEnvMap free1+ dev <- asks deviceProperties+ -- The array computation passed to 'toSeq' needs to be treated+ -- specially. We don't want the entire array to be made manifest+ -- if we can help it. In the event it is a delayed array, we make+ -- the subarrays manifest one at a time and feed them to the 'Seq'+ -- computation. --- wAcc :: Idx () t -> Idx aenv t- wAcc _ = error "I'm making a note here:"+ -- For the purposes of device configuration and launching, this can+ -- be seen to work like 'Slice', even though in reality it+ -- resembles a delayed 'Slice'.+ let acc'' = Manifest (Slice slix acc (Const (zeroSlice slix) :: DelayedExp aenv slix)) - wExp :: Idx ((),a) t -> Idx (env,a) t- wExp ZeroIdx = ZeroIdx- wExp _ = error "HUGE SUCCESS"+ kernel <- build1 acc'' (codegenToSeq slix dev acc gamma)+ return $ ExecToSeq slix acc' kernel gamma ([] :: [slix])+ StreamIn xs -> return $ ExecStreamIn xs+ MapSeq f x -> do+ f' <- compileOpenAfun f+ return $ ExecMap f' x+ ZipWithSeq f x y -> do+ f' <- compileOpenAfun f+ return $ ExecZipWith f' x y+ ScanSeq f a0 x -> do+ (_, a0') <- travE a0+ (_, f') <- travF f+ return $ ExecScanSeq f' a0' x Nothing+ ChunkedMapSeq{} -> error "TODO: @fmma needs to finish this..." - bind :: (Shape sh, Elt e) => ExecOpenAcc aenv (Array sh e) -> Free aenv- bind (ExecAcc _ _ (Avar ix)) = freevar ix- bind _ = $internalError "bind" "expected array variable"+ compileC :: forall a. Consumer DelayedOpenAcc aenv lenv a -> CIO (ExecC aenv lenv a)+ compileC c =+ case c of+ FoldSeq f a0 x -> do+ (_, a0') <- travE a0+ (_, f') <- travF f+ return $ ExecFoldSeq f' a0' x Nothing+ FoldSeqFlatten f acc x -> do+ acc' <- compileOpenAcc acc+ f' <- compileOpenAfun f+ return $ ExecFoldSeqFlatten f' acc' x Nothing+ Stuple t -> ExecStuple <$> compileCT t + compileCT :: forall t. Atuple (Consumer DelayedOpenAcc aenv lenv) t -> CIO (Atuple (ExecC aenv lenv) t)+ compileCT NilAtup = return NilAtup+ compileCT (SnocAtup t c) = SnocAtup <$> compileCT t <*> compileC c + travA :: DelayedOpenAcc aenv a -> CIO (Free aenv, ExecOpenAcc aenv a)+ travA acc = case acc of+ Manifest{} -> pure <$> compileOpenAcc acc+ Delayed{..} -> liftA2 (const EmbedAcc) <$> travF indexD <*> travE extentD++ travE :: DelayedOpenExp env aenv e+ -> CIO (Free aenv, PreOpenExp ExecOpenAcc env aenv e)+ travE = compileOpenExp++ travF :: DelayedOpenFun env aenv t -> CIO (Free aenv, PreOpenFun ExecOpenAcc env aenv t)+ travF (Body b) = liftA Body <$> travE b+ travF (Lam f) = liftA Lam <$> travF f++ zeroSlice :: SliceIndex slix sl co sh -> slix+ zeroSlice SliceNil = ()+ zeroSlice (SliceFixed sl) = (zeroSlice sl, 0)+ zeroSlice (SliceAll sl) = (zeroSlice sl, ())+--}++ -- Applicative -- ----------- --@@ -342,7 +450,7 @@ let (cta,blocks,smem) = launchConfig acc dev occ (mdl,fun,occ) = unsafePerformIO $ do m <- link context table key- f <- CUDA.getFun m entry+ f <- withLifetime m $ flip CUDA.getFun entry l <- CUDA.requires f CUDA.MaxKernelThreadsPerBlock o <- determineOccupancy acc dev f l D.when D.dump_cc (stats entry f o)@@ -366,14 +474,14 @@ -- make sure kernel/stats are printed together. Use 'intercalate' rather -- than 'unlines' to avoid a trailing newline. --- message $ intercalate "\n ... " [msg1, msg2]+ message $ intercalate "\n ... " [msg1, msg2] -- Link a compiled binary and update the associated kernel entry in the hash -- table. This may entail waiting for the external compilation process to -- complete. If successful, the temporary files are removed. ---link :: Context -> KernelTable -> KernelKey -> IO CUDA.Module+link :: Context -> KernelTable -> KernelKey -> IO (Lifetime CUDA.Module) link context table key = let intErr = $internalError "link" "missing kernel entry" ctx = deviceContext context@@ -395,12 +503,13 @@ () <- takeMVar done bin <- B.readFile cubin mdl <- CUDA.loadData bin- addFinalizer mdl (module_finalizer weak_ctx key mdl)+ lmdl <- newLifetime mdl+ addFinalizer lmdl (module_finalizer weak_ctx key lmdl) -- Update hash tables and stash the binary object into the persistent -- cache --- KT.insert table key $! KernelObject bin (FL.singleton ctx mdl)+ KT.insert table key $! KernelObject bin (FL.singleton ctx lmdl) KT.persist table cubin key -- Remove temporary build products.@@ -412,25 +521,30 @@ removeDirectory (dropFileName cufile) `catchIOError` \_ -> return () -- directory not empty - return mdl+ return lmdl -- If we get a real object back, then this will already be in the -- persistent cache, since either it was just read in from there, or we -- had to generate new code and the link step above has added it. -- KernelObject bin active- | Just mdl <- FL.lookup ctx active -> return mdl+ | Just lmdl <- FL.lookup ctx active -> return lmdl | otherwise -> do message "re-linking module for current context" mdl <- CUDA.loadData bin- addFinalizer mdl (module_finalizer weak_ctx key mdl)- KT.insert table key $! KernelObject bin (FL.cons ctx mdl active)- return mdl+ lmdl <- newLifetime mdl+ addFinalizer lmdl (module_finalizer weak_ctx key lmdl)+ KT.insert table key $! KernelObject bin (FL.cons ctx lmdl active)+ return lmdl -- Generate and compile code for a single open array expression ---compile :: KernelTable -> CUDA.DeviceProperties -> CUTranslSkel aenv a -> CIO (String, KernelKey)+compile+ :: KernelTable+ -> CUDA.DeviceProperties+ -> CUTranslSkel aenv a+ -> CIO (String, KernelKey) compile table dev cunit = do context <- asks activeContext exists <- isJust `fmap` liftIO (KT.lookup context table key)@@ -459,6 +573,8 @@ compileFlags cufile = do CUDA.Compute m n <- CUDA.computeCapability `fmap` asks deviceProperties ddir <- liftIO getDataDir+ warnings <- liftIO $ (&&) <$> D.queryFlag D.dump_cc <*> D.queryFlag D.verbose+ debug <- liftIO $ D.queryFlag D.debug_cc return $ filter (not . null) $ [ "-I", ddir </> "cubits" , "-arch=sm_" ++ show m ++ show n@@ -472,8 +588,6 @@ , machine , cufile ] where- warnings = D.mode D.dump_cc && D.mode D.verbose- debug = D.mode D.debug_cc machine = case finiteBitSize (undefined :: Int) of 32 -> "-m32" 64 -> "-m64"@@ -593,9 +707,5 @@ {-# INLINE message #-} message :: MonadIO m => String -> m ()-message msg = trace msg $ return ()--{-# INLINE trace #-}-trace :: MonadIO m => String -> m a -> m a-trace msg next = D.message D.dump_cc ("cc: " ++ msg) >> next+message msg = liftIO $ D.traceIO D.dump_cc ("cc: " ++ msg)
Data/Array/Accelerate/CUDA/Context.hs view
@@ -18,22 +18,21 @@ -- An execution context Context(..), create, push, pop, destroy,- keepAlive, fromDeviceContext+ keepAlive, fromDeviceContext, ) where -- friends-import Data.Array.Accelerate.CUDA.Debug ( message, verbose, dump_gc, showFFloatSIBase )+import Data.Array.Accelerate.CUDA.Debug ( traceIO, verbose, dump_gc, showFFloatSIBase ) import Data.Array.Accelerate.CUDA.Analysis.Device+import Data.Array.Accelerate.Lifetime -- system import Data.Function ( on ) import Control.Exception ( bracket_ ) import Control.Concurrent ( forkIO, threadDelay ) import Control.Monad ( when )-import GHC.Exts ( Ptr(..), mkWeak# )-import GHC.Base ( IO(..) )-import GHC.Weak ( Weak(..) )+import System.Mem.Weak ( Weak ) import Text.PrettyPrint import qualified Foreign.CUDA.Driver as CUDA hiding ( device ) import qualified Foreign.CUDA.Driver.Context as CUDA@@ -43,21 +42,21 @@ -- | The execution context -- data Context = Context {- deviceProperties :: {-# UNPACK #-} !CUDA.DeviceProperties, -- information on hardware resources- deviceContext :: {-# UNPACK #-} !CUDA.Context, -- device execution context- weakContext :: {-# UNPACK #-} !(Weak CUDA.Context) -- weak pointer to the context (for memory management)+ deviceProperties :: {-# UNPACK #-} !CUDA.DeviceProperties, -- information on hardware resources+ deviceContext :: {-# UNPACK #-} !(Lifetime CUDA.Context), -- device execution context+ weakContext :: {-# UNPACK #-} !(Weak (Lifetime CUDA.Context)) -- weak pointer to the context (for memory management) } instance Eq Context where (==) = (==) `on` deviceContext - -- | Create a new CUDA context associated with the calling thread -- create :: CUDA.Device -> [CUDA.ContextFlag] -> IO Context create dev flags = do- ctx <- CUDA.create dev flags >> CUDA.pop >>= keepAlive+ ctx <- CUDA.create dev flags >> CUDA.pop actx@(Context prp _ _) <- fromDeviceContext dev ctx+ _ <- keepAlive actx -- Generated code does not take particular advantage of shared memory, so -- for devices that support it use those banks as an L1 cache instead.@@ -68,7 +67,7 @@ when (CUDA.computeCapability prp >= CUDA.Compute 2 0) $ bracket_ (CUDA.push ctx) CUDA.pop (CUDA.setCache CUDA.PreferL1) - message verbose (deviceInfo dev prp)+ traceIO verbose (deviceInfo dev prp) return actx -- |Given a device context, construct a new context around it.@@ -76,21 +75,22 @@ fromDeviceContext :: CUDA.Device -> CUDA.Context -> IO Context fromDeviceContext dev ctx = do prp <- CUDA.props dev- weak <- mkWeakContext ctx $ do- message dump_gc $ "gc: finalise context #" ++ show (CUDA.useContext ctx)+ lctx <- newLifetime ctx+ addFinalizer lctx $ do+ traceIO dump_gc $ "gc: finalise context #" ++ show (CUDA.useContext ctx)+ CUDA.push ctx CUDA.destroy ctx- message dump_gc $ "gc: initialise context #" ++ show (CUDA.useContext ctx)+ weak <- mkWeakPtr lctx+ traceIO dump_gc $ "gc: initialise context #" ++ show (CUDA.useContext ctx) - return $! Context prp ctx weak+ return $! Context prp lctx weak -- | Destroy the specified context. This will fail if the context is more than -- single attachment. -- {-# INLINE destroy #-} destroy :: Context -> IO ()-destroy (deviceContext -> ctx) = do- message dump_gc ("gc: destroy context: #" ++ show (CUDA.useContext ctx))- CUDA.destroy ctx+destroy (deviceContext -> ctx) = finalize ctx -- | Push the given context onto the CPU's thread stack of current contexts. The@@ -98,8 +98,8 @@ -- {-# INLINE push #-} push :: Context -> IO ()-push (deviceContext -> ctx) = do- message dump_gc ("gc: push context: #" ++ show (CUDA.useContext ctx))+push (deviceContext -> lctx) = withLifetime lctx $ \ctx -> do+ traceIO dump_gc ("gc: push context: #" ++ show (CUDA.useContext ctx)) CUDA.push ctx @@ -109,17 +109,7 @@ pop :: IO () pop = do ctx <- CUDA.pop- message dump_gc ("gc: pop context: #" ++ show (CUDA.useContext ctx))----- Make a weak pointer to a CUDA context. We need to be careful to put the--- finaliser on the underlying pointer, rather than the box around it as--- 'mkWeak' will do, because unpacking the context will cause the finaliser to--- fire prematurely.----mkWeakContext :: CUDA.Context -> IO () -> IO (Weak CUDA.Context)-mkWeakContext c@(CUDA.Context (Ptr c#)) f = IO $ \s ->- case mkWeak# c# c f s of (# s', w #) -> (# s', Weak w #)+ traceIO dump_gc ("gc: pop context: #" ++ show (CUDA.useContext ctx)) -- Make sure the GC knows that we want to keep this thing alive past the end of
Data/Array/Accelerate/CUDA/Debug.hs view
@@ -1,9 +1,6 @@-{-# LANGUAGE CPP #-}-{-# LANGUAGE TemplateHaskell #-}-{-# LANGUAGE TypeOperators #-}-{-# OPTIONS -fno-warn-incomplete-patterns #-}-{-# OPTIONS -fno-warn-unused-binds #-}-{-# OPTIONS -fno-warn-unused-imports #-}+{-# LANGUAGE CPP #-}+{-# OPTIONS -fno-warn-unused-binds #-}+{-# OPTIONS -fno-warn-unused-imports #-} -- | -- Module : Data.Array.Accelerate.CUDA.Debug -- Copyright : [2008..2014] Manuel M T Chakravarty, Gabriele Keller@@ -20,185 +17,47 @@ module Data.Array.Accelerate.CUDA.Debug ( - showFFloatSIBase,-- message, trace, event, when, unless, mode, timed, elapsed,- verbose, flush_cache,- dump_gc, dump_cc, debug_cc, dump_exec,+ module Data.Array.Accelerate.Debug,+ module Data.Array.Accelerate.CUDA.Debug, ) where -import Numeric-import Data.List-import Data.Label-import Data.IORef-import Debug.Trace ( traceIO, traceEventIO )+import Data.Array.Accelerate.Debug hiding ( timed, elapsed )++import Control.Concurrent ( forkIO ) import Control.Monad ( void ) import Control.Monad.IO.Class ( liftIO, MonadIO )-import Control.Concurrent ( forkIO )+import GHC.Float ( float2Double ) import System.CPUTime 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 ---- -------------------------------------------------------------------------------- Pretty-printing--showFFloatSIBase :: RealFloat a => Maybe Int -> a -> a -> ShowS-showFFloatSIBase p b n- = showString- $ showFFloat p n' (' ':si_unit)- where- n' = n / (b ^^ pow)- pow = (-4) `max` floor (logBase b n) `min` 4 :: Int- si_unit = case pow of- -4 -> "p"- -3 -> "n"- -2 -> "µ"- -1 -> "m"- 0 -> ""- 1 -> "k"- 2 -> "M"- 3 -> "G"- 4 -> "T"----- -------------------------------------------------------------------------------- Internals--data Flags = Flags- {- -- debugging- _dump_gc :: !Bool -- garbage collection & memory management- , _dump_cc :: !Bool -- compilation & linking- , _debug_cc :: !Bool -- compile device code with debug symbols- , _dump_exec :: !Bool -- kernel execution- , _verbose :: !Bool -- additional status messages-- -- general options / functionality- , _flush_cache :: !Bool -- delete the persistent cache directory- , _fast_math :: !Bool -- use faster, less accurate maths library operations- }--$(mkLabels [''Flags])--allFlags :: [OptDescr (Flags -> Flags)]-allFlags =- [- -- debugging- Option [] ["ddump-gc"] (NoArg (set dump_gc True)) "print device memory management trace"- , Option [] ["ddump-cc"] (NoArg (set dump_cc True)) "print generated code and compilation information"- , Option [] ["ddebug-cc"] (NoArg (set debug_cc True)) "generate debug information for device code"- , Option [] ["ddump-exec"] (NoArg (set dump_exec True)) "print kernel execution trace"- , Option [] ["dverbose"] (NoArg (set verbose True)) "print additional information"-- -- functionality / optimisation- , Option [] ["fflush-cache"] (NoArg (set flush_cache True)) "delete the persistent cache directory"- , Option [] ["ffast-math"] (NoArg (set fast_math True)) "use faster, less accurate maths library operations"- ]--initialise :: IO Flags-initialise = parse `fmap` getArgs- where- defaults = Flags False False False False False False False- parse = foldl parse1 defaults- parse1 opts x = case filter (\(Option _ [f] _ _) -> x `isPrefixOf` ('-':f)) allFlags of- [Option _ _ (NoArg go) _] -> go opts- _ -> opts -- not specified, or ambiguous--#ifdef ACCELERATE_DEBUG-{-# NOINLINE options #-}-options :: IORef Flags-options = unsafePerformIO $ newIORef =<< initialise-#endif--{-# INLINE mode #-}-mode :: (Flags :-> Bool) -> Bool-#ifdef ACCELERATE_DEBUG-mode f = unsafePerformIO $ get f `fmap` readIORef options-#else-mode _ = False-#endif--{-# INLINE message #-}-message :: MonadIO m => (Flags :-> Bool) -> String -> m ()-#ifdef ACCELERATE_DEBUG-message f str- = when f . liftIO- $ do psec <- getCPUTime- let sec = fromIntegral psec * 1E-12 :: Double- traceIO $ showFFloat (Just 2) sec (':':str)-#else-message _ _ = return ()-#endif--{-# INLINE event #-}-event :: MonadIO m => (Flags :-> Bool) -> String -> m ()-#ifdef ACCELERATE_DEBUG-event f str = when f (liftIO $ traceEventIO str)-#else-event _ _ = return ()-#endif--{-# INLINE trace #-}-trace :: (Flags :-> Bool) -> String -> a -> a-#ifdef ACCELERATE_DEBUG-trace f str next = unsafePerformIO (message f str) `seq` next-#else-trace _ _ next = next-#endif---{-# INLINE when #-}-when :: MonadIO m => (Flags :-> Bool) -> m () -> m ()-#ifdef ACCELERATE_DEBUG-when f action- | mode f = action- | otherwise = return ()-#else-when _ _ = return ()-#endif--{-# INLINE unless #-}-unless :: MonadIO m => (Flags :-> Bool) -> m () -> m ()-#ifdef ACCELERATE_DEBUG-unless f action- | mode f = return ()- | otherwise = action-#else-unless _ action = action-#endif--{-# INLINE timed #-}-timed- :: MonadIO m- => (Flags :-> Bool)- -> (Double -> Double -> String)- -> Maybe Stream- -> m ()- -> m ()-timed _f _str _stream action+-- | Execute an action and time the results. The second argument specifies how+-- to format the output string given elapsed GPU and CPU time respectively+--+timed :: Mode -> (Double -> Double -> String) -> Maybe Stream -> IO () -> IO () #ifdef ACCELERATE_DEBUG- | mode _f- = do- gpuBegin <- liftIO $ Event.create []- gpuEnd <- liftIO $ Event.create []- cpuBegin <- liftIO getCPUTime- liftIO $ Event.record gpuBegin _stream+{-# NOINLINE timed #-}+timed f fmt stream action = do+ enabled <- queryFlag f+ if enabled+ then do+ gpuBegin <- Event.create []+ gpuEnd <- Event.create []+ cpuBegin <- getCPUTime+ Event.record gpuBegin stream action- liftIO $ Event.record gpuEnd _stream- cpuEnd <- liftIO getCPUTime+ Event.record gpuEnd stream+ cpuEnd <- getCPUTime -- Wait for the GPU to finish executing then display the timing execution -- message. Do this in a separate thread so that the remaining kernels can -- be queued asynchronously. --- _ <- liftIO . forkIO $ do+ void . forkIO $ do Event.block gpuEnd diff <- Event.elapsedTime gpuBegin gpuEnd let gpuTime = float2Double $ diff * 1E-3 -- milliseconds@@ -207,13 +66,14 @@ Event.destroy gpuBegin Event.destroy gpuEnd --- message _f (_str gpuTime cpuTime)- --- return ()+ traceIO f (fmt gpuTime cpuTime) - | otherwise+ else+ action+#else+{-# INLINE timed #-}+timed _ _ _ action = action #endif- = action {-# INLINE elapsed #-} elapsed :: Double -> Double -> String
Data/Array/Accelerate/CUDA/Execute.hs view
@@ -7,6 +7,7 @@ {-# LANGUAGE PatternGuards #-} {-# LANGUAGE ScopedTypeVariables #-} {-# LANGUAGE TemplateHaskell #-}+{-# LANGUAGE TupleSections #-} {-# LANGUAGE TypeOperators #-} {-# LANGUAGE TypeSynonymInstances #-} {-# LANGUAGE UndecidableInstances #-}@@ -24,30 +25,34 @@ module Data.Array.Accelerate.CUDA.Execute ( -- * Execute a computation under a CUDA environment- executeAcc, executeAfun1+ executeAcc, executeAfun1, + -- -- * Executing a sequence computation and streaming its output.+ -- StreamSeq(..), streamSeq,+ ) where -- friends import Data.Array.Accelerate.CUDA.AST-import Data.Array.Accelerate.CUDA.State-import Data.Array.Accelerate.CUDA.FullList ( FullList(..), List(..) )+import Data.Array.Accelerate.CUDA.Analysis.Shape import Data.Array.Accelerate.CUDA.Array.Data import Data.Array.Accelerate.CUDA.Array.Sugar-import Data.Array.Accelerate.CUDA.Foreign.Import ( canExecuteAcc ) import Data.Array.Accelerate.CUDA.CodeGen.Base ( Name, namesOfArray, groupOfInt ) import Data.Array.Accelerate.CUDA.Execute.Event ( Event ) import Data.Array.Accelerate.CUDA.Execute.Stream ( Stream )+import Data.Array.Accelerate.CUDA.Foreign.Import ( canExecuteAcc )+import Data.Array.Accelerate.CUDA.State import qualified Data.Array.Accelerate.CUDA.Array.Prim as Prim import qualified Data.Array.Accelerate.CUDA.Debug as D 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 ) import Data.Array.Accelerate.Array.Representation ( SliceIndex(..) )+import Data.Array.Accelerate.FullList ( FullList(..), List(..) )+import Data.Array.Accelerate.Lifetime ( withLifetime ) import qualified Data.Array.Accelerate.Array.Representation as R @@ -57,6 +62,7 @@ import Control.Monad.Reader ( asks ) import Control.Monad.State ( gets ) import Control.Monad.Trans ( MonadIO, liftIO )+import Control.Monad.Trans.Cont ( ContT(..) ) import System.IO.Unsafe ( unsafeInterleaveIO ) import Data.Int import Data.Word@@ -81,6 +87,8 @@ Aempty :: Aval () Apush :: Aval env -> Async t -> Aval (env, t) +-- -- A suspended sequence computation.+-- newtype StreamSeq a = StreamSeq (CIO (Maybe (a, StreamSeq a))) -- Projection of a value from a valuation using a de Bruijn index. --@@ -111,7 +119,8 @@ streaming first second = do context <- asks activeContext reservoir <- gets streamReservoir- Stream.streaming context reservoir first (\e a -> second (Async e a))+ table <- gets eventTable+ Stream.streaming context reservoir table first (\e a -> second (Async e a)) -- Array expression evaluation@@ -130,6 +139,7 @@ -- memory allocated for the result, and the kernel(s) that implement the -- skeleton are invoked --+ executeAcc :: Arrays a => ExecAcc a -> CIO a executeAcc !acc = streaming (executeOpenAcc acc Aempty) wait @@ -159,19 +169,21 @@ -> CIO arrs executeOpenAcc EmbedAcc{} _ _ = $internalError "execute" "unexpected delayed array"+-- executeOpenAcc (ExecSeq l) !aenv !stream+-- = executeSequence l aenv stream executeOpenAcc (ExecAcc (FL () kernel more) !gamma !pacc) !aenv !stream = case pacc of -- Array introduction Use arr -> return (toArr arr)- Unit x -> newArray Z . const =<< travE x+ Unit x -> fromFunction Z . const =<< travE x -- Environment manipulation Avar ix -> after stream (aprj ix aenv) 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+ Atuple tup -> toAtuple <$> travT tup+ Aprj ix tup -> evalPrj ix . fromAtuple <$> travA tup Acond p t e -> travE p >>= \x -> if x then travA t else travA e Awhile p f a -> awhile p f =<< travA a @@ -200,13 +212,15 @@ Stencil _ _ a -> stencilOp =<< travA a Stencil2 _ _ a1 _ a2 -> join $ stencil2Op <$> travA a1 <*> travA a2 - -- Removed by fusion- Replicate _ _ _ -> fusionError- Slice _ _ _ -> fusionError- ZipWith _ _ _ -> fusionError+ -- AST nodes that should be inaccessible at this point+ Replicate{} -> fusionError+ Slice{} -> fusionError+ ZipWith{} -> fusionError+ -- Collect{} -> streamingError where- fusionError = $internalError "executeOpenAcc" "unexpected fusible matter"+ fusionError = $internalError "executeOpenAcc" "unexpected fusible matter"+ -- streamingError = $internalError "executeOpenAcc" "unexpected sequence computation" -- term traversals travA :: ExecOpenAcc aenv a -> CIO a@@ -221,21 +235,24 @@ awhile :: PreOpenAfun ExecOpenAcc aenv (a -> Scalar Bool) -> PreOpenAfun ExecOpenAcc aenv (a -> a) -> a -> CIO a awhile p f a = do- nop <- liftIO Event.create -- record event never call, so this is a functional no-op+ tbl <- gets eventTable+ ctx <- asks activeContext+ nop <- liftIO $ Event.create ctx tbl -- record event never call, so this is a functional no-op r <- executeOpenAfun1 p aenv (Async nop a) ok <- indexArray r 0 -- TLM TODO: memory manager should remember what is already on the host if ok then awhile p f =<< executeOpenAfun1 f aenv (Async nop a) else return a - aforeign :: (Arrays as, Arrays bs, Foreign f) => f as bs -> PreAfun ExecOpenAcc (as -> bs) -> as -> CIO bs- aforeign ff pureFun a =+ aforeign :: (Arrays as, Arrays bs, Foreign asm) => asm (as -> bs) -> PreAfun ExecOpenAcc (as -> bs) -> as -> CIO bs+ aforeign ff next a = case canExecuteAcc ff of- Just cudaFun -> cudaFun stream a- Nothing -> executeAfun1 pureFun a+ Just asm -> asm stream a+ Nothing -> executeAfun1 next a -- get the extent of an embedded array extent :: Shape sh => ExecOpenAcc aenv (Array sh e) -> CIO sh extent ExecAcc{} = $internalError "executeOpenAcc" "expected delayed array"+ -- extent ExecSeq{} = $internalError "executeOpenAcc" "expected delayed array" extent (EmbedAcc sh) = travE sh -- Skeleton implementation@@ -273,7 +290,7 @@ foldCore :: (Shape sh, Elt e) => (sh :. Int) -> CIO (Array sh e) foldCore !(!sh :. sz)- | dim sh > 0 = executeOp sh+ | rank sh > 0 = executeOp sh | otherwise = let !numElements = size sh * sz (_,!numBlocks,_) = configure kernel numElements@@ -308,35 +325,41 @@ -- Scans, all variations on a theme. --- scanOp :: Elt e => Bool -> (Z :. Int) -> CIO (Vector e)- scanOp !left !(Z :. numElements) = do- arr@(Array _ adata) <- allocateArray (Z :. numElements + 1)- out <- devicePtrsOfArrayData adata- let (!body, !sum)- | left = (out, advancePtrsOfArrayData adata numElements out)- | otherwise = (advancePtrsOfArrayData adata 1 out, out)- --- scanCore numElements arr body sum- return arr+ scanOp :: forall sh e. (Shape sh, Elt e) => Bool -> (sh :. Int) -> CIO (Array (sh:.Int) e)+ scanOp !left !(_ :. numElements)+ | Just Refl <- matchShapeType (undefined::sh) (undefined::Z)+ = do+ arr@(Array _ adata) <- allocateArray (Z :. numElements + 1)+ withDevicePtrs adata (Just stream) $ \out -> do+ let (!body, !sum)+ | left = (out, advancePtrsOfArrayData adata numElements out)+ | otherwise = (advancePtrsOfArrayData adata 1 out, out)+ --+ scanCore numElements arr body sum+ return arr - scan1Op :: forall e. Elt e => (Z :. Int) -> CIO (Vector e)- scan1Op !(Z :. numElements) = do- arr@(Array _ adata) <- allocateArray (Z :. numElements + 1) :: CIO (Vector e)- body <- devicePtrsOfArrayData adata- let sum {- to fix type -} = advancePtrsOfArrayData adata numElements body- --- scanCore numElements arr body sum- return (Array ((),numElements) adata)+ scan1Op :: forall sh e. (Shape sh, Elt e) => (sh :. Int) -> CIO (Array (sh:.Int) e)+ scan1Op !(_ :. numElements)+ | Just Refl <- matchShapeType (undefined::sh) (undefined::Z)+ = do+ arr@(Array _ adata) <- allocateArray (Z :. numElements + 1) :: CIO (Vector e)+ withDevicePtrs adata (Just stream) $ \body -> do+ let sum {- to fix type -} = advancePtrsOfArrayData adata numElements body+ --+ scanCore numElements arr body sum+ return (Array ((),numElements) adata) - scan'Op :: forall e. Elt e => (Z :. Int) -> CIO (Vector e, Scalar e)- scan'Op !(Z :. numElements) = do- vec@(Array _ ad_vec) <- allocateArray (Z :. numElements) :: CIO (Vector e)- sum@(Array _ ad_sum) <- allocateArray Z :: CIO (Scalar e)- d_vec <- devicePtrsOfArrayData ad_vec- d_sum <- devicePtrsOfArrayData ad_sum- --- scanCore numElements vec d_vec d_sum- return (vec, sum)+ scan'Op :: forall sh e. (Shape sh, Elt e) => (sh :. Int) -> CIO (Array (sh:.Int) e, Array sh e)+ scan'Op !(_ :. numElements)+ | Just Refl <- matchShapeType (undefined::sh) (undefined::Z)+ = do+ vec@(Array _ ad_vec) <- allocateArray (Z :. numElements) :: CIO (Vector e)+ sum@(Array _ ad_sum) <- allocateArray Z :: CIO (Scalar e)+ withDevicePtrs ad_vec (Just stream) $ \d_vec ->+ withDevicePtrs ad_sum (Just stream) $ \d_sum -> do+ --+ scanCore numElements vec d_vec d_sum+ return (vec, sum) scanCore :: forall e. Elt e@@ -377,12 +400,12 @@ out <- allocateArray sh' Array _ locks <- allocateArray sh' :: CIO (Array sh' Int32)- ((), d_locks) <- devicePtrsOfArrayData locks :: CIO ((), CUDA.DevicePtr Int32)+ withDevicePtrs locks (Just stream) $ \d_locks -> do - 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+ 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 -- same as those given in the function 'mkStencil[2]'.@@ -394,9 +417,10 @@ dev <- asks deviceProperties if computeCapability dev < Compute 2 0- then marshalAccTex (namesOfArray "Stencil" (undefined :: a)) kernel arr >>- execute kernel gamma aenv (size sh) (out, sh) stream+ then marshalAccTex (namesOfArray "Stencil" (undefined :: a)) kernel arr (Just stream) $+ execute kernel gamma aenv (size sh) (out, sh) stream else execute kernel gamma aenv (size sh) (out, arr) stream+ execute kernel gamma aenv (size sh) (out, arr) stream -- return out @@ -414,10 +438,11 @@ dev <- asks deviceProperties if computeCapability dev < Compute 2 0- then marshalAccTex (namesOfArray "Stencil1" (undefined :: a)) op arr1 >>- marshalAccTex (namesOfArray "Stencil2" (undefined :: b)) op arr2 >>+ then marshalAccTex (namesOfArray "Stencil1" (undefined :: a)) op arr1 (Just stream) $+ marshalAccTex (namesOfArray "Stencil2" (undefined :: b)) op arr2 (Just stream) $ execute op gamma aenv (size sh) (out, sh1, sh2) stream else execute op gamma aenv (size sh) (out, arr1, arr2) stream+ execute op gamma aenv (size sh) (out, arr1, arr2) stream -- return out @@ -425,13 +450,269 @@ = $internalError "stencil2Op" "missing stencil specialisation kernel" +{--+-- Execute a streaming computation+--+executeSequence+ :: forall aenv arrs. Arrays arrs+ => ExecOpenSeq aenv () arrs+ -> Aval aenv+ -> Stream+ -> CIO arrs+executeSequence topSequence aenv stream+ = initializeOpenSeq topSequence aenv stream >>= loop >>= returnOut+ where+ loop :: ExecOpenSeq aenv () arrs -> CIO (ExecOpenSeq aenv () arrs)+ loop = loop'+ where+ loop' :: ExecOpenSeq aenv () arrs -> CIO (ExecOpenSeq aenv () arrs)+ loop' s = do+ ms <- runMaybeT (stepOpenSeq aenv s stream)+ case ms of+ Nothing -> return s+ Just s' -> loop' s'++ returnOut :: forall lenv arrs' . ExecOpenSeq aenv lenv arrs' -> CIO arrs'+ returnOut !l =+ case l of+ ExecP _ l' -> returnOut l'+ ExecC c -> readConsumer c+ ExecR _ ma -> case ma of+ Just a -> return [a]+ Nothing -> $internalError "executeSequence" "Expected already executed Sequence"++ where+ readConsumer :: forall a . ExecC aenv lenv a -> CIO a+ readConsumer c =+ case c of+ ExecFoldSeq _ _ _ (Just a) -> let a' = fromList Z [a] in useArray a' >> return a'+ ExecFoldSeqFlatten _ _ _ (Just a) -> return a+ ExecStuple t -> toAtuple <$> rdT t+ where+ rdT :: forall t. Atuple (ExecC aenv lenv) t -> CIO t+ rdT NilAtup = return ()+ rdT (SnocAtup t' c') = (,) <$> rdT t' <*> readConsumer c'+ _ -> $internalError "executeSequence" "Expected already executed consumer"++initializeSeq+ :: Aval aenv+ -> ExecOpenSeq aenv () arrs'+ -> CIO (ExecOpenSeq aenv () arrs')+initializeSeq aenv s = streaming (initializeOpenSeq s aenv) wait++initializeOpenSeq+ :: forall lenv aenv arrs'.+ ExecOpenSeq aenv lenv arrs'+ -> Aval aenv+ -> Stream+ -> CIO (ExecOpenSeq aenv lenv arrs')+initializeOpenSeq l aenv stream =+ case l of+ ExecP p l' -> ExecP <$> initP p <*> initializeOpenSeq l' aenv stream+ ExecC c -> ExecC <$> initC c+ ExecR ix a -> return (ExecR ix a)++ where+ initP :: forall a. ExecP aenv lenv a -> CIO (ExecP aenv lenv a)+ initP (ExecToSeq slix acc k g (_::[slix])) =+ do sh <- extent acc+ --Lazy evaluation will stop this entire list being generated.+ let slices = enumSlices slix sh :: [slix]+ return $ ExecToSeq slix acc k g slices+ initP (ExecUseLazy slix arr (_::[slix])) =+ let sh = shape arr+ -- Same as above.+ slices = enumSlices slix sh :: [slix]+ in return $ ExecUseLazy slix arr slices+ initP s@ExecStreamIn{} = return s+ initP s@ExecMap{} = return s+ initP s@ExecZipWith{} = return s+ initP (ExecScanSeq f a0 ix _) =+ do a <- executeExp a0 aenv stream+ return $ ExecScanSeq f a0 ix (Just a)++ initC :: forall a. ExecC aenv lenv a -> CIO (ExecC aenv lenv a)+ initC c =+ case c of+ ExecFoldSeq f a0 ix _ ->+ do a <- executeExp a0 aenv stream+ return $ ExecFoldSeq f a0 ix (Just a)+ ExecFoldSeqFlatten afun acc ix _ ->+ do a <- executeOpenAcc acc aenv stream+ return $ ExecFoldSeqFlatten afun acc ix (Just a)+ ExecStuple t -> ExecStuple <$> initCT t++ initCT :: forall t. Atuple (ExecC aenv lenv) t -> CIO (Atuple (ExecC aenv lenv) t)+ initCT NilAtup = return NilAtup+ initCT (SnocAtup t c) = SnocAtup <$> initCT t <*> initC c++ -- get the extent of an embedded array+ extent :: Shape sh => ExecOpenAcc aenv (Array sh e) -> CIO sh+ extent ExecAcc{} = $internalError "executeOpenAcc" "expected delayed array"+ extent ExecSeq{} = $internalError "executeOpenAcc" "expected delayed array"+ extent (EmbedAcc sh) = executeExp sh aenv stream+++-- Turn a sequence computation into a suspended computation that can be forced+-- an element at a time.+--+streamSeq :: ExecSeq [a] -> StreamSeq a+streamSeq (ExecS binds sequ) = StreamSeq $ do+ aenv <- executeExtend binds Aempty+ iseq <- initializeSeq aenv sequ+ let go s = do+ ms <- stepSeq aenv s+ case ms of+ Nothing -> return Nothing+ Just (s', a) -> return (Just (a, StreamSeq (go s')))+ go iseq+++stepSeq+ :: forall a aenv.+ Aval aenv+ -> ExecOpenSeq aenv () [a]+ -> CIO (Maybe (ExecOpenSeq aenv () [a], a))+stepSeq aenv s = streaming step wait+ where+ step :: Stream -> CIO (Maybe (ExecOpenSeq aenv () [a], a))+ step stream = do+ ms <- runMaybeT (stepOpenSeq aenv s stream)+ return $ (,) <$> ms <*> (coll <$> ms)++ coll :: ExecOpenSeq aenv lenv [a] -> a+ coll (ExecP _ s') = coll s'+ coll (ExecC _) = $internalError "stepSeq" "Unreachable"+ coll (ExecR _ ma) = case ma of+ Nothing -> $internalError "stepSeq" "Trying to collect the value of an unexecuted sequence"+ Just a -> a++stepOpenSeq+ :: forall aenv arrs'.+ Aval aenv+ -> ExecOpenSeq aenv () arrs'+ -> Stream+ -> MaybeT CIO (ExecOpenSeq aenv () arrs')+stepOpenSeq aenv !l stream = go l Empty+ where+ go :: forall lenv. ExecOpenSeq aenv lenv arrs' -> Val lenv -> MaybeT CIO (ExecOpenSeq aenv lenv arrs')+ go s lenv =+ case s of+ ExecP p s' -> do+ (p', a) <- produce p+ s'' <- go s' (lenv `Push` a)+ return $ ExecP p' s''+ ExecC c -> ExecC <$> lift (consume c)+ ExecR ix _ -> return $ ExecR ix (Just (prj ix lenv))+ where+ produce :: forall a . ExecP aenv lenv a -> MaybeT CIO (ExecP aenv lenv a, a)+ produce (ExecToSeq slix acc kernel gamma sls) =+ do+ sl : sls' <- return sls+ lift $ do+ sh <- extent acc+ out <- allocateArray (sliceShape slix sh)+ m <- marshalSlice slix sl+ execute kernel gamma aenv (size sh) (m, out) stream+ return (ExecToSeq slix acc kernel gamma sls', out)+ produce (ExecUseLazy slix arr sls) =+ do let sh = shape arr+ lift $ do+ sl : sls' <- return sls+ out <- allocateArray (sliceShape slix sh)+ useArraySlice slix sl arr out+ return (ExecUseLazy slix arr sls', out)+ produce (ExecStreamIn xs) =+ let use :: ArraysR arrs -> arrs -> CIO ()+ use ArraysRunit () = return ()+ use ArraysRarray arr = useArrayAsync arr Nothing+ use (ArraysRpair r1 r2) (a1, a2) = use r1 a1 >> use r2 a2+ in case xs of+ [] -> MaybeT (return Nothing)+ x:xs' -> lift (use (arrays x) (fromArr x)) >> return (ExecStreamIn xs', x)+ produce (ExecMap afun x) = (ExecMap afun x ,) <$> lift (travAfun1 afun (prj x lenv))+ produce (ExecZipWith afun x y) = (ExecZipWith afun x y,) <$> lift (travAfun2 afun (prj x lenv) (prj y lenv))+ produce (ExecScanSeq afun a0 x ma) =+ do+ a <- MaybeT (return ma)+ a' <- lift $ travFun2 afun a (prj x lenv ! Z)+ return (ExecScanSeq afun a0 x (Just a'), fromList Z [a'])++ consume :: forall a . ExecC aenv lenv a -> CIO (ExecC aenv lenv a)+ consume c =+ case c of+ ExecFoldSeq afun a0 x (Just a) ->+ do b <- indexArray (prj x lenv) 0+ a' <- travFun2 afun a b+ return $ ExecFoldSeq afun a0 x (Just a')+ ExecFoldSeqFlatten afun a0 x (Just a) ->+ do useArray shapes+ a' <- travAfun3 afun a shapes elems+ return $ ExecFoldSeqFlatten afun a0 x (Just a')+ where+ Array sh adata = prj x lenv+ elems = Array ((), R.size sh) adata+ shapes = fromList (Z:.1) [toElt sh]+ ExecStuple t -> ExecStuple <$> consumeT t+ _ -> $internalError "executeSequence" "Expected already executed consumer"++ consumeT :: forall t. Atuple (ExecC aenv lenv) t -> CIO (Atuple (ExecC aenv lenv) t)+ consumeT NilAtup = return NilAtup+ consumeT (SnocAtup t c) = SnocAtup <$> consumeT t <*> consume c++ travAfun1 :: forall a b. PreOpenAfun ExecOpenAcc aenv (a -> b) -> a -> CIO b+ travAfun1 (Alam (Abody afun)) a =+ do nop <- join $ liftIO <$> (Event.create <$> asks activeContext <*> gets eventTable)+ executeOpenAcc afun (aenv `Apush` (Async nop a)) stream+ travAfun1 _ _ = error "travAfun1"++ travAfun2 :: forall a b c. PreOpenAfun ExecOpenAcc aenv (a -> b -> c) -> a -> b -> CIO c+ travAfun2 (Alam (Alam (Abody afun))) a b =+ do nop <- join $ liftIO <$> (Event.create <$> asks activeContext <*> gets eventTable)+ executeOpenAcc afun (aenv `Apush` (Async nop a) `Apush` (Async nop b)) stream+ travAfun2 _ _ _ = error "travAfun2"++ travAfun3 :: forall a b c d. PreOpenAfun ExecOpenAcc aenv (a -> b -> c -> d) -> a -> b -> c -> CIO d+ travAfun3 (Alam (Alam (Alam (Abody afun)))) a b c =+ do nop <- join $ liftIO <$> (Event.create <$> asks activeContext <*> gets eventTable)+ executeOpenAcc afun (aenv `Apush` (Async nop a) `Apush` (Async nop b) `Apush` (Async nop c)) stream+ travAfun3 _ _ _ _ = error "travAfun3"++ travFun2 :: forall a b c. PreFun ExecOpenAcc aenv (a -> b -> c) -> a -> b -> CIO c+ travFun2 (Lam (Lam (Body c))) a b = executeOpenExp c (Empty `Push` a `Push` b) aenv stream+ travFun2 _ _ _ = error "travFun2"++ -- get the extent of an embedded array+ extent :: Shape sh => ExecOpenAcc aenv (Array sh e) -> CIO sh+ extent ExecAcc{} = $internalError "executeOpenAcc" "expected delayed array"+ extent ExecSeq{} = $internalError "executeOpenAcc" "expected delayed array"+ extent (EmbedAcc sh) = executeExp sh aenv stream+++-- Evaluating bindings+-- -------------------++executeExtend :: Extend ExecOpenAcc aenv aenv' -> Aval aenv -> CIO (Aval aenv')+executeExtend BaseEnv aenv = return aenv+executeExtend (PushEnv e a) aenv = do+ aenv' <- executeExtend e aenv+ streaming (executeOpenAcc a aenv') $ \a' -> return $ Apush aenv' a'+--}++ -- Scalar expression evaluation -- ---------------------------- executeExp :: ExecExp aenv t -> Aval aenv -> Stream -> CIO t executeExp !exp !aenv !stream = executeOpenExp exp Empty aenv stream -executeOpenExp :: forall env aenv exp. ExecOpenExp env aenv exp -> Val env -> Aval aenv -> Stream -> CIO exp+executeOpenExp+ :: forall env aenv exp.+ ExecOpenExp env aenv exp+ -> Val env+ -> Aval aenv+ -> Stream+ -> CIO exp executeOpenExp !rootExp !env !aenv !stream = travE rootExp where travE :: ExecOpenExp env aenv t -> CIO t@@ -455,6 +736,7 @@ ToIndex sh ix -> toIndex <$> travE sh <*> travE ix FromIndex sh ix -> fromIndex <$> travE sh <*> travE ix Intersect sh1 sh2 -> intersect <$> travE sh1 <*> travE sh2+ Union sh1 sh2 -> union <$> travE sh1 <*> travE sh2 ShapeSize sh -> size <$> travE sh Shape acc -> shape <$> travA acc Index acc ix -> join $ index <$> travA acc <*> travE ix@@ -517,45 +799,64 @@ -- Marshalling data -- ---------------- +{--+marshalSlice'+ :: SliceIndex slix sl co dim+ -> slix+ -> CIO [CUDA.FunParam]+marshalSlice' SliceNil () = return []+marshalSlice' (SliceAll sl) (sh, ()) = marshalSlice' sl sh+marshalSlice' (SliceFixed sl) (sh, n) =+ do x <- runContT (marshal n Nothing) return+ xs <- marshalSlice' sl sh+ return (xs ++ x)++marshalSlice+ :: Elt slix => SliceIndex (EltRepr slix) sl co dim+ -> slix+ -> CIO [CUDA.FunParam]+marshalSlice slix = marshalSlice' slix . fromElt+--}+ -- Data which can be marshalled as function arguments to a kernel invocation. -- class Marshalable a where- marshal :: a -> CIO [CUDA.FunParam]+ marshal :: a -> Maybe Stream -> ContT b CIO [CUDA.FunParam] instance Marshalable () where- marshal () = return []+ marshal () _ = return [] instance Marshalable CUDA.FunParam where- marshal !x = return [x]+ marshal !x _ = return [x] instance ArrayElt e => Marshalable (ArrayData e) where- marshal !ad = marshalArrayData ad+ marshal !ad ms = ContT $ marshalArrayData ad ms instance Shape sh => Marshalable sh where- marshal !sh = marshal (reverse (shapeToList sh))+ marshal !sh ms = marshal (reverse (shapeToList sh)) ms instance Marshalable a => Marshalable [a] where- marshal = concatMapM marshal+ marshal xs ms = concatMapM (flip marshal ms) xs instance (Marshalable sh, Elt e) => Marshalable (Array sh e) where- marshal !(Array sh ad) = (++) <$> marshal (toElt sh :: sh) <*> marshal ad+ marshal !(Array sh ad) ms = (++) <$> marshal (toElt sh :: sh) ms <*> marshal ad ms instance (Marshalable a, Marshalable b) => Marshalable (a, b) where- marshal (!a, !b) = (++) <$> marshal a <*> marshal b+ marshal (!a, !b) ms = (++) <$> marshal a ms <*> marshal b ms instance (Marshalable a, Marshalable b, Marshalable c) => Marshalable (a, b, c) where- marshal (!a, !b, !c)- = concat <$> sequence [marshal a, marshal b, marshal c]+ marshal (!a, !b, !c) ms+ = concat <$> sequence [marshal a ms, marshal b ms, marshal c ms] instance (Marshalable a, Marshalable b, Marshalable c, Marshalable d) => Marshalable (a, b, c, d) where- marshal (!a, !b, !c, !d)- = concat <$> sequence [marshal a, marshal b, marshal c, marshal d]+ marshal (!a, !b, !c, !d) ms+ = concat <$> sequence [marshal a ms, marshal b ms, marshal c ms, marshal d ms] #define primMarshalable(ty) \ instance Marshalable (ty) where { \- marshal !x = return [CUDA.VArg x] }+ marshal !x _ = return [CUDA.VArg x] } primMarshalable(Int) primMarshalable(Int8)@@ -589,21 +890,23 @@ -- as to use the larger available caches. -- -marshalAccEnvTex :: AccKernel a -> Aval aenv -> Gamma aenv -> Stream -> CIO [CUDA.FunParam]+marshalAccEnvTex :: AccKernel a -> Aval aenv -> Gamma aenv -> Stream -> ContT b CIO [CUDA.FunParam] marshalAccEnvTex !kernel !aenv (Gamma !gamma) !stream = flip concatMapM (Map.toList gamma) $ \(Idx_ !(idx :: Idx aenv (Array sh e)), i) -> do arr <- after stream (aprj idx aenv)- marshalAccTex (namesOfArray (groupOfInt i) (undefined :: e)) kernel arr- marshal (shape arr)+ ContT $ \f -> marshalAccTex (namesOfArray (groupOfInt i) (undefined :: e)) kernel arr (Just stream) (f ())+ marshal (shape arr) (Just stream) -marshalAccTex :: (Name,[Name]) -> AccKernel a -> Array sh e -> CIO ()-marshalAccTex (_, !arrIn) (AccKernel _ _ !mdl _ _ _ _) (Array !sh !adata)- = marshalTextureData adata (R.size sh) =<< liftIO (sequence' $ map (CUDA.getTex mdl) (reverse arrIn))+marshalAccTex :: (Name,[Name]) -> AccKernel a -> Array sh e -> Maybe Stream -> CIO b -> CIO b+marshalAccTex (_, !arrIn) (AccKernel _ _ !lmdl _ _ _ _) (Array !sh !adata) ms run+ = do+ texs <- liftIO $ withLifetime lmdl $ \mdl -> (sequence' $ map (CUDA.getTex mdl) (reverse arrIn))+ marshalTextureData adata (R.size sh) texs ms (const run) -marshalAccEnvArg :: Aval aenv -> Gamma aenv -> Stream -> CIO [CUDA.FunParam]+marshalAccEnvArg :: Aval aenv -> Gamma aenv -> Stream -> ContT b CIO [CUDA.FunParam] marshalAccEnvArg !aenv (Gamma !gamma) !stream- = concatMapM (\(Idx_ !idx) -> marshal =<< after stream (aprj idx aenv)) (Map.keys gamma)+ = concatMapM (\(Idx_ !idx) -> flip marshal (Just stream) =<< after stream (aprj idx aenv)) (Map.keys gamma) -- A lazier version of 'Control.Monad.sequence'@@ -632,45 +935,47 @@ -- texture references, and for newer devices adds the parameters to the front of -- the argument list ---arguments :: Marshalable args- => AccKernel a- -> Aval aenv- -> Gamma aenv- -> args- -> Stream- -> CIO [CUDA.FunParam]+arguments+ :: Marshalable args+ => AccKernel a+ -> Aval aenv+ -> Gamma aenv+ -> args+ -> Stream+ -> ContT b CIO [CUDA.FunParam] arguments !kernel !aenv !gamma !a !stream = do dev <- asks deviceProperties let marshaller | computeCapability dev < Compute 2 0 = marshalAccEnvTex kernel | otherwise = marshalAccEnvArg --- (++) <$> marshaller aenv gamma stream <*> marshal a+ (++) <$> marshaller aenv gamma stream <*> marshal a (Just stream) -- Link the binary object implementing the computation, configure the kernel -- launch parameters, and initiate the computation. This also handles lifting -- and binding of array references from scalar expressions. ---execute :: Marshalable args- => AccKernel a -- The binary module implementing this kernel- -> Gamma aenv -- variables of arrays embedded in scalar expressions- -> Aval aenv -- the environment- -> Int -- a "size" parameter, typically number of elements in the output- -> args -- arguments to marshal to the kernel function- -> Stream -- Compute stream to execute in- -> CIO ()-execute !kernel !gamma !aenv !n !a !stream = do- args <- arguments kernel aenv gamma a stream- launch kernel (configure kernel n) args stream+execute+ :: Marshalable args+ => AccKernel a -- The binary module implementing this kernel+ -> Gamma aenv -- variables of arrays embedded in scalar expressions+ -> Aval aenv -- the environment+ -> Int -- a "size" parameter, typically number of elements in the output+ -> args -- arguments to marshal to the kernel function+ -> Stream -- Compute stream to execute in+ -> CIO ()+execute !kernel !gamma !aenv !n !a !stream = flip runContT return $ do+ args <- arguments kernel aenv gamma a stream+ liftIO $ launch kernel (configure kernel n) args stream -- Execute a device function, with the given thread configuration and function -- parameters. The tuple contains (threads per block, grid size, shared memory) ---launch :: AccKernel a -> (Int,Int,Int) -> [CUDA.FunParam] -> Stream -> CIO ()+launch :: AccKernel a -> (Int,Int,Int) -> [CUDA.FunParam] -> Stream -> IO () launch (AccKernel entry !fn _ _ _ _ _) !(cta, grid, smem) !args !stream = D.timed D.dump_exec msg (Just stream)- $ liftIO $ CUDA.launchKernel fn (grid,1,1) (cta,1,1) smem (Just stream) args+ $ CUDA.launchKernel fn (grid,1,1) (cta,1,1) smem (Just stream) args where msg gpuTime cpuTime = "exec: " ++ entry ++ "<<< " ++ shows grid ", " ++ shows cta ", " ++ shows smem " >>> "
Data/Array/Accelerate/CUDA/Execute/Event.hs view
@@ -1,5 +1,8 @@ {-# LANGUAGE BangPatterns #-} {-# LANGUAGE CPP #-}+{-# LANGUAGE FlexibleInstances #-}+{-# LANGUAGE ViewPatterns #-}+{-# OPTIONS_GHC -fno-warn-orphans #-} -- | -- Module : Data.Array.Accelerate.CUDA.Execute.Event -- Copyright : [2013..2014] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell@@ -11,39 +14,94 @@ -- module Data.Array.Accelerate.CUDA.Execute.Event ( - Event, create, waypoint, after, block, Event.destroy,+ Event, EventTable, new, create, waypoint, after, block, query, destroy, ) where -- friends-#ifdef ACCELERATE_DEBUG+import Data.Array.Accelerate.FullList ( FullList(..) )+import Data.Array.Accelerate.Lifetime+import Data.Array.Accelerate.CUDA.Context ( Context(..) ) import qualified Data.Array.Accelerate.CUDA.Debug as D-#endif+import qualified Data.Array.Accelerate.FullList as FL -- libraries-import Foreign.CUDA.Driver.Event ( Event(..) )+import Control.Concurrent.MVar ( MVar, newMVar, withMVar, mkWeakMVar )+import Control.Exception ( bracket_ )+import Data.Hashable ( Hashable(..) ) import Foreign.CUDA.Driver.Stream ( Stream(..) )+import Foreign.Ptr ( ptrToIntPtr )+import qualified Foreign.CUDA.Driver as CUDA import qualified Foreign.CUDA.Driver.Event as Event+import qualified Data.HashTable.IO as HT +type Event = Lifetime Event.Event --- Create a new event. It will not be automatically garbage collected, and is--- not suitable for timing purposes.+type HashTable key val = HT.BasicHashTable key val++type EventTable = MVar ( HashTable (Lifetime CUDA.Context) (FullList () Event.Event) )++instance Hashable (Lifetime CUDA.Context) where+ {-# INLINE hashWithSalt #-}+ hashWithSalt salt (unsafeGetValue -> CUDA.Context ctx)+ = salt `hashWithSalt` (fromIntegral (ptrToIntPtr ctx) :: Int)+++-- Generate a new empty event table. --+new :: IO EventTable+new = do+ tbl <- HT.new+ ref <- newMVar tbl+ _ <- mkWeakMVar ref (flush tbl)+ return ref++-- Create a new event. It will be automatically garbage collected, if a recycled+-- event is available, it will be returned, else a new event is created.+-- {-# INLINE create #-}-create :: IO Event-create = do- event <- Event.create [Event.DisableTiming]- message ("create " ++ showEvent event)+create :: Context -> EventTable -> IO Event+create ctx ref = withMVar ref $ \tbl -> do+ --+ let key = deviceContext ctx+ me <- HT.lookup tbl key+ e <- case me of+ Nothing -> do+ e <- Event.create [Event.DisableTiming]+ message ("new " ++ show e)+ return e++ Just (FL () e rest) -> do+ case rest of+ FL.Nil -> HT.delete tbl key+ FL.Cons () e' es -> HT.insert tbl key (FL () e' es)+ --+ return e+ --+ event <- newLifetime e+ addFinalizer event $ do+ D.traceIO D.dump_gc ("gc: finalise event " ++ showEvent event)+ insert ref (deviceContext ctx) e return event +{-# INLINE insert #-}+-- Insert an event into the table.+--+insert :: EventTable -> Lifetime CUDA.Context -> Event.Event -> IO ()+insert ref lctx e = withMVar ref $ \tbl -> do+ me <- HT.lookup tbl lctx+ case me of+ Nothing -> HT.insert tbl lctx (FL.singleton () e)+ Just es -> HT.insert tbl lctx (FL.cons () e es)+ -- Create a new event marker that will be filled once execution in the specified -- stream has completed all previously submitted work. -- {-# INLINE waypoint #-}-waypoint :: Stream -> IO Event-waypoint stream = do- event <- create- Event.record event (Just stream)+waypoint :: Context -> EventTable -> Stream -> IO Event+waypoint ctx ref stream = do+ event <- create ctx ref+ withLifetime event (`Event.record` Just stream) message $ "waypoint " ++ showEvent event ++ " in " ++ showStream stream return event @@ -54,40 +112,47 @@ after :: Event -> Stream -> IO () after event stream = do message $ "after " ++ showEvent event ++ " in " ++ showStream stream- Event.wait event (Just stream) []+ withLifetime event $ \e -> Event.wait e (Just stream) [] -- Block the calling thread until the event is recorded -- {-# INLINE block #-} block :: Event -> IO ()-block = Event.block+block = flip withLifetime Event.block +-- Query the status of the event.+--+{-# INLINE query #-}+query :: Event -> IO Bool+query = flip withLifetime Event.query --- Add a finaliser to an event token+-- Explicitly destroy the event. ----- addEventFinalizer :: Event -> IO () -> IO ()--- addEventFinalizer e@(Event (Ptr e#)) f = IO $ \s ->--- case mkWeak# e# e f s of (# s', _w #) -> (# s', () #)+{-# INLINE destroy #-}+destroy :: Event -> IO ()+destroy = finalize +-- Destroy all events in the table.+--+flush :: HashTable (Lifetime CUDA.Context) (FullList () Event.Event) -> IO ()+flush !tbl =+ let clean (!lctx,!es) = do+ withLifetime lctx $ \ctx -> bracket_ (CUDA.push ctx) CUDA.pop (FL.mapM_ (const Event.destroy) es)+ HT.delete tbl lctx+ in+ message "flush reservoir" >> HT.mapM_ clean tbl + -- Debug -- ----- -{-# INLINE trace #-}-trace :: String -> IO a -> IO a-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 ()+message msg = D.traceIO D.dump_sched ("event: " ++ msg) {-# INLINE showEvent #-} showEvent :: Event -> String-showEvent (Event e) = show e+showEvent (unsafeGetValue -> Event.Event e) = show e {-# INLINE showStream #-} showStream :: Stream -> String
Data/Array/Accelerate/CUDA/Execute/Stream.hs view
@@ -1,5 +1,8 @@-{-# LANGUAGE BangPatterns #-}-{-# LANGUAGE CPP #-}+{-# LANGUAGE BangPatterns #-}+{-# LANGUAGE CPP #-}+{-# LANGUAGE FlexibleInstances #-}+{-# LANGUAGE ViewPatterns #-}+{-# OPTIONS_GHC -fno-warn-orphans #-} -- | -- Module : Data.Array.Accelerate.CUDA.Execute.Stream -- Copyright : [2013..2014] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell@@ -17,16 +20,13 @@ ) where -- friends-import Data.Array.Accelerate.CUDA.Array.Nursery ( ) -- hashable CUDA.Context instance import Data.Array.Accelerate.CUDA.Context ( Context(..) )-import Data.Array.Accelerate.CUDA.FullList ( FullList(..) )-import Data.Array.Accelerate.CUDA.Execute.Event ( Event )+import Data.Array.Accelerate.CUDA.Execute.Event ( Event, EventTable )+import Data.Array.Accelerate.FullList ( FullList(..) )+import Data.Array.Accelerate.Lifetime ( Lifetime, withLifetime ) 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+import qualified Data.Array.Accelerate.FullList as FL -- libraries import Control.Monad.Trans ( MonadIO, liftIO )@@ -49,7 +49,7 @@ -- type HashTable key val = HT.BasicHashTable key val -type RSV = MVar ( HashTable CUDA.Context (FullList () Stream) )+type RSV = MVar ( HashTable (Lifetime CUDA.Context) (FullList () Stream) ) data Reservoir = Reservoir {-# UNPACK #-} !RSV {-# UNPACK #-} !(Weak RSV) @@ -63,11 +63,11 @@ -- second operation completes. -- {-# INLINE streaming #-}-streaming :: MonadIO m => Context -> Reservoir -> (Stream -> m a) -> (Event -> a -> m b) -> m b-streaming !ctx !rsv@(Reservoir !_ !weak_rsv) !action !after = do+streaming :: MonadIO m => Context -> Reservoir -> EventTable -> (Stream -> m a) -> (Event -> a -> m b) -> m b+streaming !ctx !rsv@(Reservoir !_ !weak_rsv) !etbl !action !after = do stream <- liftIO $ create ctx rsv first <- action stream- end <- liftIO $ Event.waypoint stream+ end <- liftIO $ Event.waypoint ctx etbl stream final <- after end first liftIO $! destroy (weakContext ctx) weak_rsv stream liftIO $! Event.destroy end@@ -94,8 +94,8 @@ {-# INLINE create #-} create :: Context -> Reservoir -> IO Stream create !ctx (Reservoir !ref !_) = withMVar ref $ \tbl -> do- let key = deviceContext ctx --+ let key = deviceContext ctx ms <- HT.lookup tbl key case ms of Nothing -> do@@ -115,7 +115,7 @@ -- pending operations in the stream have completed. -- {-# INLINE destroy #-}-destroy :: Weak CUDA.Context -> Weak RSV -> Stream -> IO ()+destroy :: Weak (Lifetime 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'.@@ -149,11 +149,11 @@ -- Destroy all streams in the reservoir. ---flush :: HashTable CUDA.Context (FullList () Stream) -> IO ()+flush :: HashTable (Lifetime CUDA.Context) (FullList () Stream) -> IO () flush !tbl =- let clean (!ctx,!ss) = do- bracket_ (CUDA.push ctx) CUDA.pop (FL.mapM_ (const Stream.destroy) ss)- HT.delete tbl ctx+ let clean (!lctx,!ss) = do+ withLifetime lctx $ \ctx -> bracket_ (CUDA.push ctx) CUDA.pop (FL.mapM_ (const Stream.destroy) ss)+ HT.delete tbl lctx in message "flush reservoir" >> HT.mapM_ clean tbl @@ -163,15 +163,11 @@ {-# INLINE trace #-} trace :: String -> IO a -> IO a-trace _msg next = do-#ifdef ACCELERATE_DEBUG- D.when D.verbose $ D.message D.dump_exec ("stream: " ++ _msg)-#endif- next+trace msg next = message msg >> next {-# INLINE message #-} message :: String -> IO ()-message s = s `trace` return ()+message msg = D.traceIO D.dump_sched ("stream: " ++ msg) {-# INLINE showStream #-} showStream :: Stream -> String
Data/Array/Accelerate/CUDA/Foreign/Export.hs view
@@ -1,14 +1,15 @@-{-# LANGUAGE RankNTypes #-}-{-# LANGUAGE ScopedTypeVariables #-} {-# LANGUAGE BangPatterns #-}-{-# LANGUAGE GADTs #-} {-# LANGUAGE CPP #-}+{-# LANGUAGE FlexibleInstances #-} {-# LANGUAGE ForeignFunctionInterface #-}-{-# LANGUAGE TemplateHaskell #-}+{-# LANGUAGE GADTs #-}+{-# LANGUAGE ImpredicativeTypes #-} {-# LANGUAGE QuasiQuotes #-}+{-# LANGUAGE RankNTypes #-}+{-# LANGUAGE ScopedTypeVariables #-}+{-# LANGUAGE TemplateHaskell #-} {-# LANGUAGE TypeFamilies #-}-{-# LANGUAGE FlexibleInstances #-}-{-# LANGUAGE ImpredicativeTypes #-}+{-# LANGUAGE ViewPatterns #-} {-# OPTIONS_GHC -fno-warn-name-shadowing #-} {-# OPTIONS_GHC -fno-warn-orphans #-} -- |@@ -41,23 +42,23 @@ import Data.Functor import Control.Applicative+import Control.Monad.State ( liftIO ) import Foreign.StablePtr import Foreign.C.Types import Foreign.Ptr import Foreign.Storable ( Storable(..) ) import Foreign.Marshal.Array ( peekArray, pokeArray, mallocArray ) import Foreign.Marshal.Alloc ( free )-import Control.Monad.State ( liftIO ) import Language.Haskell.TH hiding ( ppr )-import Prelude as P- import qualified Foreign.CUDA.Driver as CUDA +import Prelude as P+ -- friends import Data.Array.Accelerate.Smart ( Acc ) import Data.Array.Accelerate.Type import Data.Array.Accelerate.Array.Data-import Data.Array.Accelerate.CUDA ( run1In )+import Data.Array.Accelerate.CUDA ( run1With ) import Data.Array.Accelerate.CUDA.Array.Sugar hiding ( shape, size ) import Data.Array.Accelerate.CUDA.Array.Data hiding ( pokeArray, peekArray, mallocArray ) import Data.Array.Accelerate.CUDA.State@@ -212,12 +213,13 @@ shbuf <- liftIO $ mallocArray (P.length sh') liftIO $ pokeArray shbuf (map fromIntegral sh') - dptrs <- devicePtrsToWordPtrs adata <$> devicePtrsOfArrayData adata- pbuf <- liftIO $ mallocArray (P.length dptrs)- liftIO $ pokeArray pbuf dptrs+ withDevicePtrs adata Nothing $ \dptrs -> do+ let wptrs = devicePtrsToWordPtrs adata dptrs+ pbuf <- liftIO $ mallocArray (P.length wptrs)+ liftIO $ pokeArray pbuf wptrs - sa <- liftIO $ newStablePtr (EArray a)- return (shbuf, pbuf, sa)+ sa <- liftIO $ newStablePtr (EArray a)+ return (shbuf, pbuf, sa) marshalOut (ArraysRpair aR1 aR2) (x,y) ptr = do ptr' <- marshalOut aR1 x ptr@@ -228,7 +230,11 @@ -- a function callable from foreign code with the second argument specifying it's name. exportAfun :: Name -> String -> Q [Dec] exportAfun fname ename = do+#if __GLASGOW_HASKELL__ <= 710 (VarI n ty _ _) <- reify fname+#else+ (VarI n ty _) <- reify fname+#endif -- Generate initialisation function genCompileFun n ename ty@@ -256,14 +262,14 @@ ef :: IO (StablePtr Afun) ef = do ctx <- deRefStablePtr hndl- newStablePtr (Afun (run1In ctx f) (undefined :: a) (undefined :: b))+ newStablePtr (Afun (run1With ctx f) (undefined :: a) (undefined :: b)) -- Utility functions -- ------------------ arrayFromForeignData :: forall sh e. (Shape sh, Elt e) => DevicePtrBuffer -> ShapeBuffer -> CIO (Array sh e) arrayFromForeignData ptrs shape = do- let d = dim (ignore :: sh) -- Using ignore as using dim requires a non-dummy argument+ let d = rank (ignore :: sh) -- Using ignore as using dim requires a non-dummy argument let sz = eltSize (eltType (undefined :: e)) lst <- liftIO (peekArray d shape) let sh = listToShape (map fromIntegral lst) :: sh
Data/Array/Accelerate/CUDA/Foreign/Import.hs view
@@ -44,13 +44,13 @@ -- * Manipulating arrays DevicePtrs,- devicePtrsOfArray,+ withDevicePtrs, indexArray, useArray, useArrayAsync, peekArray, peekArrayAsync, pokeArray, pokeArrayAsync, copyArray, copyArrayAsync,- allocateArray, newArray,+ allocateArray, fromFunction, -- * Running IO actions in an Accelerate context CIO, Stream, liftIO, inContext, inDefaultContext@@ -61,13 +61,15 @@ import Data.Array.Accelerate.CUDA.State import Data.Array.Accelerate.CUDA.Context import Data.Array.Accelerate.CUDA.Array.Sugar-import Data.Array.Accelerate.CUDA.Array.Data+import Data.Array.Accelerate.CUDA.Array.Data hiding ( withDevicePtrs )+import qualified Data.Array.Accelerate.CUDA.Array.Data as D import Data.Array.Accelerate.CUDA.Array.Prim ( DevicePtrs ) import Data.Array.Accelerate.CUDA.Execute.Stream ( Stream ) import Data.Typeable import Control.Exception ( bracket_ ) import Control.Monad.Trans ( liftIO )+import qualified Foreign.CUDA.Driver.Stream as CUDA -- CUDA backend representation of foreign functions@@ -75,10 +77,10 @@ -- |CUDA foreign Acc functions are just CIO functions. ---data CUDAForeignAcc as bs where+data CUDAForeignAcc f where CUDAForeignAcc :: String -- name of the function -> (Stream -> as -> CIO bs) -- operation to execute- -> CUDAForeignAcc as bs+ -> CUDAForeignAcc (as -> bs) deriving instance Typeable CUDAForeignAcc @@ -89,11 +91,11 @@ -- CUDA backend. -- canExecuteAcc- :: (Foreign f, Typeable as, Typeable bs)- => f as bs+ :: forall asm as bs. (Foreign asm, Typeable as, Typeable bs)+ => asm (as -> bs) -> Maybe (Stream -> as -> CIO bs) canExecuteAcc ff- | Just (CUDAForeignAcc _ fun) <- cast ff+ | Just (CUDAForeignAcc _ fun) <- cast ff :: Maybe (CUDAForeignAcc (as -> bs)) = Just fun | otherwise@@ -102,11 +104,11 @@ -- |CUDA foreign Exp functions consist of a list of C header files necessary to call the function -- and the name of the function to call. ---data CUDAForeignExp x y where+data CUDAForeignExp f where CUDAForeignExp :: IsScalar y => [String] -- header files to be imported -> String -- name of the foreign function- -> CUDAForeignExp x y+ -> CUDAForeignExp (x -> y) deriving instance Typeable CUDAForeignExp @@ -117,11 +119,11 @@ -- for the CUDA backend. -- canExecuteExp- :: forall f x y. (Foreign f, Typeable y, Typeable x)- => f x y+ :: forall asm x y. (Foreign asm, Typeable y, Typeable x)+ => asm (x -> y) -> Maybe ([String], String) canExecuteExp ff- | Just (CUDAForeignExp hdr fun) <- cast ff :: Maybe (CUDAForeignExp x y)+ | Just (CUDAForeignExp hdr fun) <- cast ff :: Maybe (CUDAForeignExp (x -> y)) = Just (hdr, fun) | otherwise@@ -131,10 +133,11 @@ -- User facing utility functions -- ----------------------------- --- |Get the raw CUDA device pointers associated with an array+-- |Get the raw CUDA device pointers associated with an array and call the given+-- continuation. ---devicePtrsOfArray :: Array sh e -> CIO (DevicePtrs (EltRepr e))-devicePtrsOfArray (Array _ adata) = devicePtrsOfArrayData adata+withDevicePtrs :: Array sh e -> Maybe CUDA.Stream -> (DevicePtrs (EltRepr e) -> CIO b) -> CIO b+withDevicePtrs (Array _ adata) = D.withDevicePtrs adata -- |Run an IO action within the given Acclerate context --
− Data/Array/Accelerate/CUDA/FullList.hs
@@ -1,118 +0,0 @@-{-# LANGUAGE BangPatterns #-}-{-# LANGUAGE PatternGuards #-}--- |--- Module : Data.Array.Accelerate.CUDA.FullList--- 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>--- Stability : experimental--- Portability : non-portable (GHC extensions)------ Non-empty lists of key/value pairs. The lists are strict in the key and lazy--- in the values. We assume that keys only occur once.-----module Data.Array.Accelerate.CUDA.FullList (-- FullList(..),- List(..),-- singleton,- cons,- size,- mapM_,- lookup,- lookupDelete,--) where--import Prelude hiding ( lookup, mapM_ )---data FullList k v = FL !k v !(List k v)-data List k v = Nil | Cons !k v !(List k v)--infixr 5 `Cons`--instance (Eq k, Eq v) => Eq (FullList k v) where- (FL k1 v1 xs) == (FL k2 v2 ys) = k1 == k2 && v1 == v2 && xs == ys- (FL k1 v1 xs) /= (FL k2 v2 ys) = k1 /= k2 || v1 /= v2 || xs /= ys--instance (Eq k, Eq v) => Eq (List k v) where- (Cons k1 v1 xs) == (Cons k2 v2 ys) = k1 == k2 && v1 == v2 && xs == ys- Nil == Nil = True- _ == _ = False-- (Cons k1 v1 xs) /= (Cons k2 v2 ys) = k1 /= k2 || v1 /= v2 || xs /= ys- Nil /= Nil = False- _ /= _ = True----- List-like operations----infixr 5 `cons`-cons :: k -> v -> FullList k v -> FullList k v-cons k v (FL k' v' xs) = FL k v (Cons k' v' xs)--singleton :: k -> v -> FullList k v-singleton k v = FL k v Nil--size :: FullList k v -> Int-size (FL _ _ xs) = 1 + sizeL xs--sizeL :: List k v -> Int-sizeL Nil = 0-sizeL (Cons _ _ xs) = 1 + sizeL xs--lookup :: Eq k => k -> FullList k v -> Maybe v-lookup key (FL k v xs)- | key == k = Just v- | otherwise = lookupL key xs-{-# INLINABLE lookup #-}-{-# SPECIALISE lookup :: () -> FullList () v -> Maybe v #-}--lookupL :: Eq k => k -> List k v -> Maybe v-lookupL !key = go- where- go Nil = Nothing- go (Cons k v xs)- | key == k = Just v- | otherwise = go xs-{-# INLINABLE lookupL #-}-{-# SPECIALISE lookupL :: () -> List () v -> Maybe v #-}--lookupDelete :: Eq k => k -> FullList k v -> (Maybe v, Maybe (FullList k v))-lookupDelete key (FL k v xs)- | key == k- = case xs of- Nil -> (Just v, Nothing)- Cons k' v' xs' -> (Just v, Just $ FL k' v' xs')-- | (r, xs') <- lookupDeleteL k xs- = (r, Just $ FL k v xs')-{-# INLINABLE lookupDelete #-}-{-# SPECIALISE lookupDelete :: () -> FullList () v -> (Maybe v, Maybe (FullList () v)) #-}--lookupDeleteL :: Eq k => k -> List k v -> (Maybe v, List k v)-lookupDeleteL !key = go- where- go Nil = (Nothing, Nil)- go (Cons k v xs)- | key == k = (Just v, xs)- | (r, xs') <- go xs = (r, Cons k v xs')-{-# INLINABLE lookupDeleteL #-}-{-# SPECIALISE lookupDeleteL :: () -> List () v -> (Maybe v, List () v) #-}--mapM_ :: Monad m => (k -> v -> m a) -> FullList k v -> m ()-mapM_ !f (FL k v xs) = f k v >> mapML_ f xs-{-# INLINABLE mapM_ #-}--mapML_ :: Monad m => (k -> v -> m a) -> List k v -> m ()-mapML_ !f = go- where- go Nil = return ()- go (Cons k v xs) = f k v >> go xs-{-# INLINABLE mapML_ #-}-
Data/Array/Accelerate/CUDA/Persistent.hs view
@@ -24,24 +24,25 @@ -- friends import Data.Array.Accelerate.Error+import Data.Array.Accelerate.FullList ( FullList )+import Data.Array.Accelerate.Lifetime ( Lifetime, withLifetime, newLifetime ) import Data.Array.Accelerate.CUDA.Context-import Data.Array.Accelerate.CUDA.FullList ( FullList ) import qualified Data.Array.Accelerate.CUDA.Debug as D-import qualified Data.Array.Accelerate.CUDA.FullList as FL+import qualified Data.Array.Accelerate.FullList as FL -- libraries import Numeric import Control.Applicative import Control.Concurrent import Control.Exception-import Control.Monad.Trans+import Control.Monad ( when ) import Data.Binary import Data.Binary.Get import Data.ByteString ( ByteString ) import Data.ByteString.Internal ( w2c ) import Data.Char import Data.Hashable-import Data.Maybe+import Data.Maybe ( fromMaybe ) import Data.Version import System.Directory import System.FilePath@@ -117,8 +118,9 @@ cubin <- (</>) <$> cacheDirectory <*> pure (cacheFilePath key) 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)+ lmdl <- newLifetime mdl+ let obj = KernelObject bin (FL.singleton (deviceContext context) lmdl)+ addFinalizer lmdl (module_finalizer (weakContext context) key lmdl) HT.insert kt key obj return $! Just obj @@ -137,13 +139,14 @@ -- Unload a kernel module from the specified context ---module_finalizer :: Weak CUDA.Context -> KernelKey -> CUDA.Module -> IO ()-module_finalizer weak_ctx key mdl = do+module_finalizer :: Weak (Lifetime CUDA.Context) -> KernelKey -> Lifetime CUDA.Module -> IO ()+module_finalizer weak_ctx key lmdl = do mc <- deRefWeak weak_ctx case mc of- Nothing -> D.message D.dump_gc ("gc: finalise module/dead context: " ++ cacheFilePath key)- Just ctx -> D.message D.dump_gc ("gc: finalise module: " ++ cacheFilePath key)- >> bracket_ (CUDA.push ctx) CUDA.pop (CUDA.unload mdl)+ Nothing -> D.traceIO D.dump_gc ("gc: finalise module/dead context: " ++ cacheFilePath key)+ Just fctx -> D.traceIO D.dump_gc ("gc: finalise module: " ++ cacheFilePath key)+ >> withLifetime fctx (\ctx -> withLifetime lmdl (\mdl ->+ bracket_ (CUDA.push ctx) CUDA.pop (CUDA.unload mdl))) -- Local cache -----------------------------------------------------------------@@ -183,7 +186,8 @@ -- re-link into the current context. -- | KernelObject {-# UNPACK #-} !ByteString- {-# UNPACK #-} !(FullList CUDA.Context CUDA.Module)+ {-# UNPACK #-} !(FullList (Lifetime CUDA.Context)+ (Lifetime CUDA.Module)) -- Persistent cache ------------------------------------------------------------@@ -300,7 +304,8 @@ -- restore :: FilePath -> IO PersistentCache restore !db = do- D.when D.flush_cache $ do+ mflush <- D.queryFlag D.flush_cache+ when (fromMaybe False mflush) $ do message "deleting persistent cache" cacheDir <- cacheDirectory removeDirectoryRecursive cacheDir@@ -377,10 +382,6 @@ -- ----- {-# INLINE message #-}-message :: MonadIO m => String -> m ()-message msg = trace msg $ return ()--{-# INLINE trace #-}-trace :: MonadIO m => String -> m a -> m a-trace msg next = D.message D.dump_cc ("cc: " ++ msg) >> next+message :: String -> IO ()+message msg = D.traceIO D.dump_cc ("cc: " ++ msg)
Data/Array/Accelerate/CUDA/State.hs view
@@ -20,20 +20,22 @@ module Data.Array.Accelerate.CUDA.State ( -- Evaluating computations- CIO, Context, evalCUDA,+ CIO, Context, evalCUDA, evalCUDAState, -- Querying execution state defaultContext, deviceProperties, activeContext, kernelTable, memoryTable, streamReservoir,+ eventTable, ) 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.Debug ( traceIO, dump_gc ) import Data.Array.Accelerate.CUDA.Persistent as KT ( KernelTable, new )-import Data.Array.Accelerate.CUDA.Array.Table as MT ( MemoryTable, new )+import Data.Array.Accelerate.CUDA.Array.Remote as MT ( MemoryTable, new ) import Data.Array.Accelerate.CUDA.Execute.Stream as ST ( Reservoir, new )+import Data.Array.Accelerate.CUDA.Execute.Event as ET ( EventTable, new) import Data.Array.Accelerate.CUDA.Analysis.Device -- library@@ -59,7 +61,8 @@ data State = State { memoryTable :: {-# UNPACK #-} !MemoryTable, -- host/device memory associations kernelTable :: {-# UNPACK #-} !KernelTable, -- compiled kernel object code- streamReservoir :: {-# UNPACK #-} !Reservoir -- kernel execution streams+ streamReservoir :: {-# UNPACK #-} !Reservoir, -- kernel execution streams+ eventTable :: {-# UNPACK #-} !EventTable -- CUDA events } newtype CIO a = CIO {@@ -88,6 +91,14 @@ teardown = pop action = evalStateT (runReaderT (runCIO acc) ctx) theState +-- |Evaluate a CUDA array computation with the specific state. Exceptions are+-- not caught.+--+-- RCE: This is unfortunately hacky, but necessary to stop device pointers+-- leaking.+evalCUDAState :: Context -> MemoryTable -> KernelTable -> Reservoir -> EventTable -> CIO a -> IO a+evalCUDAState ctx mt kt rsv etbl acc = evalStateT (runReaderT (runCIO acc) ctx)+ (State mt kt rsv etbl) -- Top-level mutable state -- -----------------------@@ -100,11 +111,12 @@ theState :: State theState = unsafePerformIO- $ do message dump_gc "gc: initialise CUDA state"- mtb <- keepAlive =<< MT.new+ $ do message "initialise CUDA state"+ etbl <- keepAlive =<< ET.new+ mtb <- keepAlive =<< MT.new etbl ktb <- keepAlive =<< KT.new rsv <- keepAlive =<< ST.new- return $! State mtb ktb rsv+ return $! State mtb ktb rsv etbl -- Select and initialise a default CUDA device, and create a new execution@@ -114,8 +126,16 @@ {-# NOINLINE defaultContext #-} defaultContext :: Context defaultContext = unsafePerformIO $ do- message dump_gc "gc: initialise default context"+ message "initialise default context" CUDA.initialise [] (dev,_) <- selectBestDevice create dev [CUDA.SchedAuto]+++-- Debug+-- -----++{-# INLINE message #-}+message :: String -> IO ()+message msg = traceIO dump_gc ("gc: " ++ msg)
accelerate-cuda.cabal view
@@ -1,11 +1,13 @@ Name: accelerate-cuda-Version: 0.16.0.0-Cabal-version: >= 1.8-Tested-with: GHC == 7.8.*+Version: 0.17.0.0+Cabal-version: >= 1.6+Tested-with: GHC >= 7.8 Build-type: Simple Synopsis: Accelerate backend for NVIDIA GPUs Description:+ __This backend has been deprecated in favour of <http://hackage.haskell.org/package/accelerate-llvm-ptx accelerate-llvm-ptx>.__+ . This library implements a backend for the /Accelerate/ language instrumented for parallel execution on CUDA-capable NVIDIA GPUs. For further information, refer to the main /Accelerate/ package:@@ -42,6 +44,7 @@ Data-files: cubits/accelerate_cuda.h cubits/accelerate_cuda_assert.h+ cubits/accelerate_cuda_exceptional.h cubits/accelerate_cuda_function.h cubits/accelerate_cuda_texture.h cubits/accelerate_cuda_type.h@@ -81,19 +84,20 @@ Default: False Library- Build-depends: accelerate == 0.15.1.*,+ Build-depends: accelerate == 1.0.*, array >= 0.3,- base >= 4.7 && < 4.9,+ base >= 4.7 && < 4.9.1, binary >= 0.7, bytestring >= 0.9,+ containers >= 0.3, cryptohash >= 0.7,- cuda == 0.7.0.0,+ cuda >= 0.7, directory >= 1.0, fclabels >= 2.0, filepath >= 1.0, hashable >= 1.1, hashtables >= 1.0.1,- language-c-quote >= 0.4.4 && < 0.9 || > 0.10.1,+ language-c-quote >= 0.4.4, mainland-pretty >= 0.2, mtl >= 2.0, old-time >= 1.0,@@ -102,7 +106,7 @@ SafeSemaphore >= 0.9, srcloc >= 0.2, text >= 0.11,- template-haskell >= 2.2,+ template-haskell, transformers >= 0.2, unordered-containers >= 0.1.4 @@ -119,14 +123,16 @@ Other-modules: Data.Array.Accelerate.CUDA.AST Data.Array.Accelerate.CUDA.Analysis.Device Data.Array.Accelerate.CUDA.Analysis.Launch+ Data.Array.Accelerate.CUDA.Analysis.Shape+ Data.Array.Accelerate.CUDA.Array.Remote Data.Array.Accelerate.CUDA.Array.Data- Data.Array.Accelerate.CUDA.Array.Nursery Data.Array.Accelerate.CUDA.Array.Prim+ Data.Array.Accelerate.CUDA.Array.Slice Data.Array.Accelerate.CUDA.Array.Sugar- Data.Array.Accelerate.CUDA.Array.Table- Data.Array.Accelerate.CUDA.Async Data.Array.Accelerate.CUDA.CodeGen+ Data.Array.Accelerate.CUDA.CodeGen.Arithmetic Data.Array.Accelerate.CUDA.CodeGen.Base+ Data.Array.Accelerate.CUDA.CodeGen.Constant Data.Array.Accelerate.CUDA.CodeGen.IndexSpace Data.Array.Accelerate.CUDA.CodeGen.Mapping Data.Array.Accelerate.CUDA.CodeGen.Monad@@ -134,6 +140,7 @@ Data.Array.Accelerate.CUDA.CodeGen.Reduction Data.Array.Accelerate.CUDA.CodeGen.Stencil Data.Array.Accelerate.CUDA.CodeGen.Stencil.Extra+ Data.Array.Accelerate.CUDA.CodeGen.Streaming Data.Array.Accelerate.CUDA.CodeGen.Type Data.Array.Accelerate.CUDA.Compile Data.Array.Accelerate.CUDA.Context@@ -143,7 +150,6 @@ Data.Array.Accelerate.CUDA.Execute.Stream Data.Array.Accelerate.CUDA.Foreign.Export Data.Array.Accelerate.CUDA.Foreign.Import- Data.Array.Accelerate.CUDA.FullList Data.Array.Accelerate.CUDA.Persistent Data.Array.Accelerate.CUDA.State Paths_accelerate_cuda@@ -172,9 +178,14 @@ -- -- Extensions: +source-repository head+ type: git+ location: https://github.com/AccelerateHS/accelerate-cuda+ source-repository this type: git+ tag: 0.17.0.0 location: https://github.com/AccelerateHS/accelerate-cuda- branch: release/0.16- tag: 0.16.0.0++-- vim: nospell
cubits/accelerate_cuda.h view
@@ -14,6 +14,7 @@ #define __ACCELERATE_CUDA_H__ #include "accelerate_cuda_assert.h"+#include "accelerate_cuda_exceptional.h" #include "accelerate_cuda_function.h" #include "accelerate_cuda_texture.h" #include "accelerate_cuda_type.h"
+ cubits/accelerate_cuda_exceptional.h view
@@ -0,0 +1,26 @@+/* -----------------------------------------------------------------------------+ *+ * Module : Exceptional+ * 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>+ * Stability : experimental+ *+ * ---------------------------------------------------------------------------*/++#ifndef __ACCELERATE_CUDA_EXCEPTIONAL_H__+#define __ACCELERATE_CUDA_EXCEPTIONAL_H__++/*+ * Exceptional values have slightly different names than what are produced by+ * the code generator.+ */+#include <math.h>++#define Infinity INFINITY+#define NaN NAN++#endif+
cubits/accelerate_cuda_function.h view
@@ -23,102 +23,6 @@ * -------------------------------------------------------------------------- */ /*- * Left/Right bitwise rotation- */-template <typename T>-static __inline__ __device__ T rotateL(const T x, const Int32 i)-{- const Int32 i8 = i & 8 * sizeof(x) - 1;- return i8 == 0 ? x : x << i8 | x >> 8 * sizeof(x) - i8;-}--template <typename T>-static __inline__ __device__ T rotateR(const T x, const Int32 i)-{- const Int32 i8 = i & 8 * sizeof(x) - 1;- return i8 == 0 ? x : x >> i8 | x << 8 * sizeof(x) - i8;-}--/*- * Integer division, truncated towards negative infinity- */-template <typename T>-static __inline__ __device__ T idiv(const T x, const T y)-{- // Stolen from GHC.Classes- if ( x > 0 && y < 0 ) {- return ((x-1) / y) - 1;- } else- if ( x < 0 && y > 0 ) {- return ((x+1) / y) - 1;- }- else {- return x / y;- }-}--template <>-__inline__ __device__ Word8 idiv(const Word8 x, const Word8 y)-{- return x / y;-}--template <>-__inline__ __device__ Word16 idiv(const Word16 x, const Word16 y)-{- return x / y;-}--template <>-__inline__ __device__ Word32 idiv(const Word32 x, const Word32 y)-{- return x / y;-}--template <>-__inline__ __device__ Word64 idiv(const Word64 x, const Word64 y)-{- return x / y;-}---/*- * Integer modulus, Haskell style- */-template <typename T>-static __inline__ __device__ T mod(const T x, const T y)-{- const T r = x % y;- return x > 0 && y < 0 || x < 0 && y > 0 ? (r != 0 ? r + y : 0) : r;-}--template <>-__inline__ __device__ Word8 mod(const Word8 x, const Word8 y)-{- return x % y;-}--template <>-__inline__ __device__ Word16 mod(const Word16 x, const Word16 y)-{- return x % y;-}--template <>-__inline__ __device__ Word32 mod(const Word32 x, const Word32 y)-{- return x % y;-}--template <>-__inline__ __device__ Word64 mod(const Word64 x, const Word64 y)-{- return x % y;-}----/* * Type coercion */ template <typename T>
cubits/accelerate_cuda_type.h view
@@ -30,3 +30,4 @@ typedef unsigned long long Word64; #endif+