diff --git a/Data/Array/Accelerate/CUDA.hs b/Data/Array/Accelerate/CUDA.hs --- a/Data/Array/Accelerate/CUDA.hs +++ b/Data/Array/Accelerate/CUDA.hs @@ -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 diff --git a/Data/Array/Accelerate/CUDA/AST.hs b/Data/Array/Accelerate/CUDA/AST.hs --- a/Data/Array/Accelerate/CUDA/AST.hs +++ b/Data/Array/Accelerate/CUDA/AST.hs @@ -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 "" + +{-- +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 +--} diff --git a/Data/Array/Accelerate/CUDA/Analysis/Launch.hs b/Data/Array/Accelerate/CUDA/Analysis/Launch.hs --- a/Data/Array/Accelerate/CUDA/Analysis/Launch.hs +++ b/Data/Array/Accelerate/CUDA/Analysis/Launch.hs @@ -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 diff --git a/Data/Array/Accelerate/CUDA/Analysis/Shape.hs b/Data/Array/Accelerate/CUDA/Analysis/Shape.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/CUDA/Analysis/Shape.hs @@ -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 +-- 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 + diff --git a/Data/Array/Accelerate/CUDA/Array/Data.hs b/Data/Array/Accelerate/CUDA/Array/Data.hs --- a/Data/Array/Accelerate/CUDA/Array/Data.hs +++ b/Data/Array/Accelerate/CUDA/Array/Data.hs @@ -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 diff --git a/Data/Array/Accelerate/CUDA/Array/Nursery.hs b/Data/Array/Accelerate/CUDA/Array/Nursery.hs deleted file mode 100644 --- a/Data/Array/Accelerate/CUDA/Array/Nursery.hs +++ /dev/null @@ -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 --- 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 () - diff --git a/Data/Array/Accelerate/CUDA/Array/Prim.hs b/Data/Array/Accelerate/CUDA/Array/Prim.hs --- a/Data/Array/Accelerate/CUDA/Array/Prim.hs +++ b/Data/Array/Accelerate/CUDA/Array/Prim.hs @@ -1,4 +1,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 () diff --git a/Data/Array/Accelerate/CUDA/Array/Remote.hs b/Data/Array/Accelerate/CUDA/Array/Remote.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/CUDA/Array/Remote.hs @@ -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 +-- 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 + diff --git a/Data/Array/Accelerate/CUDA/Array/Slice.hs b/Data/Array/Accelerate/CUDA/Array/Slice.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/CUDA/Array/Slice.hs @@ -0,0 +1,212 @@ +{-# LANGUAGE GADTs #-} +-- | +-- Module : Data.Array.Accelerate.Array.Slice +-- +-- Maintainer : Frederik Meisner Madsen +-- 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 } diff --git a/Data/Array/Accelerate/CUDA/Array/Sugar.hs b/Data/Array/Accelerate/CUDA/Array/Sugar.hs --- a/Data/Array/Accelerate/CUDA/Array/Sugar.hs +++ b/Data/Array/Accelerate/CUDA/Array/Sugar.hs @@ -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 diff --git a/Data/Array/Accelerate/CUDA/Array/Table.hs b/Data/Array/Accelerate/CUDA/Array/Table.hs deleted file mode 100644 --- a/Data/Array/Accelerate/CUDA/Array/Table.hs +++ /dev/null @@ -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 --- 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 () - diff --git a/Data/Array/Accelerate/CUDA/Async.hs b/Data/Array/Accelerate/CUDA/Async.hs deleted file mode 100644 --- a/Data/Array/Accelerate/CUDA/Async.hs +++ /dev/null @@ -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 --- 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 - diff --git a/Data/Array/Accelerate/CUDA/CodeGen.hs b/Data/Array/Accelerate/CUDA/CodeGen.hs --- a/Data/Array/Accelerate/CUDA/CodeGen.hs +++ b/Data/Array/Accelerate/CUDA/CodeGen.hs @@ -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" diff --git a/Data/Array/Accelerate/CUDA/CodeGen/Arithmetic.hs b/Data/Array/Accelerate/CUDA/CodeGen/Arithmetic.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/CUDA/CodeGen/Arithmetic.hs @@ -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 +-- 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) + diff --git a/Data/Array/Accelerate/CUDA/CodeGen/Base.hs b/Data/Array/Accelerate/CUDA/CodeGen/Base.hs --- a/Data/Array/Accelerate/CUDA/CodeGen/Base.hs +++ b/Data/Array/Accelerate/CUDA/CodeGen/Base.hs @@ -1,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 diff --git a/Data/Array/Accelerate/CUDA/CodeGen/Constant.hs b/Data/Array/Accelerate/CUDA/CodeGen/Constant.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/CUDA/CodeGen/Constant.hs @@ -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 +-- 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 + diff --git a/Data/Array/Accelerate/CUDA/CodeGen/IndexSpace.hs b/Data/Array/Accelerate/CUDA/CodeGen/IndexSpace.hs --- a/Data/Array/Accelerate/CUDA/CodeGen/IndexSpace.hs +++ b/Data/Array/Accelerate/CUDA/CodeGen/IndexSpace.hs @@ -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); |] ] diff --git a/Data/Array/Accelerate/CUDA/CodeGen/PrefixSum.hs b/Data/Array/Accelerate/CUDA/CodeGen/PrefixSum.hs --- a/Data/Array/Accelerate/CUDA/CodeGen/PrefixSum.hs +++ b/Data/Array/Accelerate/CUDA/CodeGen/PrefixSum.hs @@ -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 diff --git a/Data/Array/Accelerate/CUDA/CodeGen/Reduction.hs b/Data/Array/Accelerate/CUDA/CodeGen/Reduction.hs --- a/Data/Array/Accelerate/CUDA/CodeGen/Reduction.hs +++ b/Data/Array/Accelerate/CUDA/CodeGen/Reduction.hs @@ -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) diff --git a/Data/Array/Accelerate/CUDA/CodeGen/Streaming.hs b/Data/Array/Accelerate/CUDA/CodeGen/Streaming.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/CUDA/CodeGen/Streaming.hs @@ -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 +-- 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 ") + $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." + diff --git a/Data/Array/Accelerate/CUDA/CodeGen/Type.hs b/Data/Array/Accelerate/CUDA/CodeGen/Type.hs --- a/Data/Array/Accelerate/CUDA/CodeGen/Type.hs +++ b/Data/Array/Accelerate/CUDA/CodeGen/Type.hs @@ -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 diff --git a/Data/Array/Accelerate/CUDA/Compile.hs b/Data/Array/Accelerate/CUDA/Compile.hs --- a/Data/Array/Accelerate/CUDA/Compile.hs +++ b/Data/Array/Accelerate/CUDA/Compile.hs @@ -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) diff --git a/Data/Array/Accelerate/CUDA/Context.hs b/Data/Array/Accelerate/CUDA/Context.hs --- a/Data/Array/Accelerate/CUDA/Context.hs +++ b/Data/Array/Accelerate/CUDA/Context.hs @@ -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 diff --git a/Data/Array/Accelerate/CUDA/Debug.hs b/Data/Array/Accelerate/CUDA/Debug.hs --- a/Data/Array/Accelerate/CUDA/Debug.hs +++ b/Data/Array/Accelerate/CUDA/Debug.hs @@ -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 diff --git a/Data/Array/Accelerate/CUDA/Execute.hs b/Data/Array/Accelerate/CUDA/Execute.hs --- a/Data/Array/Accelerate/CUDA/Execute.hs +++ b/Data/Array/Accelerate/CUDA/Execute.hs @@ -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 " >>> " diff --git a/Data/Array/Accelerate/CUDA/Execute/Event.hs b/Data/Array/Accelerate/CUDA/Execute/Event.hs --- a/Data/Array/Accelerate/CUDA/Execute/Event.hs +++ b/Data/Array/Accelerate/CUDA/Execute/Event.hs @@ -1,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 diff --git a/Data/Array/Accelerate/CUDA/Execute/Stream.hs b/Data/Array/Accelerate/CUDA/Execute/Stream.hs --- a/Data/Array/Accelerate/CUDA/Execute/Stream.hs +++ b/Data/Array/Accelerate/CUDA/Execute/Stream.hs @@ -1,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 diff --git a/Data/Array/Accelerate/CUDA/Foreign/Export.hs b/Data/Array/Accelerate/CUDA/Foreign/Export.hs --- a/Data/Array/Accelerate/CUDA/Foreign/Export.hs +++ b/Data/Array/Accelerate/CUDA/Foreign/Export.hs @@ -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 diff --git a/Data/Array/Accelerate/CUDA/Foreign/Import.hs b/Data/Array/Accelerate/CUDA/Foreign/Import.hs --- a/Data/Array/Accelerate/CUDA/Foreign/Import.hs +++ b/Data/Array/Accelerate/CUDA/Foreign/Import.hs @@ -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 -- diff --git a/Data/Array/Accelerate/CUDA/FullList.hs b/Data/Array/Accelerate/CUDA/FullList.hs deleted file mode 100644 --- a/Data/Array/Accelerate/CUDA/FullList.hs +++ /dev/null @@ -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 --- 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_ #-} - diff --git a/Data/Array/Accelerate/CUDA/Persistent.hs b/Data/Array/Accelerate/CUDA/Persistent.hs --- a/Data/Array/Accelerate/CUDA/Persistent.hs +++ b/Data/Array/Accelerate/CUDA/Persistent.hs @@ -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) diff --git a/Data/Array/Accelerate/CUDA/State.hs b/Data/Array/Accelerate/CUDA/State.hs --- a/Data/Array/Accelerate/CUDA/State.hs +++ b/Data/Array/Accelerate/CUDA/State.hs @@ -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) diff --git a/accelerate-cuda.cabal b/accelerate-cuda.cabal --- a/accelerate-cuda.cabal +++ b/accelerate-cuda.cabal @@ -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 .__ + . 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 diff --git a/cubits/accelerate_cuda.h b/cubits/accelerate_cuda.h --- a/cubits/accelerate_cuda.h +++ b/cubits/accelerate_cuda.h @@ -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" diff --git a/cubits/accelerate_cuda_exceptional.h b/cubits/accelerate_cuda_exceptional.h new file mode 100644 --- /dev/null +++ b/cubits/accelerate_cuda_exceptional.h @@ -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 + * 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 + +#define Infinity INFINITY +#define NaN NAN + +#endif + diff --git a/cubits/accelerate_cuda_function.h b/cubits/accelerate_cuda_function.h --- a/cubits/accelerate_cuda_function.h +++ b/cubits/accelerate_cuda_function.h @@ -23,102 +23,6 @@ * -------------------------------------------------------------------------- */ /* - * Left/Right bitwise rotation - */ -template -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 -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 -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 -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 diff --git a/cubits/accelerate_cuda_type.h b/cubits/accelerate_cuda_type.h --- a/cubits/accelerate_cuda_type.h +++ b/cubits/accelerate_cuda_type.h @@ -30,3 +30,4 @@ typedef unsigned long long Word64; #endif +