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 @@ -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 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 @@ -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 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 @@ -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 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 @@ -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 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 @@ -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 diff --git a/Data/Array/Accelerate/CUDA/Array/Table.hs b/Data/Array/Accelerate/CUDA/Array/Table.hs --- a/Data/Array/Accelerate/CUDA/Array/Table.hs +++ b/Data/Array/Accelerate/CUDA/Array/Table.hs @@ -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 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 @@ -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 = 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 @@ -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] 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 @@ -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) 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 @@ -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 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 @@ -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 ) 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 @@ -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 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 @@ -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 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 @@ -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 diff --git a/accelerate-cuda.cabal b/accelerate-cuda.cabal --- a/accelerate-cuda.cabal +++ b/accelerate-cuda.cabal @@ -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 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 @@ -45,7 +45,16 @@ template 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 <>