accelerate-cuda 0.15.0.0 → 0.15.1.0
raw patch · 16 files changed
+193/−78 lines, 16 filesdep ~acceleratedep ~basedep ~cuda
Dependency ranges changed: accelerate, base, cuda, language-c-quote
Files
- Data/Array/Accelerate/CUDA.hs +35/−0
- Data/Array/Accelerate/CUDA/AST.hs +2/−1
- Data/Array/Accelerate/CUDA/Array/Data.hs +23/−3
- Data/Array/Accelerate/CUDA/Array/Prim.hs +1/−1
- Data/Array/Accelerate/CUDA/Array/Sugar.hs +6/−5
- Data/Array/Accelerate/CUDA/Array/Table.hs +27/−14
- Data/Array/Accelerate/CUDA/CodeGen.hs +19/−6
- Data/Array/Accelerate/CUDA/CodeGen/Base.hs +11/−11
- Data/Array/Accelerate/CUDA/Compile.hs +20/−8
- Data/Array/Accelerate/CUDA/Execute.hs +8/−3
- Data/Array/Accelerate/CUDA/Foreign/Export.hs +5/−3
- Data/Array/Accelerate/CUDA/Foreign/Import.hs +4/−3
- Data/Array/Accelerate/CUDA/Persistent.hs +10/−9
- Data/Array/Accelerate/CUDA/State.hs +4/−4
- accelerate-cuda.cabal +8/−6
- cubits/accelerate_cuda_function.h +10/−1
Data/Array/Accelerate/CUDA.hs view
@@ -194,6 +194,7 @@ -- * Execution contexts Context, create, destroy,+ unsafeFree, unsafeFreeIn, performGC, performGCIn, ) where @@ -202,6 +203,7 @@ import Control.Applicative import Control.Monad.Trans import System.IO.Unsafe+import Prelude -- friends import Data.Array.Accelerate.Trafo@@ -390,4 +392,37 @@ #else dumpStats next = return next #endif+++-- Device memory management+-- ------------------------+--+-- Temporarily defining here, until we can define the interface for it.+++-- Deallocate the device arrays corresponding to the given host side arrays.+-- This is unsafe in the sense that it is possible to call this function while+-- the array is currently in use.+--+unsafeFree :: Arrays arrs => arrs -> IO ()+unsafeFree = unsafeFreeIn defaultContext++unsafeFreeIn :: forall arrs. Arrays arrs => Context -> arrs -> IO ()+unsafeFreeIn !ctx !arrs+ = evalCUDA ctx+ $ freeR (arrays (undefined :: arrs)) (fromArr arrs)+ where+ freeR :: ArraysR a -> a -> CIO ()+ freeR ArraysRunit () = return ()+ freeR ArraysRarray arr = freeArray arr+ freeR (ArraysRpair aeR1 aeR2) (arrs1, arrs2) = freeR aeR1 arrs1 >> freeR aeR2 arrs2+++-- Release any unused device memory+--+performGC :: IO ()+performGC = performGCIn defaultContext++performGCIn :: Context -> IO ()+performGCIn !ctx = evalCUDA ctx cleanupArrayData
Data/Array/Accelerate/CUDA/AST.hs view
@@ -35,9 +35,10 @@ -- system import Text.PrettyPrint import Data.Hashable-import Data.Monoid ( Monoid(..) )+import Data.Monoid hiding ( (<>) ) import qualified Data.HashSet as Set import qualified Data.HashMap.Strict as Map+import Prelude -- A non-empty list of binary objects will be used to execute a kernel. We keep
Data/Array/Accelerate/CUDA/Array/Data.hs view
@@ -19,7 +19,8 @@ module Data.Array.Accelerate.CUDA.Array.Data ( -- * Array operations and representations- mallocArray, indexArray,+ mallocArray, freeArray,+ indexArray, useArray, useArrayAsync, useDevicePtrs, copyArray, copyArrayAsync, copyArrayPeer, copyArrayPeerAsync,@@ -35,14 +36,14 @@ ) where -- libraries-import Prelude hiding ( fst, snd )-import qualified Prelude as P import Control.Applicative import Control.Monad.Reader ( asks ) import Control.Monad.State ( gets ) import Control.Monad.Trans ( liftIO ) import Foreign.C.Types import Foreign.Ptr+import Prelude hiding ( fst, snd )+import qualified Prelude as P -- friends import Data.Array.Accelerate.Error@@ -131,6 +132,25 @@ -- mallocPrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> Int -> IO () mkPrimDispatch(mallocPrim,Prim.mallocArray)+++-- |Deallocate the device array accompanying the given host-side array.+--+-- Note that this does not take into account whether or not the data is still+-- required by the current (or future) computation.+--+freeArray :: Array dim e -> CIO ()+freeArray (Array !_ !adata) = run doFree+ where+ doFree !ctx !mt = freeR arrayElt adata+ where+ freeR :: ArrayEltR e -> ArrayData e -> IO ()+ freeR ArrayEltRunit _ = return ()+ freeR (ArrayEltRpair aeR1 aeR2) ad = freeR aeR1 (fst ad) >> freeR aeR2 (snd ad)+ freeR aer ad = freePrim aer ctx mt ad+ --+ freePrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> IO ()+ mkPrimDispatch(freePrim,free) -- |Upload an existing array to the device
Data/Array/Accelerate/CUDA/Array/Prim.hs view
@@ -31,7 +31,6 @@ ) where -- libraries-import Prelude hiding ( lookup ) import Data.Int import Data.Word import Data.Maybe@@ -47,6 +46,7 @@ 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
Data/Array/Accelerate/CUDA/Array/Sugar.hs view
@@ -16,6 +16,8 @@ ) where +import Control.Monad.Trans+ import Data.Array.Accelerate.CUDA.State import Data.Array.Accelerate.CUDA.Array.Data import Data.Array.Accelerate.Array.Sugar hiding (newArray, allocateArray)@@ -36,9 +38,8 @@ -- Allocate a new, uninitialised Accelerate array on host and device -- allocateArray :: (Shape dim, Elt e) => dim -> CIO (Array dim e)-allocateArray sh =- let arr = Sugar.allocateArray sh- in do- mallocArray arr- return arr+allocateArray sh = do+ arr <- liftIO $ Sugar.allocateArray sh+ mallocArray arr+ return arr
Data/Array/Accelerate/CUDA/Array/Table.hs view
@@ -17,26 +17,26 @@ module Data.Array.Accelerate.CUDA.Array.Table ( -- Tables for host/device memory associations- MemoryTable, new, lookup, malloc, insert, insertRemote, reclaim+ MemoryTable, new, lookup, malloc, free, insert, insertRemote, reclaim ) where -import Prelude hiding ( lookup )-import Data.Maybe ( isJust )-import Data.Hashable ( Hashable(..) )-import Data.Typeable ( Typeable, gcast )-import Control.Monad ( unless )+import Control.Applicative import Control.Concurrent ( yield ) import Control.Concurrent.MVar ( MVar, newMVar, withMVar, mkWeakMVar ) import Control.Exception ( bracket_, catch, throwIO )-import Control.Applicative ( (<$>) )+import Control.Monad ( unless )+import Data.Hashable ( Hashable(..) )+import Data.Maybe ( isJust )+import Data.Typeable ( Typeable, gcast ) import System.Mem ( performGC )-import System.Mem.Weak ( Weak, mkWeak, deRefWeak, finalize ) 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@@ -156,7 +156,7 @@ 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 = 128+ chunk = 1024 !n' = chunk * multiple (fromIntegral n) (fromIntegral chunk) !bytes = n' * sizeOf (undefined :: b)@@ -173,6 +173,19 @@ 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. --@@ -204,7 +217,7 @@ -- reclaim :: MemoryTable -> IO () reclaim (MemoryTable _ weak_ref (Nursery nrs _)) = do- (free, total) <- CUDA.getMemInfo+ (before, total) <- CUDA.getMemInfo performGC yield withMVar nrs N.flush@@ -217,9 +230,9 @@ unless alive $ finalize w -- D.when D.dump_gc $ do- (free', _) <- CUDA.getMemInfo- message $ "reclaim: freed " ++ showBytes (fromIntegral (free - free'))- ++ ", " ++ showBytes (fromIntegral free')+ (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
Data/Array/Accelerate/CUDA/CodeGen.hs view
@@ -23,16 +23,16 @@ ) where -- libraries-import Prelude hiding ( id, exp, replicate )-import Control.Applicative ( (<$>), (<*>) )-import Control.Monad.State.Strict import Data.Loc import Data.Char import Data.HashSet ( HashSet )+import Control.Monad.State.Strict import Foreign.CUDA.Analysis import Language.C.Quote.CUDA import qualified Language.C as C import qualified Data.HashSet as Set+import Control.Applicative hiding ( Const )+import Prelude hiding ( id, exp, replicate ) -- friends import Data.Array.Accelerate.Error@@ -824,10 +824,23 @@ codegenSig (FloatingNumType ty) = codegenFloatingSig ty codegenIntegralSig :: IntegralType a -> C.Exp -> C.Exp-codegenIntegralSig ty x = [cexp|$exp:x == $exp:zero ? $exp:zero : $exp:(ccall "copysign" [one,x]) |]+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- zero | IntegralDict <- integralDict ty = codegenIntegralScalar ty 0- one | IntegralDict <- integralDict ty = codegenIntegralScalar ty 1+ 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 =
Data/Array/Accelerate/CUDA/CodeGen/Base.hs view
@@ -268,11 +268,11 @@ -- writeArray :: forall sh e. (Shape sh, Elt e)- => Name -- group names- -> Array sh e -- dummy to fix types- -> ( [C.Param] -- function parameters to marshal the output array- , [C.Exp] -- the shape of the output array- , Rvalue x => x -> [C.Exp] ) -- write an element at a given index+ => Name -- group names+ -> Array sh e -- dummy to fix types+ -> ( [C.Param] -- function parameters to marshal the output array+ , [C.Exp] -- the shape of the output array+ , forall x. Rvalue x => x -> [C.Exp] ) -- write an element at a given index writeArray grp _ = let (sh, arrs) = namesOfArray grp (undefined :: e) dim = expDim (undefined :: Exp aenv sh)@@ -293,12 +293,12 @@ -- shared :: forall e. Elt e- => e -- dummy type- -> Name -- group name- -> C.Exp -- how much shared memory per type- -> Maybe C.Exp -- (optional) initialise from this base address- -> ( [C.InitGroup] -- shared memory declaration and...- , Rvalue x => x -> [C.Exp]) -- ...indexing function+ => e -- dummy type+ -> Name -- group name+ -> C.Exp -- how much shared memory per type+ -> Maybe C.Exp -- (optional) initialise from this base address+ -> ( [C.InitGroup] -- shared memory declaration and...+ , forall x. Rvalue x => x -> [C.Exp]) -- ...indexing function shared _ grp size mprev = let e:es = eltType (undefined :: e) x:xs = let k = length es in map (\n -> grp ++ show n) [k, k-1 .. 0]
Data/Array/Accelerate/CUDA/Compile.hs view
@@ -40,7 +40,6 @@ -- libraries import Numeric-import Prelude hiding ( exp, scanl, scanr ) import Control.Applicative hiding ( Const ) import Control.Exception import Control.Monad@@ -59,8 +58,8 @@ import System.IO import System.IO.Error import System.IO.Unsafe-import System.Process import System.Mem.Weak+import System.Process import Text.PrettyPrint.Mainland ( ppr, renderCompact, displayLazyText ) import qualified Data.ByteString as B import qualified Data.Text.Lazy as T@@ -69,16 +68,21 @@ import qualified Control.Concurrent.MSem as Q import qualified Foreign.CUDA.Driver as CUDA import qualified Foreign.CUDA.Analysis as CUDA+import Prelude hiding ( exp, scanl, scanr ) import GHC.Conc ( getNumProcessors ) #ifdef ACCELERATE_DEBUG import System.Time #endif-#ifdef VERSION_unix++-- Multiplatform support for dealing with external process spawning+#if defined(UNIX) import System.Posix.Process+#elif defined(WIN32)+import System.Win32.Process hiding (ProcessHandle) #else-import System.Win32.Process+#error "I don't know what operating system I am" #endif import Paths_accelerate_cuda ( getDataDir )@@ -387,8 +391,8 @@ -- the binary object. -- message "waiting for nvcc..."- takeMVar done let cubin = replaceExtension cufile ".cubin"+ () <- takeMVar done bin <- B.readFile cubin mdl <- CUDA.loadData bin addFinalizer mdl (module_finalizer weak_ctx key mdl)@@ -486,9 +490,13 @@ createDirectoryIfMissing True dir openTempFile dir template -#ifndef VERSION_unix-getProcessID :: ProcessHandle -> IO ProcessId-getProcessID = getProcessId+#if defined(WIN32)+-- TLM: On windows, how do we get either the ProcessID or ProcessHandle of the+-- current process? For new, just use a dummy value (the sound of+-- disappearing down a rabbit hole...)+--+getProcessID :: IO ProcessId+getProcessID = return 0xaaaa #endif @@ -518,6 +526,10 @@ -- ... and wait for it to complete waitFor pid+ -- If compilation fails for some reason, fill the MVar by re-throwing+ -- the exception. This prevents the host thread from waiting+ -- indefinitely, which then requires the program to be killed manually.+ `catch` \(e :: SomeException) -> do putMVar mvar (throw e) ccEnd <- getTime return (diffTime ccBegin ccEnd)
Data/Array/Accelerate/CUDA/Execute.hs view
@@ -52,7 +52,6 @@ -- standard library-import Prelude hiding ( exp, sum, iterate ) import Control.Applicative hiding ( Const ) import Control.Monad ( join, when, liftM ) import Control.Monad.Reader ( asks )@@ -61,7 +60,7 @@ import System.IO.Unsafe ( unsafeInterleaveIO ) import Data.Int import Data.Word-import Data.Maybe+import Prelude hiding ( exp, sum, iterate ) import Foreign.CUDA.Analysis.Device ( computeCapability, Compute(..) ) import qualified Foreign.CUDA.Driver as CUDA@@ -177,7 +176,7 @@ Awhile p f a -> awhile p f =<< travA a -- Foreign- Aforeign ff afun a -> fromMaybe (executeAfun1 afun) (canExecuteAcc ff) =<< travA a+ Aforeign ff afun a -> aforeign ff afun =<< travA a -- Producers Map _ a -> executeOp =<< extent a@@ -227,6 +226,12 @@ 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 =+ case canExecuteAcc ff of+ Just cudaFun -> cudaFun stream a+ Nothing -> executeAfun1 pureFun a -- get the extent of an embedded array extent :: Shape sh => ExecOpenAcc aenv (Array sh e) -> CIO sh
Data/Array/Accelerate/CUDA/Foreign/Export.hs view
@@ -9,7 +9,8 @@ {-# LANGUAGE TypeFamilies #-} {-# LANGUAGE FlexibleInstances #-} {-# LANGUAGE ImpredicativeTypes #-}-{-# OPTIONS_GHC -fno-warn-name-shadowing -fno-warn-orphans #-}+{-# OPTIONS_GHC -fno-warn-name-shadowing #-}+{-# OPTIONS_GHC -fno-warn-orphans #-} -- | -- Module : Data.Array.Accelerate.CUDA.Foreign.Export -- Copyright : [2013..2014] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell, Robert Clifton-Everest@@ -38,7 +39,6 @@ ) where -import Prelude as P import Data.Functor import Control.Applicative import Foreign.StablePtr@@ -48,8 +48,10 @@ import Foreign.Marshal.Array ( peekArray, pokeArray, mallocArray ) import Foreign.Marshal.Alloc ( free ) import Control.Monad.State ( liftIO )-import qualified Foreign.CUDA.Driver as CUDA import Language.Haskell.TH hiding ( ppr )+import Prelude as P++import qualified Foreign.CUDA.Driver as CUDA -- friends import Data.Array.Accelerate.Smart ( Acc )
Data/Array/Accelerate/CUDA/Foreign/Import.hs view
@@ -53,7 +53,7 @@ allocateArray, newArray, -- * Running IO actions in an Accelerate context- CIO, liftIO, inContext, inDefaultContext+ CIO, Stream, liftIO, inContext, inDefaultContext ) where @@ -63,6 +63,7 @@ import Data.Array.Accelerate.CUDA.Array.Sugar import Data.Array.Accelerate.CUDA.Array.Data import Data.Array.Accelerate.CUDA.Array.Prim ( DevicePtrs )+import Data.Array.Accelerate.CUDA.Execute.Stream ( Stream ) import Data.Typeable import Control.Exception ( bracket_ )@@ -76,7 +77,7 @@ -- data CUDAForeignAcc as bs where CUDAForeignAcc :: String -- name of the function- -> (as -> CIO bs) -- operation to execute+ -> (Stream -> as -> CIO bs) -- operation to execute -> CUDAForeignAcc as bs deriving instance Typeable CUDAForeignAcc@@ -90,7 +91,7 @@ canExecuteAcc :: (Foreign f, Typeable as, Typeable bs) => f as bs- -> Maybe (as -> CIO bs)+ -> Maybe (Stream -> as -> CIO bs) canExecuteAcc ff | Just (CUDAForeignAcc _ fun) <- cast ff = Just fun
Data/Array/Accelerate/CUDA/Persistent.hs view
@@ -30,28 +30,29 @@ import qualified Data.Array.Accelerate.CUDA.FullList as FL -- libraries-import Prelude hiding ( lookup ) import Numeric-import Data.Char-import System.IO-import System.FilePath-import System.Directory-import System.IO.Error-import System.Mem.Weak import Control.Applicative import Control.Concurrent import Control.Exception import Control.Monad.Trans-import Data.Version import Data.Binary-import Data.Hashable import Data.Binary.Get import Data.ByteString ( ByteString ) import Data.ByteString.Internal ( w2c )+import Data.Char+import Data.Hashable+import Data.Maybe+import Data.Version+import System.Directory+import System.FilePath+import System.IO+import System.IO.Error+import System.Mem.Weak import qualified Data.ByteString as BS import qualified Data.ByteString.Lazy as BL import qualified Data.ByteString.Lazy.Internal as BL import qualified Data.HashTable.IO as HT+import Prelude hiding ( lookup ) import qualified Foreign.CUDA.Driver as CUDA
Data/Array/Accelerate/CUDA/State.hs view
@@ -37,16 +37,16 @@ import Data.Array.Accelerate.CUDA.Analysis.Device -- library-import Control.Applicative ( Applicative )+import Control.Applicative import Control.Concurrent ( runInBoundThread ) import Control.Exception ( catch, bracket_ )-import Control.Monad.Trans ( MonadIO ) import Control.Monad.Reader ( MonadReader, ReaderT(..), runReaderT ) import Control.Monad.State.Strict ( MonadState, StateT(..), evalStateT )-import System.Mem ( performGC )+import Control.Monad.Trans ( MonadIO ) import System.IO.Unsafe ( unsafePerformIO ) import Foreign.CUDA.Driver.Error import qualified Foreign.CUDA.Driver as CUDA+import Prelude -- Execution State@@ -85,7 +85,7 @@ \e -> $internalError "unhandled" (show (e :: CUDAException)) where setup = push ctx- teardown = pop >> performGC+ teardown = pop action = evalStateT (runReaderT (runCIO acc) ctx) theState
accelerate-cuda.cabal view
@@ -1,6 +1,6 @@ Name: accelerate-cuda-Version: 0.15.0.0-Cabal-version: >= 1.6+Version: 0.15.1.0+Cabal-version: >= 1.8 Tested-with: GHC == 7.8.* Build-type: Simple @@ -81,9 +81,9 @@ Default: False Library- Build-depends: accelerate == 0.15.*,+ Build-depends: accelerate == 0.15.1.*, array >= 0.3,- base == 4.7.*,+ base >= 4.7 && < 4.9, binary >= 0.7, bytestring >= 0.9, cryptohash >= 0.7,@@ -93,7 +93,7 @@ filepath >= 1.0, hashable >= 1.1, hashtables >= 1.0.1,- language-c-quote >= 0.4.4,+ language-c-quote >= 0.4.4 && < 0.9 || > 0.10.1, mainland-pretty >= 0.2, mtl >= 2.0, old-time >= 1.0,@@ -172,7 +172,9 @@ -- -- Extensions: -source-repository head+source-repository this type: git location: https://github.com/AccelerateHS/accelerate-cuda+ branch: release/0.15+ tag: 0.15.1.0
cubits/accelerate_cuda_function.h view
@@ -45,7 +45,16 @@ template <typename T> static __inline__ __device__ T idiv(const T x, const T y) {- return x > 0 && y < 0 ? (x - y - 1) / y : (x < 0 && y > 0 ? (x - y + 1) / y : x / 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 <>