accelerate-cuda 0.12.0.0 → 0.12.1.0
raw patch · 27 files changed
+1124/−681 lines, 27 filesdep −symboldep ~accelerate
Dependencies removed: symbol
Dependency ranges changed: accelerate
Files
- Data/Array/Accelerate/CUDA.hs +47/−8
- Data/Array/Accelerate/CUDA/Array/Data.hs +81/−56
- Data/Array/Accelerate/CUDA/Array/Prim.hs +87/−68
- Data/Array/Accelerate/CUDA/Array/Table.hs +55/−34
- Data/Array/Accelerate/CUDA/CodeGen.hs +8/−4
- Data/Array/Accelerate/CUDA/CodeGen/Base.hs +2/−2
- Data/Array/Accelerate/CUDA/CodeGen/IndexSpace.hs +1/−0
- Data/Array/Accelerate/CUDA/CodeGen/Monad.hs +7/−2
- Data/Array/Accelerate/CUDA/CodeGen/PrefixSum.hs +14/−8
- Data/Array/Accelerate/CUDA/CodeGen/Reduction.hs +19/−16
- Data/Array/Accelerate/CUDA/CodeGen/Stencil.hs +2/−2
- Data/Array/Accelerate/CUDA/CodeGen/Type.hs +36/−50
- Data/Array/Accelerate/CUDA/Compile.hs +46/−46
- Data/Array/Accelerate/CUDA/Debug.hs +17/−10
- Data/Array/Accelerate/CUDA/Execute.hs +39/−37
- Data/Array/Accelerate/CUDA/Persistent.hs +274/−0
- Data/Array/Accelerate/CUDA/State.hs +50/−93
- accelerate-cuda.buildinfo.in +3/−0
- accelerate-cuda.cabal +29/−14
- configure +223/−152
- cubits/accelerate_cuda_extras.h +1/−0
- cubits/accelerate_cuda_function.h +20/−18
- cubits/accelerate_cuda_shape.h +3/−3
- cubits/accelerate_cuda_shape.h.in +3/−3
- cubits/accelerate_cuda_stencil.h +3/−3
- cubits/accelerate_cuda_texture.h +22/−52
- cubits/accelerate_cuda_type.h +32/−0
Data/Array/Accelerate/CUDA.hs view
@@ -13,9 +13,34 @@ -- Accelerate. Expressions are on-line translated into CUDA code, compiled, and -- executed in parallel on the GPU. ----- The accelerate-cuda library is hosted at <https://github.com/tmcdonell/accelerate-cuda>.--- Comments, bug reports, and patches are always welcome.+-- The accelerate-cuda library is hosted at: <https://github.com/AccelerateHS/accelerate-cuda>.+-- Comments, bug reports, and patches, are always welcome. --+--+-- /NOTES:/+--+-- CUDA devices are categorised into different \'compute capabilities\',+-- indicating what operations are supported by the hardware. For example, double+-- precision arithmetic is only supported on devices of compute capability 1.3+-- or higher.+--+-- Devices generally perform best when dealing with (tuples of) 32-bit types, so+-- be cautious when introducing 8-, 16-, or 64-bit elements. Keep in mind the+-- size of 'Int' and 'Data.Word.Word' changes depending on the architecture GHC+-- runs on.+--+-- Additional notes:+--+-- * 'Double' precision requires compute-1.3.+--+-- * 'Bool' is represented internally using 'Data.Word.Word8', 'Char' by+-- 'Data.Word.Word32'.+--+-- * If the permutation function to 'Data.Array.Accelerate.permute' resolves to+-- non-unique indices, the combination function requires compute-1.1 to+-- combine 32-bit types, or compute-1.2 for 64-bit types. Tuple components+-- are resolved separately.+-- module Data.Array.Accelerate.CUDA ( @@ -57,10 +82,18 @@ -- This will select the fastest device available on which to execute -- computations, based on compute capability and estimated maximum GFLOPS. --+-- /NOTE:/+-- GPUs typically have their own attached memory, which is separate from the+-- computer's main memory. Hence, every 'Data.Array.Accelerate.use' operation+-- implies copying data to the device, and every 'run' operation must copy the+-- results of a computation back to the host. Thus, it is best to keep all+-- computations in the 'Acc' meta-language form and only 'run' the computation+-- once at the end, to avoid transferring (unused) intermediate results.+-- run :: Arrays a => Acc a -> a run a = unsafePerformIO- $ withMVar defaultContext $ \ctx -> evaluate (runIn ctx a)+ $ 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@@ -72,7 +105,7 @@ runAsync :: Arrays a => Acc a -> Async a runAsync a = unsafePerformIO- $ withMVar defaultContext $ \ctx -> evaluate (runAsyncIn ctx a)+ $ evaluate (runAsyncIn defaultContext a) -- | As 'run', but execute using the specified device context rather than using -- the default, automatically selected device.@@ -111,15 +144,21 @@ -- This function can be used to improve performance in cases where the array -- program is constant between invocations, because it allows us to bypass all -- front-end conversion stages and move directly to the execution phase. If you--- have a computation applied repeatedly to different input data, use this.+-- have a computation applied repeatedly to different input data, use this. If+-- the function is only evaluated once, this is equivalent to 'run'. --+-- > let step :: Vector a -> Vector b+-- > step = run1 f+-- > in+-- > simulate step ...+-- -- See the Crystal demo, part of the 'accelerate-examples' package, for an -- example. -- run1 :: (Arrays a, Arrays b) => (Acc a -> Acc b) -> a -> b run1 f = unsafePerformIO- $ withMVar defaultContext $ \ctx -> evaluate (run1In ctx f)+ $ evaluate (run1In defaultContext f) -- | As 'run1', but the computation is executed asynchronously.@@ -127,7 +166,7 @@ run1Async :: (Arrays a, Arrays b) => (Acc a -> Acc b) -> a -> Async b run1Async f = unsafePerformIO- $ withMVar defaultContext $ \ctx -> evaluate (run1AsyncIn ctx f)+ $ evaluate (run1AsyncIn defaultContext f) -- | As 'run1', but execute in the specified context. --@@ -156,7 +195,7 @@ stream :: (Arrays a, Arrays b) => (Acc a -> Acc b) -> [a] -> [b] stream f arrs = unsafePerformIO- $ withMVar defaultContext $ \ctx -> evaluate (streamIn ctx f arrs)+ $ evaluate (streamIn defaultContext f arrs) -- | As 'stream', but execute in the specified context. --
Data/Array/Accelerate/CUDA/Array/Data.hs view
@@ -1,4 +1,7 @@-{-# LANGUAGE CPP, GADTs, TypeFamilies, ScopedTypeVariables #-}+{-# LANGUAGE CPP #-}+{-# LANGUAGE GADTs #-}+{-# LANGUAGE ScopedTypeVariables #-}+{-# LANGUAGE TypeFamilies #-} -- | -- Module : Data.Array.Accelerate.CUDA.Array.Data -- Copyright : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee@@ -26,15 +29,15 @@ ) where -- libraries-import Prelude hiding (fst, snd)+import Prelude hiding ( fst, snd ) import Data.Label.PureM import Control.Applicative import Control.Monad.Trans -- friends import Data.Array.Accelerate.Array.Data-import Data.Array.Accelerate.Array.Sugar (Array(..), Shape, Elt, fromElt, toElt)-import Data.Array.Accelerate.Array.Representation (size, index)+import Data.Array.Accelerate.Array.Sugar ( Array(..), Shape, Elt, fromElt, toElt )+import Data.Array.Accelerate.Array.Representation ( size, index ) import Data.Array.Accelerate.CUDA.State import Data.Array.Accelerate.CUDA.Array.Table import qualified Data.Array.Accelerate.CUDA.Array.Prim as Prim@@ -61,6 +64,14 @@ snd :: ArrayData (a,b) -> ArrayData b snd = sndArrayData +-- Extract the state information to pass along to the primitive data handlers+--+run :: (Context -> MemoryTable -> IO a) -> CIO a+run f = do+ ctx <- gets activeContext+ mt <- gets memoryTable+ liftIO $ f ctx mt+ -- CPP hackery to generate the cases where we dispatch to the worker function handling -- elementary types. --@@ -77,53 +88,53 @@ ; dispatcher ArrayEltRword64 = worker \ ; dispatcher ArrayEltRfloat = worker \ ; dispatcher ArrayEltRdouble = worker \-; dispatcher ArrayEltRbool = error "mkPrimDispatcher: ArrayEltRbool" \-; dispatcher ArrayEltRchar = error "mkPrimDispatcher: ArrayEltRchar" \+; dispatcher ArrayEltRbool = worker \+; dispatcher ArrayEltRchar = worker \ ; dispatcher _ = error "mkPrimDispatcher: not primitive" -- |Allocate a new device array to accompany the given host-side array. -- mallocArray :: (Shape dim, Elt e) => Array dim e -> CIO ()-mallocArray (Array sh adata) = doMalloc =<< gets memoryTable+mallocArray (Array sh adata) = run doMalloc where- doMalloc mt = liftIO $ mallocR arrayElt adata+ doMalloc ctx mt = mallocR arrayElt adata where mallocR :: ArrayEltR e -> ArrayData e -> IO () mallocR ArrayEltRunit _ = return () mallocR (ArrayEltRpair aeR1 aeR2) ad = mallocR aeR1 (fst ad) >> mallocR aeR2 (snd ad)- mallocR aer ad = mallocPrim aer mt ad (size sh)+ mallocR aer ad = mallocPrim aer ctx mt ad (size sh) --- mallocPrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> Int -> IO ()+ mallocPrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> Int -> IO () mkPrimDispatch(mallocPrim,Prim.mallocArray) -- |Upload an existing array to the device -- useArray :: (Shape dim, Elt e) => Array dim e -> CIO ()-useArray (Array sh adata) = doUse =<< gets memoryTable+useArray (Array sh adata) = run doUse where- doUse mt = liftIO $ useR arrayElt adata+ doUse ctx mt = useR arrayElt adata where useR :: ArrayEltR e -> ArrayData e -> IO () useR ArrayEltRunit _ = return () useR (ArrayEltRpair aeR1 aeR2) ad = useR aeR1 (fst ad) >> useR aeR2 (snd ad)- useR aer ad = usePrim aer mt ad (size sh)+ useR aer ad = usePrim aer ctx mt ad (size sh) --- usePrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> Int -> IO ()+ usePrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> Int -> IO () mkPrimDispatch(usePrim,Prim.useArray) useArrayAsync :: (Shape dim, Elt e) => Array dim e -> Maybe CUDA.Stream -> CIO ()-useArrayAsync (Array sh adata) ms = doUse =<< gets memoryTable+useArrayAsync (Array sh adata) ms = run doUse where- doUse mt = liftIO $ useR arrayElt adata+ doUse ctx mt = useR arrayElt adata where useR :: ArrayEltR e -> ArrayData e -> IO () useR ArrayEltRunit _ = return () useR (ArrayEltRpair aeR1 aeR2) ad = useR aeR1 (fst ad) >> useR aeR2 (snd ad)- useR aer ad = usePrim aer mt ad (size sh) ms+ useR aer ad = usePrim aer ctx mt ad (size sh) ms --- usePrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> Int -> Maybe CUDA.Stream -> IO ()+ usePrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> Int -> Maybe CUDA.Stream -> IO () mkPrimDispatch(usePrim,Prim.useArrayAsync) @@ -131,19 +142,33 @@ -- synchronous operation. -- indexArray :: (Shape dim, Elt e) => Array dim e -> dim -> CIO e-indexArray (Array sh adata) ix = doIndex =<< gets memoryTable+indexArray (Array sh adata) ix = run doIndex where- i = index sh (fromElt ix)- doIndex mt = toElt <$> (liftIO $ indexR arrayElt adata)+ i = index sh (fromElt ix)+ doIndex ctx mt = toElt <$> indexR arrayElt adata where indexR :: ArrayEltR e -> ArrayData e -> IO e indexR ArrayEltRunit _ = return () indexR (ArrayEltRpair aeR1 aeR2) ad = (,) <$> indexR aeR1 (fst ad) <*> indexR aeR2 (snd ad)- indexR aer ad = indexPrim aer mt ad i --- indexPrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> Int -> IO e- mkPrimDispatch(indexPrim,Prim.indexArray)+ indexR ArrayEltRbool ad = toBool <$> Prim.indexArray ctx mt ad i+ where toBool 0 = False+ toBool _ = True+ --+ indexR ArrayEltRint ad = Prim.indexArray ctx mt ad i+ indexR ArrayEltRint8 ad = Prim.indexArray ctx mt ad i+ indexR ArrayEltRint16 ad = Prim.indexArray ctx mt ad i+ indexR ArrayEltRint32 ad = Prim.indexArray ctx mt ad i+ indexR ArrayEltRint64 ad = Prim.indexArray ctx mt ad i+ indexR ArrayEltRword ad = Prim.indexArray ctx mt ad i+ indexR ArrayEltRword8 ad = Prim.indexArray ctx mt ad i+ indexR ArrayEltRword16 ad = Prim.indexArray ctx mt ad i+ indexR ArrayEltRword32 ad = Prim.indexArray ctx mt ad i+ indexR ArrayEltRword64 ad = Prim.indexArray ctx mt ad i+ indexR ArrayEltRfloat ad = Prim.indexArray ctx mt ad i+ indexR ArrayEltRdouble ad = Prim.indexArray ctx mt ad i+ indexR ArrayEltRchar ad = Prim.indexArray ctx mt ad i -- |Copy data between two device arrays. The operation is asynchronous with@@ -152,75 +177,75 @@ copyArray :: (Shape dim, Elt e) => Array dim e -> Array dim e -> CIO () copyArray (Array sh1 adata1) (Array sh2 adata2) = BOUNDS_CHECK(check) "copyArray" "shape mismatch" (sh1 == sh2)- $ doCopy =<< gets memoryTable+ $ run doCopy where- doCopy mt = liftIO $ copyR arrayElt adata1 adata2+ doCopy ctx mt = copyR arrayElt adata1 adata2 where copyR :: ArrayEltR e -> ArrayData e -> ArrayData e -> IO () copyR ArrayEltRunit _ _ = return () copyR (ArrayEltRpair aeR1 aeR2) ad1 ad2 = copyR aeR1 (fst ad1) (fst ad2) >> copyR aeR2 (snd ad1) (snd ad2)- copyR aer ad1 ad2 = copyPrim aer mt ad1 ad2 (size sh1)+ copyR aer ad1 ad2 = copyPrim aer ctx mt ad1 ad2 (size sh1) --- copyPrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> ArrayData e -> Int -> IO ()+ copyPrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> ArrayData e -> Int -> IO () mkPrimDispatch(copyPrim,Prim.copyArray) -- Copy data from the device into the associated Accelerate host-side array -- peekArray :: (Shape dim, Elt e) => Array dim e -> CIO ()-peekArray (Array sh adata) = doPeek =<< gets memoryTable+peekArray (Array sh adata) = run doPeek where- doPeek mt = liftIO $ peekR arrayElt adata+ doPeek ctx mt = peekR arrayElt adata where peekR :: ArrayEltR e -> ArrayData e -> IO () peekR ArrayEltRunit _ = return () peekR (ArrayEltRpair aeR1 aeR2) ad = peekR aeR1 (fst ad) >> peekR aeR2 (snd ad)- peekR aer ad = peekPrim aer mt ad (size sh)+ peekR aer ad = peekPrim aer ctx mt ad (size sh) --- peekPrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> Int -> IO ()+ peekPrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> Int -> IO () mkPrimDispatch(peekPrim,Prim.peekArray) peekArrayAsync :: (Shape dim, Elt e) => Array dim e -> Maybe CUDA.Stream -> CIO ()-peekArrayAsync (Array sh adata) ms = doPeek =<< gets memoryTable+peekArrayAsync (Array sh adata) ms = run doPeek where- doPeek mt = liftIO $ peekR arrayElt adata+ doPeek ctx mt = peekR arrayElt adata where peekR :: ArrayEltR e -> ArrayData e -> IO () peekR ArrayEltRunit _ = return () peekR (ArrayEltRpair aeR1 aeR2) ad = peekR aeR1 (fst ad) >> peekR aeR2 (snd ad)- peekR aer ad = peekPrim aer mt ad (size sh) ms+ peekR aer ad = peekPrim aer ctx mt ad (size sh) ms --- peekPrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> Int -> Maybe CUDA.Stream -> IO ()+ peekPrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> Int -> Maybe CUDA.Stream -> IO () mkPrimDispatch(peekPrim,Prim.peekArrayAsync) -- Copy data from an Accelerate array into the associated device array -- pokeArray :: (Shape dim, Elt e) => Array dim e -> CIO ()-pokeArray (Array sh adata) = doPoke =<< gets memoryTable+pokeArray (Array sh adata) = run doPoke where- doPoke mt = liftIO $ pokeR arrayElt adata+ doPoke ctx mt = pokeR arrayElt adata where pokeR :: ArrayEltR e -> ArrayData e -> IO () pokeR ArrayEltRunit _ = return () pokeR (ArrayEltRpair aeR1 aeR2) ad = pokeR aeR1 (fst ad) >> pokeR aeR2 (snd ad)- pokeR aer ad = pokePrim aer mt ad (size sh)+ pokeR aer ad = pokePrim aer ctx mt ad (size sh) --- pokePrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> Int -> IO ()+ pokePrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> Int -> IO () mkPrimDispatch(pokePrim,Prim.pokeArray) pokeArrayAsync :: (Shape dim, Elt e) => Array dim e -> Maybe CUDA.Stream -> CIO ()-pokeArrayAsync (Array sh adata) ms = doPoke =<< gets memoryTable+pokeArrayAsync (Array sh adata) ms = run doPoke where- doPoke mt = liftIO $ pokeR arrayElt adata+ doPoke ctx mt = pokeR arrayElt adata where pokeR :: ArrayEltR e -> ArrayData e -> IO () pokeR ArrayEltRunit _ = return () pokeR (ArrayEltRpair aeR1 aeR2) ad = pokeR aeR1 (fst ad) >> pokeR aeR2 (snd ad)- pokeR aer ad = pokePrim aer mt ad (size sh) ms+ pokeR aer ad = pokePrim aer ctx mt ad (size sh) ms --- pokePrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> Int -> Maybe CUDA.Stream -> IO ()+ pokePrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> Int -> Maybe CUDA.Stream -> IO () mkPrimDispatch(pokePrim,Prim.pokeArrayAsync) @@ -244,17 +269,17 @@ -- that can be passed to a kernel upon invocation. -- marshalArrayData :: ArrayElt e => ArrayData e -> CIO [CUDA.FunParam]-marshalArrayData adata = doMarshal =<< gets memoryTable+marshalArrayData adata = run doMarshal where- doMarshal mt = liftIO $ marshalR arrayElt adata+ doMarshal ctx mt = marshalR arrayElt adata where marshalR :: ArrayEltR e -> ArrayData e -> 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 mt ad+ marshalR aer ad = return <$> marshalPrim aer ctx mt ad --- marshalPrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> IO CUDA.FunParam+ marshalPrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> IO CUDA.FunParam mkPrimDispatch(marshalPrim,Prim.marshalArrayData) @@ -263,9 +288,9 @@ -- consumed, in projection index order --- i.e. right-to-left -- marshalTextureData :: ArrayElt e => ArrayData e -> Int -> [CUDA.Texture] -> CIO ()-marshalTextureData adata n texs = doMarshal =<< gets memoryTable+marshalTextureData adata n texs = run doMarshal where- doMarshal mt = liftIO $ marshalR arrayElt adata texs >> return ()+ doMarshal ctx mt = marshalR arrayElt adata texs >> return () where marshalR :: ArrayEltR e -> ArrayData e -> [CUDA.Texture] -> IO Int marshalR ArrayEltRunit _ _ = return 0@@ -274,27 +299,27 @@ l <- marshalR aeR1 (fst ad) (drop r t) return (l + r) marshalR aer ad t- = do marshalPrim aer mt ad n (head t)+ = do marshalPrim aer ctx mt ad n (head t) return 1 --- marshalPrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> Int -> CUDA.Texture -> IO ()+ marshalPrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> Int -> CUDA.Texture -> IO () mkPrimDispatch(marshalPrim,Prim.marshalTextureData) -- |Raw device pointers associated with a host-side array -- devicePtrsOfArrayData :: ArrayElt e => ArrayData e -> CIO (Prim.DevicePtrs e)-devicePtrsOfArrayData adata = ptrs =<< gets memoryTable+devicePtrsOfArrayData adata = run ptrs where- ptrs mt = liftIO $ ptrsR arrayElt adata+ ptrs ctx mt = ptrsR arrayElt adata where ptrsR :: ArrayEltR e -> ArrayData e -> 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 mt ad+ ptrsR aer ad = ptrsPrim aer ctx mt ad --- ptrsPrim :: ArrayEltR e -> MemoryTable -> ArrayData e -> IO (Prim.DevicePtrs e)+ ptrsPrim :: ArrayEltR e -> Context -> MemoryTable -> ArrayData e -> IO (Prim.DevicePtrs e) mkPrimDispatch(ptrsPrim,Prim.devicePtrsOfArrayData)
Data/Array/Accelerate/CUDA/Array/Prim.hs view
@@ -1,4 +1,8 @@-{-# LANGUAGE BangPatterns, CPP, GADTs, TypeFamilies, ScopedTypeVariables #-}+{-# LANGUAGE BangPatterns #-}+{-# LANGUAGE CPP #-}+{-# LANGUAGE GADTs #-}+{-# LANGUAGE ScopedTypeVariables #-}+{-# LANGUAGE TypeFamilies #-} -- | -- Module : Data.Array.Accelerate.CUDA.Array.Prim -- Copyright : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee@@ -88,14 +92,10 @@ -- CFloat -- CDouble --- FIXME:--- No concrete implementation in Data.Array.Accelerate.Array.Data----type instance HostPtrs Bool = ()-type instance DevicePtrs Bool = ()+type instance HostPtrs Bool = CUDA.HostPtr Word8+type instance DevicePtrs Bool = CUDA.DevicePtr Word8 -type instance HostPtrs Char = ()-type instance DevicePtrs Char = ()+primArrayElt(Char) -- FIXME: -- CChar@@ -126,19 +126,20 @@ instance TextureData Word64 where format _ = (CUDA.Word32, 2) instance TextureData Float where format _ = (CUDA.Float, 1) instance TextureData Double where format _ = (CUDA.Int32, 2)--instance TextureData Int where- format _ = case sizeOf (undefined :: Int) of- 4 -> (CUDA.Int32, 1)- 8 -> (CUDA.Int32, 2)- _ -> error "we can never get here"--instance TextureData Word where- format _ = case sizeOf (undefined :: Word) of- 4 -> (CUDA.Word32, 1)- 8 -> (CUDA.Word32, 2)- _ -> error "we can never get here"-+instance TextureData Bool where format _ = (CUDA.Word8, 1)+#if SIZEOF_HSINT == 4+instance TextureData Int where format _ = (CUDA.Int32, 1)+#elif SIZEOF_HSINT == 8+instance TextureData Int where format _ = (CUDA.Int32, 2)+#endif+#if SIZEOF_HSINT == 4+instance TextureData Word where format _ = (CUDA.Word32, 1)+#elif SIZEOF_HSINT == 8+instance TextureData Word where format _ = (CUDA.Word32, 2)+#endif+#if SIZEOF_HSCHAR == 4+instance TextureData Char where format _ = (CUDA.Word32, 1)+#endif -- Primitive array operations@@ -150,20 +151,21 @@ -- mallocArray :: forall e a. (ArrayElt e, DevicePtrs e ~ CUDA.DevicePtr a, Typeable e, Typeable a, Storable a)- => MemoryTable+ => Context+ -> MemoryTable -> ArrayData e -> Int -> IO ()-mallocArray !mt !ad !n0 = do+mallocArray !ctx !mt !ad !n0 = do let !n = 1 `max` n0- exists <- isJust <$> (lookup mt ad :: IO (Maybe (CUDA.DevicePtr a)))+ exists <- isJust <$> (lookup ctx mt ad :: IO (Maybe (CUDA.DevicePtr a))) unless exists $ do message $ "mallocArray: " ++ showBytes (n * sizeOf (undefined::a)) ptr <- CUDA.mallocArray n `catch` \(e :: CUDAException) -> case e of ExitCode OutOfMemory -> reclaim mt >> CUDA.mallocArray n _ -> throwIO e- insert mt ad (ptr :: CUDA.DevicePtr a)+ insert ctx mt ad (ptr :: CUDA.DevicePtr a) -- A combination of 'mallocArray' and 'pokeArray' to allocate space on the@@ -172,15 +174,16 @@ -- useArray :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr a, Typeable e, Typeable a, Storable a)- => MemoryTable+ => Context+ -> MemoryTable -> ArrayData e -> Int -> IO ()-useArray !mt !ad !n0 =+useArray !ctx !mt !ad !n0 = let src = ptrsOfArrayData ad !n = 1 `max` n0 in do- exists <- isJust <$> (lookup mt ad :: IO (Maybe (CUDA.DevicePtr a)))+ exists <- isJust <$> (lookup ctx mt ad :: IO (Maybe (CUDA.DevicePtr a))) unless exists $ do message $ "useArray/malloc: " ++ showBytes (n * sizeOf (undefined::a)) dst <- CUDA.mallocArray n `catch` \(e :: CUDAException) ->@@ -188,21 +191,22 @@ ExitCode OutOfMemory -> reclaim mt >> CUDA.mallocArray n _ -> throwIO e CUDA.pokeArray n src dst- insert mt ad dst+ insert ctx mt ad dst useArrayAsync :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr a, Typeable e, Typeable a, Storable a)- => MemoryTable+ => Context+ -> MemoryTable -> ArrayData e -> Int -> Maybe CUDA.Stream -> IO ()-useArrayAsync !mt !ad !n0 !ms =+useArrayAsync !ctx !mt !ad !n0 !ms = let src = CUDA.HostPtr (ptrsOfArrayData ad) !n = 1 `max` n0 in do- exists <- isJust <$> (lookup mt ad :: IO (Maybe (CUDA.DevicePtr a)))+ exists <- isJust <$> (lookup ctx mt ad :: IO (Maybe (CUDA.DevicePtr a))) unless exists $ do message $ "useArrayAsync/malloc: " ++ showBytes (n * sizeOf (undefined::a)) dst <- CUDA.mallocArray n `catch` \(e :: CUDAException) ->@@ -210,20 +214,22 @@ ExitCode OutOfMemory -> reclaim mt >> CUDA.mallocArray n _ -> throwIO e CUDA.pokeArrayAsync n src dst ms- insert mt ad dst+ insert ctx mt ad dst -- Read a single element from an array at the given row-major index -- indexArray- :: (ArrayElt e, DevicePtrs e ~ CUDA.DevicePtr b, Typeable e, Typeable b, Storable b)- => MemoryTable+ :: forall e a. (ArrayElt e, DevicePtrs e ~ CUDA.DevicePtr a, Typeable e, Typeable a, Storable a)+ => Context+ -> MemoryTable -> ArrayData e -> Int- -> IO b-indexArray !mt !ad !i =- alloca $ \dst ->- devicePtrsOfArrayData mt ad >>= \src -> do+ -> IO a+indexArray !ctx !mt !ad !i =+ alloca $ \dst ->+ devicePtrsOfArrayData ctx mt ad >>= \src -> do+ message $ "indexArray: " ++ showBytes (sizeOf (undefined::a)) CUDA.peekArray 1 (src `CUDA.advanceDevPtr` i) dst peek dst @@ -232,63 +238,73 @@ -- respect to the host, but will never overlap kernel execution. -- copyArray- :: (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr b, Typeable a, Typeable b, Typeable e, Storable b)- => MemoryTable+ :: forall e a b. (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr b, Typeable a, Typeable b, Typeable e, Storable b)+ => Context+ -> MemoryTable -> ArrayData e -- source array -> ArrayData e -- destination array -> Int -- number of array elements -> IO ()-copyArray !mt !from !to !n = do- src <- devicePtrsOfArrayData mt from- dst <- devicePtrsOfArrayData mt to+copyArray !ctx !mt !from !to !n = do+ message $ "copyArrayAsync: " ++ showBytes (n * sizeOf (undefined :: b))+ src <- devicePtrsOfArrayData ctx mt from+ dst <- devicePtrsOfArrayData ctx mt to CUDA.copyArrayAsync n src dst -- Copy data from the device into the associated Accelerate host-side array -- peekArray- :: (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr a, Typeable a, Typeable e, Storable a)- => MemoryTable+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr a, Typeable a, Typeable e, Storable a)+ => Context+ -> MemoryTable -> ArrayData e -> Int -> IO ()-peekArray !mt !ad !n =- devicePtrsOfArrayData mt ad >>= \src ->+peekArray !ctx !mt !ad !n =+ devicePtrsOfArrayData ctx mt ad >>= \src -> do+ message $ "peekArray: " ++ showBytes (n * sizeOf (undefined :: a)) CUDA.peekArray n src (ptrsOfArrayData ad) peekArrayAsync- :: (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr a, Typeable a, Typeable e, Storable a)- => MemoryTable+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr a, Typeable a, Typeable e, Storable a)+ => Context+ -> MemoryTable -> ArrayData e -> Int -> Maybe CUDA.Stream -> IO ()-peekArrayAsync !mt !ad !n !st =- devicePtrsOfArrayData mt ad >>= \src ->+peekArrayAsync !ctx !mt !ad !n !st =+ devicePtrsOfArrayData ctx mt ad >>= \src -> do+ message $ "peekArrayAsync: " ++ showBytes (n * sizeOf (undefined :: a)) CUDA.peekArrayAsync n src (CUDA.HostPtr $ ptrsOfArrayData ad) st -- Copy data from an Accelerate array into the associated device array -- pokeArray- :: (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr a, Typeable a, Typeable e, Storable a)- => MemoryTable+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr a, Typeable a, Typeable e, Storable a)+ => Context+ -> MemoryTable -> ArrayData e -> Int -> IO ()-pokeArray !mt !ad !n =- devicePtrsOfArrayData mt ad >>= \dst ->+pokeArray !ctx !mt !ad !n =+ devicePtrsOfArrayData ctx mt ad >>= \dst -> do+ message $ "pokeArrayAsync: " ++ showBytes (n * sizeOf (undefined :: a)) CUDA.pokeArray n (ptrsOfArrayData ad) dst pokeArrayAsync- :: (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr a, Typeable a, Typeable e, Storable a)- => MemoryTable+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, DevicePtrs e ~ CUDA.DevicePtr a, Typeable a, Typeable e, Storable a)+ => Context+ -> MemoryTable -> ArrayData e -> Int -> Maybe CUDA.Stream -> IO ()-pokeArrayAsync !mt !ad !n !st =- devicePtrsOfArrayData mt ad >>= \dst ->+pokeArrayAsync !ctx !mt !ad !n !st =+ devicePtrsOfArrayData ctx mt ad >>= \dst -> do+ message $ "pokeArrayAsync: " ++ showBytes (n * sizeOf (undefined :: a)) CUDA.pokeArrayAsync n (CUDA.HostPtr $ ptrsOfArrayData ad) dst st @@ -307,24 +323,26 @@ -- marshalArrayData :: (ArrayElt e, DevicePtrs e ~ CUDA.DevicePtr b, Typeable b, Typeable e)- => MemoryTable+ => Context+ -> MemoryTable -> ArrayData e -> IO CUDA.FunParam-marshalArrayData !mt !ad = marshalDevicePtrs ad <$> devicePtrsOfArrayData mt ad+marshalArrayData !ctx !mt !ad = marshalDevicePtrs ad <$> devicePtrsOfArrayData ctx mt ad -- Bind device memory to the given texture reference, setting appropriate type -- marshalTextureData :: forall a e. (ArrayElt e, DevicePtrs e ~ CUDA.DevicePtr a, Typeable a, Typeable e, Storable a, TextureData a)- => MemoryTable+ => Context+ -> MemoryTable -> ArrayData e -- host array -> Int -- number of elements -> CUDA.Texture -- texture reference to bind array to -> IO ()-marshalTextureData !mt !ad !n !tex =+marshalTextureData !ctx !mt !ad !n !tex = let (fmt, c) = format (undefined :: a)- in devicePtrsOfArrayData mt ad >>= \ptr -> do+ in devicePtrsOfArrayData ctx mt ad >>= \ptr -> do CUDA.setFormat tex fmt c CUDA.bind tex ptr (fromIntegral $ n * sizeOf (undefined :: a)) @@ -333,11 +351,12 @@ -- devicePtrsOfArrayData :: (ArrayElt e, DevicePtrs e ~ CUDA.DevicePtr b, Typeable e, Typeable b)- => MemoryTable+ => Context+ -> MemoryTable -> ArrayData e -> IO (DevicePtrs e)-devicePtrsOfArrayData !mt !ad = do- mv <- lookup mt ad+devicePtrsOfArrayData !ctx !mt !ad = do+ mv <- lookup ctx mt ad case mv of Just v -> return v Nothing -> do
Data/Array/Accelerate/CUDA/Array/Table.hs view
@@ -1,4 +1,7 @@-{-# LANGUAGE BangPatterns, CPP, GADTs, PatternGuards #-}+{-# LANGUAGE BangPatterns #-}+{-# LANGUAGE CPP #-}+{-# LANGUAGE GADTs #-}+{-# LANGUAGE PatternGuards #-} -- | -- Module : Data.Array.Accelerate.CUDA.Array.Table -- Copyright : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee@@ -13,7 +16,7 @@ module Data.Array.Accelerate.CUDA.Array.Table ( -- Tables for host/device memory associations- MemoryTable, new, lookup, insert, reclaim+ MemoryTable, Context(..), new, lookup, insert, reclaim ) where @@ -24,7 +27,7 @@ import Data.Typeable ( Typeable, gcast ) import Control.Monad ( unless ) import Control.Exception ( bracket_ )-import Control.Applicative ( (<$>), (<*>) )+import Control.Applicative ( (<$>) ) import System.Mem ( performGC ) import System.Mem.Weak ( Weak, mkWeak, deRefWeak, finalize ) import System.Mem.StableName ( StableName, makeStableName, hashStableName )@@ -35,7 +38,7 @@ import qualified Data.HashTable.IO as HT import Data.Array.Accelerate.Array.Data ( ArrayData )-import qualified Data.Array.Accelerate.CUDA.Debug as D ( message, dump_gc )+import qualified Data.Array.Accelerate.CUDA.Debug as D #include "accelerate.h" @@ -53,15 +56,22 @@ -- 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 = IORef ( HashTable HostArray DeviceArray )-data MemoryTable = MemoryTable {-# UNPACK #-} !MT- {-# UNPACK #-} !(Weak MT)+type HashTable key val = HT.BasicHashTable key val+type MT = IORef ( HashTable HostArray DeviceArray )+data MemoryTable = MemoryTable {-# UNPACK #-} !MT+ {-# UNPACK #-} !(Weak MT) +-- The currently active context. Finaliser threads need to check if the context+-- is still active before attempting to release their associated memory.+--+data Context = Context {-# UNPACK #-} !CUDA.Context+ {-# UNPACK #-} !(Weak CUDA.Context) +-- Arrays on the host and device+-- data HostArray where HostArray :: Typeable e- => {-# UNPACK #-} !CUDA.Context+ => {-# UNPACK #-} !Int -- unique ID relating to the parent context -> {-# UNPACK #-} !(StableName (ArrayData e)) -> HostArray @@ -75,8 +85,7 @@ = maybe False (== a2) (gcast a1) instance Hashable HostArray where- hash (HostArray (CUDA.Context p) sn) =- fromIntegral (ptrToIntPtr p) `hashWithSalt` sn+ hash (HostArray cid sn) = hashWithSalt cid sn instance Show HostArray where show (HostArray _ sn) = "Array #" ++ show (hashStableName sn)@@ -98,28 +107,28 @@ -- Look for the device memory corresponding to a given host-side array. ---lookup :: (Typeable a, Typeable b) => MemoryTable -> ArrayData a -> IO (Maybe (DevicePtr b))-lookup (MemoryTable ref _) !arr = do- sa <- makeStableArray arr+lookup :: (Typeable a, Typeable b) => Context -> MemoryTable -> ArrayData a -> IO (Maybe (DevicePtr b))+lookup ctx (MemoryTable ref _) !arr = do+ sa <- makeStableArray ctx arr mw <- withIORef 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 -> INTERNAL_ERROR(error) "lookup" $ "type mismatch"- Nothing ->- makeStableArray arr >>= \x -> INTERNAL_ERROR(error) "lookup" $ "dead weak pair: " ++ show x+ Just v | Just p <- gcast v -> trace ("lookup/found: " ++ show sa) $ return (Just p)+ | otherwise -> INTERNAL_ERROR(error) "lookup" $ "type mismatch"+ Nothing ->+ makeStableArray ctx arr >>= \x -> INTERNAL_ERROR(error) "lookup" $ "dead weak pair: " ++ show x -- 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) => MemoryTable -> ArrayData a -> DevicePtr b -> IO ()-insert (MemoryTable ref weak_ref) !arr !ptr = do- key <- makeStableArray arr- dev <- DeviceArray `fmap` mkWeak arr ptr (Just $ finalizer weak_ref key ptr)+insert :: (Typeable a, Typeable b) => Context -> MemoryTable -> ArrayData a -> DevicePtr b -> IO ()+insert ctx@(Context _ weak_ctx) (MemoryTable ref weak_ref) !arr !ptr = do+ key <- makeStableArray ctx arr+ dev <- DeviceArray `fmap` mkWeak arr ptr (Just $ finalizer weak_ctx weak_ref key ptr) tbl <- readIORef ref message $ "insert: " ++ show key HT.insert tbl key dev@@ -133,7 +142,8 @@ -- reclaim :: MemoryTable -> IO () reclaim (MemoryTable _ weak_ref) = do- trace "reclaim" performGC+ (free, total) <- CUDA.getMemInfo+ performGC mr <- deRefWeak weak_ref case mr of Nothing -> return ()@@ -141,7 +151,12 @@ flip HT.mapM_ tbl $ \(_,DeviceArray w) -> do alive <- isJust `fmap` deRefWeak w unless alive $ finalize w-+ --+ D.when D.dump_gc $ do+ (free', _) <- CUDA.getMemInfo+ message $ "reclaim: freed " ++ showBytes (fromIntegral (free - free'))+ ++ ", " ++ showBytes (fromIntegral free')+ ++ " 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.@@ -152,17 +167,17 @@ -- the hash tables --- but we must do this first before failing to use a dead -- context. ---finalizer :: Weak MT -> HostArray -> DevicePtr b -> IO ()-finalizer !weak_ref !key@(HostArray ctx _) !ptr = do+finalizer :: Weak CUDA.Context -> Weak MT -> HostArray -> DevicePtr b -> IO ()+finalizer !weak_ctx !weak_ref !key !ptr = do mr <- deRefWeak weak_ref case mr of- Nothing -> trace ("finalise/dead table: " ++ show key) $ return ()- Just ref -> trace ("finalise: " ++ show key) $ withIORef ref (`HT.delete` key)+ Nothing -> message ("finalise/dead table: " ++ show key)+ Just ref -> trace ("finalise: " ++ show key) $ withIORef ref (`HT.delete` key) --- bracket_- (CUDA.push ctx)- (CUDA.pop)- (CUDA.free ptr)+ mc <- deRefWeak weak_ctx+ case mc of+ Nothing -> message ("finalise/dead context: " ++ show key)+ Just ctx -> bracket_ (CUDA.push ctx) CUDA.pop (CUDA.free ptr) table_finalizer :: HashTable HostArray DeviceArray -> IO ()@@ -175,8 +190,10 @@ -- ------------- {-# INLINE makeStableArray #-}-makeStableArray :: Typeable a => ArrayData a -> IO HostArray-makeStableArray !arr = HostArray <$> CUDA.get <*> makeStableName arr+makeStableArray :: Typeable a => Context -> ArrayData a -> IO HostArray+makeStableArray (Context (CUDA.Context !p) !_) !arr =+ let cid = fromIntegral (ptrToIntPtr p)+ in HostArray cid <$> makeStableName arr {-# INLINE withIORef #-} withIORef :: IORef a -> (a -> IO b) -> IO b@@ -185,6 +202,10 @@ -- 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
Data/Array/Accelerate/CUDA/CodeGen.hs view
@@ -184,7 +184,7 @@ idx' = show $ idxToInt idx sh = cshape ("sh" ++ idx') (accDim avar) ty = accTypeTex avar- arr n = "arr" ++ idx' ++ "_a" ++ show (n::Int)+ arr n = "avar" ++ idx' ++ "_a" ++ show (n::Int) in sh : zipWith (\t n -> cglobal t (arr n)) (reverse ty) [0..] @@ -322,7 +322,9 @@ [x] -> bind [cty| typename bool |] x _ -> INTERNAL_ERROR(error) "codegenOpenExp" "expected conditional predicate" --- return $ zipWith (\a b -> [cexp| $exp:p' ? $exp:a : $exp:b|]) t' e'+ let cond ty a b = addVar ty a >> addVar ty b >>+ return [cexp| $exp:p' ? $exp:a : $exp:b|]+ sequence $ zipWith3 cond (expType t) t' e' -- Array indices and shapes --@@ -361,8 +363,8 @@ IndexScalar arr ix | OpenAcc (Avar a) <- arr -> let avar = show (idxToInt a)- sh = cvar ("sh" ++ avar)- array x = cvar ("arr" ++ avar ++ "_a" ++ show x)+ sh = cvar ("sh" ++ avar)+ array x = cvar ("avar" ++ avar ++ "_a" ++ show x) elt = accTypeTex arr n = length elt in do@@ -408,6 +410,8 @@ addVar ty exp = case show exp of ('x':v:'_':'a':n) | [(v',[])] <- reads [v], [(n',[])] <- reads n -> use v' n' ty exp >> return True+ ('v':n) | [(_ :: Int,[])] <- reads n+ -> return True _ -> return False
Data/Array/Accelerate/CUDA/CodeGen/Base.hs view
@@ -116,8 +116,8 @@ , String -> [InitGroup] ) -- const declarations and initialisation from index getters base arrElt expElt = let n = length arrElt- arr x = "d_in" ++ shows base "_a" ++ show x arrParams = zipWith (\t x -> [cparam| const $ty:(cptr t) $id:(arr x) |]) arrElt [n-1, n-2 .. 0]+ arr x = "arrIn" ++ shows base "_a" ++ show x expVars = map (\(_,_,v) -> v) expElt expDecls = map (\(_,t,v) -> [cdecl| $ty:t $id:(show v) ; |]) expElt in@@ -139,7 +139,7 @@ , String -> [Exp] -> [Stm]) -- store a value to the given index setters arrElt = let n = length arrElt- arrVars = map (\x -> "d_out_a" ++ show x) [n-1, n-2 .. 0]+ arrVars = map (\x -> "arrOut_a" ++ show x) [n-1, n-2 .. 0] arrParams = zipWith (\t x -> [cparam| $ty:(cptr t) $id:x |]) arrElt arrVars set ix a x = [cstm| $id:a [$id:ix] = $exp:x; |] in
Data/Array/Accelerate/CUDA/CodeGen/IndexSpace.hs view
@@ -379,6 +379,7 @@ -- project :: Int -> String -> [Exp] -> [Stm] project n sh idx+ | n == 0 = [[cstm| $id:sh = 0; |]] | [e] <- idx = [[cstm| $id:sh = $exp:e; |]] | otherwise = zipWith (\i c -> [cstm| $id:sh . $id:('a':show c) = $exp:i; |]) idx [n-1,n-2..0]
Data/Array/Accelerate/CUDA/CodeGen/Monad.hs view
@@ -79,10 +79,15 @@ use :: Int -> Int -> Type -> Exp -> CGM () use base prj ty var = modify variables (S.adjust (IM.insert prj (ty,var)) base) --- Return the tuple components of a given variable that are actually used+-- Return the tuple components of a given variable that are actually used. These+-- in snoc-list ordering, i.e. with variable zero on the right. -- subscripts :: Int -> CGM [(Int, Type, Exp)]-subscripts base = map swizzle . IM.toList . flip S.index base <$> gets variables+subscripts base+ = reverse+ . map swizzle+ . IM.toList+ . flip S.index base <$> gets variables where swizzle (i, (t,e)) = (i,t,e)
Data/Array/Accelerate/CUDA/CodeGen/PrefixSum.hs view
@@ -134,7 +134,7 @@ $decls:decl2 /*- * Read in previous result partial sum. We store the carry value in x0+ * Read in previous result partial sum. We store the carry value in x2 * and read new values from the input array into x1, since 'scanBlock' * will store its results into x1 on completion. */@@ -154,19 +154,26 @@ $stms:(x1 .=. getIn0 "j") if ( $exp:carry_in ) {- $decls:env $stms:(x0 .=. x2)+ $decls:env $stms:(x1 .=. combine) } /*- * Store our input into shared memory and cooperatively scan+ * Store our input into shared memory and perform a cooperative+ * inclusive left scan. */ $stms:(sdata "threadIdx.x" .=. x1) __syncthreads(); $stms:(scanBlock dev elt Nothing (cvar "blockDim.x") sdata env combine) + /*+ * Exclusive scans write the result of the previous thread to global+ * memory. The first thread must reinstate the carry-in value which+ * is the result of the last thread from the previous interval, or+ * the carry-in/seed value for multi-block scans.+ */ if ( $exp:(cbool exclusive) ) { if ( threadIdx.x == 0 ) { $stms:(x1 .=. x2)@@ -177,9 +184,9 @@ $stms:(setOut "j" x1) /*- * Carry the final result from this block through x0. If this is the- * last section of the interval, this is the value to write out as- * the final (reduction) result.+ * Carry the final result of this block through the set x2. If this+ * is the final interval, this is the value to write out as the+ * reduction result */ if ( threadIdx.x == 0 ) { const int last = min(interval_size - i, blockDim.x) - 1;@@ -189,8 +196,7 @@ } /*- * for exclusive scans, set the overall scan result and reapply the- * initial element at the boundaries of each interval+ * for exclusive scans, set the overall scan result (reduction value) */ if ( $exp:(cbool exclusive) && threadIdx.x == 0 && blockIdx.x == $id:lastBlock ) { $stms:(setSum .=. x2)
Data/Array/Accelerate/CUDA/CodeGen/Reduction.hs view
@@ -84,8 +84,8 @@ for (i += gridSize; i < num_elements; i += gridSize) {- $decls:env $stms:(x0 .=. getIn0 "i")+ $decls:env $stms:(x1 .=. combine) } }@@ -124,6 +124,7 @@ if (gridDim.x == 1) { $decls:env' $stms:(x0 .=. seed)+ $decls:env $stms:(x1 .=. combine) } $stms:(setOut "blockIdx.x" x1)@@ -171,7 +172,6 @@ $decls:smem $decls:decl1 $decls:decl0- $decls:env /* * If the intervals of an exclusive fold are empty, use all threads to@@ -228,8 +228,8 @@ */ for (i += 2 * blockDim.x; i < end; i += blockDim.x) {- $decls:env $stms:(x0 .=. getIn0 "i")+ $decls:env $stms:(x1 .=. combine) } }@@ -252,11 +252,7 @@ * exclusive reductions, we also combine with the seed element here. */ if (threadIdx.x == 0)- {- $decls:final_decls- $stms:final_stms- $stms:(setOut "seg" x1)- }+ $stm:(maybe inclusive_finish exclusive_finish mseed) } } |]@@ -268,10 +264,16 @@ (x1, decl1) = locals "x1" elt (smem, sdata) = shared 0 Nothing [cexp| blockDim.x |] elt --- (final_decls, final_stms) =- case mseed of- Nothing -> ([], [])- Just (CUExp env' seed) -> (env', concat [x0 .=. seed, x1 .=. combine])+ inclusive_finish = [cstm| {+ $stms:(setOut "seg" x1)+ } |]+ exclusive_finish (CUExp env' seed) = [cstm| {+ $decls:env'+ $stms:(x0 .=. seed)+ $decls:env+ $stms:(x1 .=. combine)+ $stms:(setOut "seg" x1)+ } |] -- mapseed (CUExp env' seed) = [cstm| if (interval_size == 0)@@ -411,8 +413,8 @@ */ for (i += 2 * warpSize; i < end; i += warpSize) {- $decls:env $stms:(x0 .=. getIn0 "i")+ $decls:env $stms:(x1 .=. combine) } }@@ -451,6 +453,7 @@ if (num_elements > 0) { $decls:env' $stms:(x0 .=. seed)+ $decls:env $stms:(x1 .=. combine) } else { $decls:env'@@ -486,8 +489,8 @@ reduce i | i > 1 = [cstm| if ( $id:tid + $int:i < $id:n ) {- $decls:env $stms:(x0 .=. sdata ("threadIdx.x + " ++ show i))+ $decls:env $stms:(x1 .=. combine) $stms:(sdata "threadIdx.x" .=. x1) }@@ -495,8 +498,8 @@ -- | otherwise = [cstm| if ( $id:tid + $int:i < $id:n ) {- $decls:env $stms:(x0 .=. sdata "threadIdx.x + 1")+ $decls:env $stms:(x1 .=. combine) } |]@@ -525,8 +528,8 @@ | i > warpSize dev = [cstm| if ( $id:n > $int:i ) { if ( threadIdx.x + $int:i < $id:n ) {- $decls:env $stms:(x0 .=. sdata ("threadIdx.x + " ++ show i))+ $decls:env $stms:(x1 .=. combine) $stms:(sdata "threadIdx.x" .=. x1) }
Data/Array/Accelerate/CUDA/CodeGen/Stencil.hs view
@@ -184,8 +184,8 @@ , \ix -> concatMap (get ix) subs ) where n = length stencil- sh = "shIn" ++ show base- arr x = "stencil" ++ shows base "_a" ++ show (x `mod` n)+ sh = "shIn" ++ show base+ arr x = "arrIn" ++ shows base "_a" ++ show (x `mod` n) textures = zipWith cglobal stencil (map arr [n-1, n-2 .. 0]) -- offsets :: IArray.Array Int [Int]
Data/Array/Accelerate/CUDA/CodeGen/Type.hs view
@@ -35,7 +35,6 @@ -- libraries import Language.C.Quote.CUDA import qualified Language.C as C-import qualified Foreign.Storable as F #include "accelerate.h"@@ -56,14 +55,6 @@ | otherwise = INTERNAL_ERROR(error) "accType" "non-scalar segment type" ---sizeOfAccTypes :: OpenAcc aenv (Sugar.Array dim e) -> [Int]---sizeOfAccTypes = sizeOf' . Sugar.accType--- where--- sizeOf' :: TupleType a -> [Int]--- sizeOf' UnitTuple = []--- sizeOf' x@(SingleTuple _) = [Sugar.sizeOf x]--- sizeOf' (PairTuple a b) = sizeOf' a ++ sizeOf' b- eltType :: Sugar.Elt a => a {- dummy -} -> [C.Type] eltType = codegenTupleType . Sugar.eltType @@ -95,14 +86,14 @@ codegenNumType (FloatingNumType ty) = codegenFloatingType ty codegenIntegralType :: IntegralType a -> C.Type-codegenIntegralType (TypeInt8 _) = typename "int8_t"-codegenIntegralType (TypeInt16 _) = typename "int16_t"-codegenIntegralType (TypeInt32 _) = typename "int32_t"-codegenIntegralType (TypeInt64 _) = typename "int64_t"-codegenIntegralType (TypeWord8 _) = typename "uint8_t"-codegenIntegralType (TypeWord16 _) = typename "uint16_t"-codegenIntegralType (TypeWord32 _) = typename "uint32_t"-codegenIntegralType (TypeWord64 _) = typename "uint64_t"+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|]@@ -111,19 +102,16 @@ codegenIntegralType (TypeCULong _) = [cty|unsigned long int|] codegenIntegralType (TypeCLLong _) = [cty|long long int|] codegenIntegralType (TypeCULLong _) = [cty|unsigned long long int|]--codegenIntegralType (TypeInt _) =- case F.sizeOf (undefined::Int) of- 4 -> typename "int32_t"- 8 -> typename "int64_t"- _ -> error "we can never get here"--codegenIntegralType (TypeWord _) =- case F.sizeOf (undefined::Int) of- 4 -> typename "uint32_t"- 8 -> typename "uint64_t"- _ -> error "we can never get here"-+#if SIZEOF_HSINT == 4+codegenIntegralType (TypeInt _) = typename "Int32"+#elif SIZEOF_HSINT == 8+codegenIntegralType (TypeInt _) = typename "Int64"+#endif+#if SIZEOF_HSINT == 4+codegenIntegralType (TypeWord _) = typename "Word32"+#elif SIZEOF_HSINT == 8+codegenIntegralType (TypeWord _) = typename "Word64"+#endif codegenFloatingType :: FloatingType a -> C.Type codegenFloatingType (TypeFloat _) = [cty|float|]@@ -132,8 +120,10 @@ codegenFloatingType (TypeCDouble _) = [cty|double|] codegenNonNumType :: NonNumType a -> C.Type-codegenNonNumType (TypeBool _) = error "codegenNonNum :: Bool"-codegenNonNumType (TypeChar _) = error "codegenNonNum :: Char"+codegenNonNumType (TypeBool _) = typename "Word8"+#if SIZEOF_HSCHAR == 4+codegenNonNumType (TypeChar _) = typename "Word32"+#endif codegenNonNumType (TypeCChar _) = [cty|char|] codegenNonNumType (TypeCSChar _) = [cty|signed char|] codegenNonNumType (TypeCUChar _) = [cty|unsigned char|]@@ -178,18 +168,16 @@ codegenIntegralTex (TypeCULong _) = typename "TexCULong" codegenIntegralTex (TypeCLLong _) = typename "TexCLLong" codegenIntegralTex (TypeCULLong _) = typename "TexCULLong"--codegenIntegralTex (TypeInt _) =- case F.sizeOf (undefined::Int) of- 4 -> typename "TexInt32"- 8 -> typename "TexInt64"- _ -> error "we can never get here"--codegenIntegralTex (TypeWord _) =- case F.sizeOf (undefined::Word) of- 4 -> typename "TexWord32"- 8 -> typename "TexWord64"- _ -> error "we can never get here"+#if SIZEOF_HSINT == 4+codegenIntegralTex (TypeInt _) = typename "TexInt32"+#elif SIZEOF_HSINT == 8+codegenIntegralTex (TypeInt _) = typename "TexInt64"+#endif+#if SIZEOF_HSINT == 4+codegenIntegralTex (TypeWord _) = typename "TexWord32"+#elif SIZEOF_HSINT == 8+codegenIntegralTex (TypeWord _) = typename "TexWord64"+#endif codegenFloatingTex :: FloatingType a -> C.Type@@ -199,13 +187,11 @@ codegenFloatingTex (TypeCDouble _) = typename "TexCDouble" --- TLM 2010-06-29:--- Bool and Char can be implemented once the array types in--- Data.Array.Accelerate.[CUDA.]Array.Data are made concrete.--- codegenNonNumTex :: NonNumType a -> C.Type-codegenNonNumTex (TypeBool _) = error "codegenNonNumTex :: Bool"-codegenNonNumTex (TypeChar _) = error "codegenNonNumTex :: Char"+codegenNonNumTex (TypeBool _) = typename "TexWord8"+#if SIZEOF_HSCHAR == 4+codegenNonNumTex (TypeChar _) = typename "TexWord32"+#endif codegenNonNumTex (TypeCChar _) = typename "TexCChar" codegenNonNumTex (TypeCSChar _) = typename "TexCSChar" codegenNonNumTex (TypeCUChar _) = typename "TexCUChar"
Data/Array/Accelerate/CUDA/Compile.hs view
@@ -24,11 +24,12 @@ import Data.Array.Accelerate.Tuple import Data.Array.Accelerate.CUDA.AST-import Data.Array.Accelerate.CUDA.FullList as FL import Data.Array.Accelerate.CUDA.State import Data.Array.Accelerate.CUDA.CodeGen import Data.Array.Accelerate.CUDA.Array.Sugar import Data.Array.Accelerate.CUDA.Analysis.Launch+import Data.Array.Accelerate.CUDA.FullList as FL+import Data.Array.Accelerate.CUDA.Persistent as KT import qualified Data.Array.Accelerate.CUDA.Debug as D -- libraries@@ -45,7 +46,6 @@ import Data.List import Data.Maybe import Data.Monoid-import Foreign.Storable import System.Directory import System.Exit ( ExitCode(..) ) import System.FilePath@@ -55,7 +55,6 @@ import Text.PrettyPrint.Mainland ( RDoc(..), ppr, renderCompact ) import Data.ByteString.Internal ( w2c ) import qualified Data.HashSet as Set-import qualified Data.HashTable.IO as HT import qualified Data.ByteString as B import qualified Data.ByteString.Lazy as L import qualified Foreign.CUDA.Driver as CUDA@@ -300,53 +299,57 @@ -- -- make sure kernel/stats are printed together --- message $ intercalate "\n" [msg1, "cc: " ++ 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 successfully, the temporary files are removed.+-- complete. If successful, the temporary files are removed. -- link :: KernelTable -> KernelKey -> IO CUDA.Module link table key = let intErr = INTERNAL_ERROR(error) "link" "missing kernel entry" in do- ctx <- CUDA.get- (KernelEntry cufile stat) <- fromMaybe intErr `fmap` HT.lookup table key- case stat of- Right (KernelObject bin active)- | Just mdl <- FL.lookup ctx active -> return mdl- | otherwise -> do- message "re-linking module for current context"- mdl <- CUDA.loadData bin- let obj = KernelObject bin (FL.cons ctx mdl active)- HT.insert table key (KernelEntry cufile (Right obj))- return mdl- --- Left pid -> do- -- wait for compiler to finish and load binary object+ ctx <- CUDA.get+ entry <- fromMaybe intErr `fmap` KT.lookup table key+ case entry of+ CompileProcess cufile pid -> do+ -- Wait for the compiler to finish and load the binary object into the+ -- current context -- message "waiting for nvcc..." waitFor pid- bin <- B.readFile (replaceExtension cufile ".cubin")- mdl <- CUDA.loadData bin- let obj = KernelObject bin (FL.singleton ctx mdl)+ let cubin = replaceExtension cufile ".cubin"+ bin <- B.readFile cubin+ mdl <- CUDA.loadData bin -#ifndef ACCELERATE_CUDA_PERSISTENT_CACHE- -- remove build products+ -- Update hash tables and stash the binary object into the persistent+ -- cache --+ KT.insert table key $! KernelObject bin (FL.singleton ctx mdl)+ KT.persist cubin key++ -- Remove temporary build products+ -- removeFile cufile- removeFile (replaceExtension cufile ".cubin") removeDirectory (dropFileName cufile) `catch` \(_ :: IOError) -> return () -- directory not empty-#endif - -- update hash table- --- HT.insert table key (KernelEntry cufile (Right obj)) return mdl + -- 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+ | otherwise -> do+ message "re-linking module for current context"+ mdl <- CUDA.loadData bin+ KT.insert table key $! KernelObject bin (FL.cons ctx mdl active)+ return mdl + -- Generate and compile code for a single open array expression -- compile :: KernelTable@@ -355,17 +358,17 @@ -> AccBindings aenv -> CIO (String, KernelKey) compile table dev acc fvar = do- exists <- isJust `fmap` liftIO (HT.lookup table key)+ exists <- isJust `fmap` liftIO (KT.lookup table key) unless exists $ do message $ unlines [ show key, map w2c (L.unpack code) ] nvcc <- fromMaybe (error "nvcc: command not found") <$> liftIO (findExecutable "nvcc")- (file,hdl) <- openOutputFile "dragon.cu" -- rawr!+ (file,hdl) <- openTemporaryFile "dragon.cu" -- rawr! flags <- compileFlags file (_,_,_,pid) <- liftIO $ do L.hPut hdl code `finally` hClose hdl createProcess (proc nvcc flags) `onException` removeFile file --- liftIO $ HT.insert table key (KernelEntry file (Left pid))+ liftIO $ KT.insert table key (CompileProcess file pid) -- return (entry, key) where@@ -405,28 +408,25 @@ , "-arch=sm_" ++ show (round (arch * 10) :: Int) , "-cubin" , "-o", cufile `replaceExtension` "cubin"- , if D.mode D.verbose then "" else "--disable-warnings"- , if D.mode D.debug then "-G" else "-O3"+ , if D.mode D.dump_cc then "" else "--disable-warnings"+ , if D.mode D.debug_cc then "-G" else "-O3" , machine , cufile ] where- wordSize = sizeOf (undefined::Int)- machine | wordSize == 4 = "-m32"- | wordSize == 8 = "-m64"- | otherwise = error "recreational scolding?"+#if SIZEOF_HSINT == 4+ machine = "-m32"+#elif SIZEOF_HSINT == 8+ machine = "-m64"+#endif -- Open a unique file in the temporary directory used for compilation -- by-products. The directory will be created if it does not exist. ---openOutputFile :: String -> CIO (FilePath, Handle)-openOutputFile template = liftIO $ do-#ifdef ACCELERATE_CUDA_PERSISTENT_CACHE- dir <- (</>) <$> getDataDir <*> pure "cache"-#else+openTemporaryFile :: String -> CIO (FilePath, Handle)+openTemporaryFile template = liftIO $ do pid <- getProcessID dir <- (</>) <$> getTemporaryDirectory <*> pure ("accelerate-cuda-" ++ show pid)-#endif createDirectoryIfMissing True dir openTempFile dir template @@ -440,9 +440,9 @@ {-# INLINE message #-} message :: MonadIO m => String -> m ()-message msg = trace ("cc: " ++ msg) $ return ()+message msg = trace msg $ return () {-# INLINE trace #-} trace :: MonadIO m => String -> m a -> m a-trace msg next = D.message D.dump_cc msg >> next+trace msg next = D.message D.dump_cc ("cc: " ++ msg) >> next
Data/Array/Accelerate/CUDA/Debug.hs view
@@ -20,8 +20,8 @@ showFFloatSIBase, message, event, when, mode,- verbose, debug,- dump_gc, dump_cc, dump_exec,+ verbose, flush_cache,+ dump_gc, dump_cc, debug_cc, dump_exec, ) where @@ -30,6 +30,7 @@ import Data.Label import Data.IORef import Control.Monad.IO.Class+import System.CPUTime import System.IO.Unsafe import System.Environment import System.Console.GetOpt@@ -67,13 +68,14 @@ data Flags = Flags { -- phase control- _dump_gc :: !Bool -- garbage collection & memory management- , _dump_cc :: !Bool -- compilation & linking- , _dump_exec :: !Bool -- kernel execution+ _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 -- general options- , _verbose :: !Bool -- additional status messages- , _debug :: !Bool -- generate device code suitable for debugging+ , _verbose :: !Bool -- additional status messages+ , _flush_cache :: !Bool -- delete the persistent cache directory } $(mkLabels [''Flags])@@ -82,15 +84,16 @@ flags = [ 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"- , Option [] ["ddebug"] (NoArg (set debug True)) "generate debug information for device code"+ , Option [] ["fflush-cache"] (NoArg (set flush_cache True)) "delete the persistent cache directory" ] initialise :: IO Flags initialise = parse `fmap` getArgs where- defaults = Flags False False False False False+ defaults = Flags False False False False False False parse = foldl parse1 defaults parse1 opts x = case filter (\(Option _ [f] _ _) -> x `isPrefixOf` ('-':f)) flags of [Option _ _ (NoArg go) _] -> go opts@@ -113,7 +116,11 @@ {-# INLINE message #-} message :: MonadIO m => (Flags :-> Bool) -> String -> m () #ifdef ACCELERATE_DEBUG-message f str = when f (liftIO $ traceIO str)+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
Data/Array/Accelerate/CUDA/Execute.hs view
@@ -641,9 +641,9 @@ -> Val aenv -> Array dim a -> CIO (Array dim b)-stencilOp kernel@(Kernel _ mdl _ _ _) bindings aenv in0@(Array sh0 _) = do+stencilOp kernel bindings aenv in0@(Array sh0 _) = do res@(Array _ out) <- allocateArray (toElt sh0)- bindStencil 0 mdl in0+ bindAcc 0 kernel in0 execute kernel bindings aenv (size sh0) (((), out) , toElt sh0 :: dim)@@ -657,10 +657,10 @@ -> Array dim a -> Array dim b -> CIO (Array dim c)-stencil2Op kernel@(Kernel _ mdl _ _ _) bindings aenv in1@(Array sh1 _) in0@(Array sh0 _) = do+stencil2Op kernel bindings aenv in1@(Array sh1 _) in0@(Array sh0 _) = do res@(Array sh out) <- allocateArray $ toElt (sh1 `intersect` sh0)- bindStencil 1 mdl in1- bindStencil 0 mdl in0+ bindAcc 1 kernel in1+ bindAcc 0 kernel in0 execute kernel bindings aenv (size sh) (((((), out) , toElt sh :: dim)@@ -756,40 +756,42 @@ -- Array references in scalar code -- ------------------------------- -bindLifted :: CUDA.Module -> Val aenv -> AccBindings aenv -> CIO ()-bindLifted mdl aenv (AccBindings vars) = mapM_ (bindAcc mdl aenv) (Set.toList vars)-+-- All CUDA devices have between 6-8KB of read-only texture memory per+-- multiprocessor. Since all arrays in Accelerate are immutable, we can always+-- access input arrays through the texture cache to reduce global memory demand+-- when accesses do not follow the regular patterns required for coalescing.+--+-- This is great for older 1.x series devices, but compute 2.x devices have a+-- dedicated 768KB L2 cache, as well as a configurable L1 cache of 16/48KB+-- (combined with shared memory). What we really want is for the code generator+-- to pass all inputs either as textures or global arrays, depending on what+-- device we are currently targeting.+-- -bindAcc- :: CUDA.Module- -> Val aenv- -> ArrayVar aenv- -> CIO ()-bindAcc mdl aenv (ArrayVar idx) =- let idx' = show $ idxToInt idx- Array sh ad = prj idx aenv- --- bindDim = liftIO $- CUDA.getPtr mdl ("sh" ++ idx') >>=- CUDA.pokeListArray (convertSh sh) . fst- --- arr n = "arr" ++ idx' ++ "_a" ++ show (n::Int)- tex = CUDA.getTex mdl . arr- bindTex = marshalTextureData ad (size sh) =<< liftIO (sequence' $ map tex [0..])+bindAcc :: Int -> AccKernel a -> Array dim a' -> CIO ()+bindAcc base (Kernel _ mdl _ _ _) (Array sh ad) =+ let arr n = "arrIn" ++ show base ++ "_a" ++ show (n::Int)+ tex = CUDA.getTex mdl . arr in- bindDim >> bindTex+ marshalTextureData ad (size sh) =<< liftIO (sequence' $ map tex [0..]) -bindStencil- :: Int- -> CUDA.Module- -> Array dim e- -> CIO ()-bindStencil s mdl (Array sh ad) =- let sten n = "stencil" ++ show s ++ "_a" ++ show (n::Int)- tex = CUDA.getTex mdl . sten- in- marshalTextureData ad (size sh) =<< liftIO (sequence' $ map tex [0..])+bindAccEnv :: AccKernel a -> Val aenv -> AccBindings aenv -> CIO ()+bindAccEnv (Kernel _ mdl _ _ _) aenv (AccBindings vars) = mapM_ bindAvar (Set.toList vars)+ where+ bindAvar (ArrayVar idx) =+ let idx' = show $ idxToInt idx+ Array sh ad = prj idx aenv+ --+ bindDim = liftIO $+ CUDA.getPtr mdl ("sh" ++ idx') >>=+ CUDA.pokeListArray (convertSh sh) . fst+ --+ arr n = "avar" ++ idx' ++ "_a" ++ show (n::Int)+ tex = CUDA.getTex mdl . arr+ bindTex = marshalTextureData ad (size sh) =<< liftIO (sequence' $ map tex [0..])+ in+ bindDim >> bindTex -- Kernel execution@@ -864,8 +866,8 @@ -> Int -> args -> CIO ()-execute kernel@(Kernel _ !mdl !_ !_ !_) !bindings !aenv !n !args = do- bindLifted mdl aenv bindings+execute kernel@(Kernel _ !_ !_ !_ !_) !bindings !aenv !n !args = do+ bindAccEnv kernel aenv bindings launch kernel (configure kernel n) args
+ Data/Array/Accelerate/CUDA/Persistent.hs view
@@ -0,0 +1,274 @@+{-# LANGUAGE BangPatterns #-}+{-# LANGUAGE ScopedTypeVariables #-}+-- |+-- Module : Data.Array.Accelerate.CUDA.Persistent+-- Copyright : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee+-- [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-partable (GHC extensions)+--++module Data.Array.Accelerate.CUDA.Persistent (++ KernelTable, KernelKey, KernelEntry(..),+ new, lookup, insert, persist++) where++-- friends+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++-- libraries+import Prelude hiding ( lookup, catch )+import Data.Char+import System.IO+import System.FilePath+import System.Directory+import System.Process ( ProcessHandle )+import Control.Exception+import Control.Applicative+import Control.Monad.Trans+import Data.Binary+import Data.Binary.Get+import Data.ByteString ( ByteString )+import Data.ByteString.Internal ( w2c )+import qualified Data.ByteString as B+import qualified Data.ByteString.Lazy as L+import qualified Data.HashTable.IO as HT++import qualified Foreign.CUDA.Driver as CUDA+import qualified Foreign.CUDA.Analysis as CUDA++import Paths_accelerate_cuda+++-- Interface -------------------------------------------------------------------+-- --------- --++data KernelTable = KT {-# UNPACK #-} !ProgramCache -- first level cache+ {-# UNPACK #-} !PersistentCache -- second level cache++new :: IO KernelTable+new = do+ cacheDir <- cacheDirectory+ createDirectoryIfMissing True cacheDir+ --+ local <- HT.new+ persistent <- restore (cacheDir </> "persistent.db")+ --+ return $! KT local persistent+++-- Lookup a kernel through the two-level cache system. If the kernel is found in+-- the persistent cache, it is loaded and linked into the current context.+--+lookup :: KernelTable -> KernelKey -> IO (Maybe KernelEntry)+lookup (KT kt pt) key = do+ -- First check the local cache. If we get a hit, this could be:+ -- a) currently compiling+ -- b) compiled, but not linked into the current context+ -- c) compiled & linked+ --+ v1 <- HT.lookup kt key+ case v1 of+ Just _ -> return v1+ Nothing -> do++ -- Check the persistent cache. If found, read in the associated object file+ -- and link it into the current context. Also add to the first-level cache.+ --+ -- TLM: maybe we should change KernelObject to hold a possibly empty list,+ -- so we don't have to mess with the CUDA context here.+ --+ v2 <- HT.lookup pt key+ case v2 of+ Nothing -> return Nothing+ Just () -> do+ message "found/persistent"+ cubin <- (</>) <$> cacheDirectory <*> pure (cacheFilePath key)+ ctx <- CUDA.get+ bin <- B.readFile cubin+ mdl <- CUDA.loadData bin+ let obj = KernelObject bin (FL.singleton ctx mdl)+ HT.insert kt key obj+ return $! Just obj+++-- Insert a key/value pair into the first-level cache. This does not add the+-- entry to the persistent database.+--+-- TLM: Also add to the persistent cache, or return a boolean as to whether it+-- exists there already? Would require updating that hash table as new+-- entries are added, which the functions currently do not do.+--+insert :: KernelTable -> KernelKey -> KernelEntry -> IO ()+insert (KT kt _) key val = HT.insert kt key val+++-- Local cache -----------------------------------------------------------------+-- ----------- --+--+-- Kernel code that has been generated and linked into the currently running+-- program.++-- An exact association between an accelerate computation and its+-- implementation, which is either a reference to the external compiler (nvcc)+-- or the resulting binary module.+--+-- Note that since we now support running in multiple contexts, we also need to+-- keep track of+-- a) the compute architecture the code was compiled for+-- b) which contexts have linked the code+--+-- We aren't concerned with true (typed) equality of an OpenAcc expression,+-- since we largely want to disregard the array environment; we really only want+-- to assert the type and index of those variables that are accessed by the+-- computation and no more, but we can not do that. Instead, this is keyed to+-- the generated kernel code.+--+type ProgramCache = HT.BasicHashTable KernelKey KernelEntry++type KernelKey = (CUDA.Compute, ByteString)+data KernelEntry+ -- A currently compiling external process. We record the process ID and the+ -- path of the .cu file being compiled+ --+ = CompileProcess !FilePath !ProcessHandle++ -- The raw compiled data, and the list of contexts that the object has already+ -- been linked into. If we locate this entry in the ProgramCache, it may have+ -- been inserted by an alternate but compatible device context, so just+ -- re-link into the current context.+ --+ | KernelObject {-# UNPACK #-} !ByteString+ {-# UNPACK #-} !(FullList CUDA.Context CUDA.Module)+++-- Persistent cache ------------------------------------------------------------+-- ---------------- --+--+-- Stash compiled code into the user's home directory so that they are available+-- across separate runs of the program.+--+-- TLM: we don't have any migration or versioning policy here, so cache files+-- will be kept around indefinitely. This can easily clutter the cache by+-- generating many similar kernels that differ only by, for example, an+-- embedded constant value.++type PersistentCache = HT.BasicHashTable KernelKey ()+++-- The root directory of where the various persistent cache files live; the+-- database and each individual binary object.+--+-- TLM: Is this writeable, even at a 'cabal instal --global'? Maybe we should+-- specifically choose something in the user's home directory.+--+cacheDirectory :: IO FilePath+cacheDirectory = do+ dir <- canonicalizePath =<< getDataDir+ return $ dir </> "cache"++-- A relative path to be appended to (presumably) 'cacheDirectory'.+--+cacheFilePath :: KernelKey -> FilePath+cacheFilePath (cap, key) = show cap </> foldl (flip (mangle . w2c)) ".cubin" (B.unpack key)+ where+ -- TODO: complete z-encoding? see: compiler/utils/Encoding.hs+ --+ mangle '\\' = ("zr" ++)+ mangle '/' = ("zs" ++)+ mangle c = showLitChar c+++-- The default Binary instance for lists is (necessarily) spine and value+-- strict for efficiency. For us it is better if we just lazily consume elements+-- and add them directly to the hash table so they can be collected as we go.+--+{-# INLINE getMany #-}+getMany :: Binary a => Int -> Get [a]+getMany n = go n []+ where+ go 0 xs = return xs+ go i xs = do+ x <- get+ go (i-1) (x:xs)+++-- Load the entire persistent cache index file. If it does not exist, an empty+-- file is created, so that 'persist' can always append elements.+--+restore :: FilePath -> IO PersistentCache+restore db = do+ D.when D.flush_cache $ do+ message $ "deleting persistent cache"+ cacheDir <- cacheDirectory+ removeDirectoryRecursive cacheDir+ createDirectoryIfMissing True cacheDir+ --+ exists <- doesFileExist db+ case exists of+ False -> encodeFile db (0::Int) >> HT.new+ True -> do+ store <- L.readFile db+ let (n,rest,_) = runGetState get store 0+ pt <- HT.newSized n+ --+ let go [] = return ()+ go (!k:xs) = HT.insert pt k () >> go xs+ --+ message $ "persist/restore: " ++ shows n " entries"+ go (runGet (getMany n) rest)+ pt `seq` return pt+++-- Append a single value to the persistent cache.+--+-- This moves the compiled object file (first argument) to the appropriate+-- location, and updates the database on disk.+--+persist :: FilePath -> KernelKey -> IO ()+persist cubin key = do+ cacheDir <- cacheDirectory+ let db = cacheDir </> "persistent.db"+ cacheFile = cacheDir </> cacheFilePath key+ --+ message $ "persist/save: " ++ cacheFile+ createDirectoryIfMissing True (dropFileName cacheFile)+ renameFile cubin cacheFile+ -- If the temporary and cache directories are on different disks, we must+ -- copy the file instead. Unsupported operation: (Cross-device link)+ --+ `catch` \(_ :: IOError) -> do+ copyFile cubin cacheFile+ removeFile cubin+ --+ withBinaryFile db ReadWriteMode $ \h -> do+ -- The file opens with the cursor at the beginning of the file+ --+ n <- runGet (get :: Get Int) `fmap` L.hGet h 8+ hSeek h AbsoluteSeek 0+ L.hPut h $ encode (n+1)++ -- Append the new entry to the end of file+ --+ hSeek h SeekFromEnd 0+ L.hPut h (encode key)+++-- Debug+-- -----++{-# 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+
Data/Array/Accelerate/CUDA/State.hs view
@@ -2,7 +2,7 @@ {-# LANGUAGE TemplateHaskell #-} {-# LANGUAGE TupleSections #-} {-# LANGUAGE TypeOperators #-}-{-# OPTIONS_GHC -fno-warn-orphans #-} -- CUDA.Context+{-# OPTIONS_GHC -fno-warn-orphans #-} -- Eq CUDA.Context -- | -- Module : Data.Array.Accelerate.CUDA.State -- Copyright : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee@@ -20,80 +20,40 @@ module Data.Array.Accelerate.CUDA.State ( - -- Types- CIO, KernelTable, KernelKey, KernelEntry(KernelEntry), KernelObject(KernelObject),- -- Evaluating computations- evalCUDA, defaultContext, deviceProps,- memoryTable, kernelTable, kernelName, kernelStatus+ CIO, evalCUDA, + -- Querying execution state+ defaultContext, deviceProps, activeContext, kernelTable, memoryTable+ ) where -- friends-import Data.Array.Accelerate.CUDA.FullList ( FullList ) import Data.Array.Accelerate.CUDA.Debug ( message, verbose, dump_gc, showFFloatSIBase )+import Data.Array.Accelerate.CUDA.Persistent as KT import Data.Array.Accelerate.CUDA.Array.Table as MT import Data.Array.Accelerate.CUDA.Analysis.Device -- library import Data.Label import Control.Exception-import Data.ByteString ( ByteString )-import Control.Concurrent.MVar ( MVar, newMVar )+import Control.Concurrent ( forkIO, threadDelay ) import Control.Monad.State.Strict ( StateT(..), evalStateT )-import System.Process ( ProcessHandle ) import System.Mem ( performGC )-import System.Mem.Weak ( addFinalizer )-import System.IO.Unsafe+import System.Mem.Weak ( mkWeakPtr, addFinalizer )+import System.IO.Unsafe ( unsafePerformIO ) import Text.PrettyPrint import qualified Foreign.CUDA.Driver as CUDA hiding ( device ) import qualified Foreign.CUDA.Driver.Context as CUDA-import qualified Foreign.CUDA.Analysis as CUDA-import qualified Data.HashTable.IO as HT -#ifdef ACCELERATE_CUDA_PERSISTENT_CACHE-import Data.Binary ( encodeFile, decodeFile )-import Control.Arrow ( second )-import Paths_accelerate ( getDataDir )-#endif ---- An exact association between an accelerate computation and its--- implementation, which is either a reference to the external compiler (nvcc)--- or the resulting binary module.------ Note that since we now support running in multiple contexts, we also need to--- keep track of--- a) the compute architecture the code was compiled for--- b) which contexts have linked the code------ We aren't concerned with true (typed) equality of an OpenAcc expression,--- since we largely want to disregard the array environment; we really only want--- to assert the type and index of those variables that are accessed by the--- computation and no more, but we can not do that. Instead, this is keyed to--- the generated kernel code.----type KernelTable = HT.BasicHashTable KernelKey KernelEntry--type KernelKey = (CUDA.Compute, ByteString)-data KernelEntry = KernelEntry- {- _kernelName :: !FilePath,- _kernelStatus :: !(Either ProcessHandle KernelObject)- }--data KernelObject = KernelObject- {- _binaryData :: !ByteString,- _activeContexts :: {-# UNPACK #-} !(FullList CUDA.Context CUDA.Module)- }- -- The state token for CUDA accelerated array operations -- type CIO = StateT CUDAState IO data CUDAState = CUDAState { _deviceProps :: !CUDA.DeviceProperties,+ _activeContext :: {-# UNPACK #-} !Context, _kernelTable :: {-# UNPACK #-} !KernelTable, _memoryTable :: {-# UNPACK #-} !MemoryTable }@@ -101,7 +61,7 @@ instance Eq CUDA.Context where CUDA.Context p1 == CUDA.Context p2 = p1 == p2 -$(mkLabels [''CUDAState, ''KernelEntry])+$(mkLabels [''CUDAState]) -- Execution State@@ -117,37 +77,66 @@ CUDA.push ctx dev <- CUDA.device prp <- CUDA.props dev- return $! CUDAState prp knl mem+ weak_ctx <- mkWeakPtr ctx Nothing+ return $! CUDAState prp (Context ctx weak_ctx) theKernelTable theMemoryTable - -- one-shot top-level mutable state- {-# NOINLINE mem #-}- {-# NOINLINE knl #-}- mem = unsafePerformIO MT.new- knl = unsafePerformIO HT.new +-- Top-level mutable state+-- -----------------------+--+-- It is important to keep some information alive for the entire run of the+-- program, not just a single execution. These tokens use unsafePerformIO to+-- ensure they are executed only once, and reused for subsequent invocations.+-- +{-# NOINLINE theMemoryTable #-}+theMemoryTable :: MemoryTable+theMemoryTable = unsafePerformIO $ do+ message dump_gc "gc: initialise memory table"+ keepAlive =<< MT.new+++{-# NOINLINE theKernelTable #-}+theKernelTable :: KernelTable+theKernelTable = unsafePerformIO $ do+ message dump_gc "gc: initialise kernel table"+ keepAlive =<< KT.new++ -- Select and initialise a default CUDA device, and create a new execution -- context. The device is selected based on compute capability and estimated -- maximum throughput. -- {-# NOINLINE defaultContext #-}-defaultContext :: MVar CUDA.Context+defaultContext :: CUDA.Context defaultContext = unsafePerformIO $ do CUDA.initialise [] (dev,prp) <- selectBestDevice ctx <- CUDA.create dev [CUDA.SchedAuto] >> CUDA.pop- ref <- newMVar ctx -- message dump_gc $ "gc: initialise context" message verbose $ deviceInfo dev prp -- addFinalizer ctx $ do- message dump_gc $ "gc: finalise context"+ message dump_gc $ "gc: finalise context" -- should never happen! CUDA.destroy ctx --- return ref+ keepAlive ctx +-- Make sure the GC knows that we want to keep this thing alive past the end of+-- 'evalCUDA'.+--+-- We may want to introduce some way to actually shut this down if, for example,+-- the object has not been accessed in a while, and so let it be collected.+--+keepAlive :: a -> IO a+keepAlive x = forkIO (caffeine x) >> return x+ where+ caffeine hit = do threadDelay 5000000 -- microseconds = 5 seconds+ caffeine hit++ -- Debugging -- --------- @@ -172,36 +161,4 @@ clock = showFFloatSIBase (Just 2) 1000 (fromIntegral $ CUDA.clockRate prp * 1000 :: Double) "Hz" mem = showFFloatSIBase (Just 0) 1024 (fromIntegral $ CUDA.totalGlobalMem prp :: Double) "B" at = char '@'----- Persistent caching (deprecated)--- ---------------------------------#ifdef ACCELERATE_CUDA_PERSISTENT_CACHE--- Load and save the persistent kernel index file----indexFileName :: IO FilePath-indexFileName = do- tmp <- (</> "cache") `fmap` getDataDir- dir <- createDirectoryIfMissing True tmp >> canonicalizePath tmp- return (dir </> "_index")--saveIndexFile :: CUDAState -> IO ()-saveIndexFile s = do- ind <- indexFileName- encodeFile ind . map (second _kernelName) =<< HT.toList (_kernelTable s)---- Read the kernel index map file (if it exists), loading modules into the--- current context----loadIndexFile :: IO (KernelTable, Int)-loadIndexFile = do- f <- indexFileName- x <- doesFileExist f- e <- if x then mapM reload =<< decodeFile f- else return []- (,length e) <$> HT.fromList hashAccKey e- where- reload (k,n) = (k,) . KernelEntry n . Right <$> CUDA.loadFile (n `replaceExtension` ".cubin")-#endif
+ accelerate-cuda.buildinfo.in view
@@ -0,0 +1,3 @@+ghc-options: @ghc_flags@+cc-options: @cpp_flags@+
accelerate-cuda.cabal view
@@ -1,5 +1,5 @@ Name: accelerate-cuda-Version: 0.12.0.0+Version: 0.12.1.0 Cabal-version: >= 1.6 Tested-with: GHC >= 7.4 Build-type: Configure@@ -9,9 +9,11 @@ This library implements a backend for the Accelerate language instrumented for parallel execution on CUDA-capable NVIDIA GPUs. .- To use this backend you need CUDA version 3.x or later installed. Note that- currently there is no support for 'Char' and 'Bool' arrays (this is a- limitation of the front-end language).+ To use this backend you need CUDA version 3.x or later installed, which you+ can find at the NVIDIA Developer Zone.+ .+ <http://developer.nvidia.com/cuda-downloads>+ . License: BSD3 License-file: LICENSE@@ -35,22 +37,38 @@ cubits/accelerate_cuda_shape.h cubits/accelerate_cuda_stencil.h cubits/accelerate_cuda_texture.h+ cubits/accelerate_cuda_type.h cubits/accelerate_cuda_util.h Extra-tmp-files: config.status config.log+ accelerate-cuda.buildinfo -- generated by configure cubits/accelerate_cuda_shape.h -- generated by configure Extra-source-files: configure+ accelerate-cuda.buildinfo.in cubits/accelerate_cuda_shape.h.in include/accelerate.h --- Flag pcache--- Description: Enable the persistent caching of the compiled CUDA modules (experimental)--- Default: False- Flag debug- Description: Enable tracing message flags+ Description:+ Enable tracing message flags. These are read from the command-line+ arguments, which is convenient but may cause problems interacting with the+ user program, so is not enabled by default. The available options:+ .+ * -ddump-cc: print the generated code, kernel table management information,+ nvcc compiler warnings, and thread & resource statistics+ .+ * -ddebug-cc: compile code with debugging symbols, suitable for 'cuda-gdb'+ .+ * -ddump-exec: print each kernel name as it is invoked+ .+ * -ddump-gc: print memory management information+ .+ * -dverbose: other, uncategorised messages+ .+ * -fflush-cache: delete the persistent kernel cache+ . Default: False Flag bounds-checks@@ -68,7 +86,7 @@ Library Include-Dirs: include - Build-depends: accelerate == 0.12.*,+ Build-depends: accelerate >= 0.12.1 && < 0.13, array >= 0.3, base == 4.*, binary >= 0.5,@@ -88,7 +106,6 @@ pretty >= 1.0, process >= 1.0, srcloc >= 0.1,- symbol >= 0.1, transformers >= 0.2, unordered-containers >= 0.1.4 @@ -119,12 +136,10 @@ Data.Array.Accelerate.CUDA.Debug Data.Array.Accelerate.CUDA.Execute Data.Array.Accelerate.CUDA.FullList+ Data.Array.Accelerate.CUDA.Persistent Data.Array.Accelerate.CUDA.State Data.Array.Accelerate.Internal.Check Paths_accelerate_cuda---- if flag(pcache)--- CPP-options: -DACCELERATE_CUDA_PERSISTENT_CACHE if flag(debug) cpp-options: -DACCELERATE_DEBUG
configure view
@@ -1,13 +1,13 @@ #! /bin/sh # Guess values for system-dependent variables and create Makefiles.-# Generated by GNU Autoconf 2.65 for accelerate-cuda 0.9.0.0.+# Generated by GNU Autoconf 2.68 for accelerate-cuda 0.13.0.0. #-# Report bugs to <accelerate@projects.haskell.org>.+# Report bugs to <accelerate-haskell@googlegroups.com>. # # # Copyright (C) 1992, 1993, 1994, 1995, 1996, 1998, 1999, 2000, 2001,-# 2002, 2003, 2004, 2005, 2006, 2007, 2008, 2009 Free Software Foundation,-# Inc.+# 2002, 2003, 2004, 2005, 2006, 2007, 2008, 2009, 2010 Free Software+# Foundation, Inc. # # # This configure script is free software; the Free Software Foundation@@ -91,6 +91,7 @@ IFS=" "" $as_nl" # Find who we are. Look in the path if we contain no directory separator.+as_myself= case $0 in #(( *[\\/]* ) as_myself=$0 ;; *) as_save_IFS=$IFS; IFS=$PATH_SEPARATOR@@ -215,11 +216,18 @@ # We cannot yet assume a decent shell, so we have to provide a # neutralization value for shells without unset; and this also # works around shells that cannot unset nonexistent variables.+ # Preserve -v and -x to the replacement shell. BASH_ENV=/dev/null ENV=/dev/null (unset BASH_ENV) >/dev/null 2>&1 && unset BASH_ENV ENV export CONFIG_SHELL- exec "$CONFIG_SHELL" "$as_myself" ${1+"$@"}+ case $- in # ((((+ *v*x* | *x*v* ) as_opts=-vx ;;+ *v* ) as_opts=-v ;;+ *x* ) as_opts=-x ;;+ * ) as_opts= ;;+ esac+ exec "$CONFIG_SHELL" $as_opts "$as_myself" ${1+"$@"} fi if test x$as_have_required = xno; then :@@ -230,7 +238,7 @@ $as_echo "$0: be upgraded to zsh 4.3.4 or later." else $as_echo "$0: Please tell bug-autoconf@gnu.org and-$0: accelerate@projects.haskell.org about your system,+$0: accelerate-haskell@googlegroups.com about your system, $0: including any error possibly output before this $0: message. Then install a modern shell, or manually run $0: the script under such a shell if you do have one."@@ -318,7 +326,7 @@ test -d "$as_dir" && break done test -z "$as_dirs" || eval "mkdir $as_dirs"- } || test -d "$as_dir" || as_fn_error "cannot create directory $as_dir"+ } || test -d "$as_dir" || as_fn_error $? "cannot create directory $as_dir" } # as_fn_mkdir_p@@ -358,19 +366,19 @@ fi # as_fn_arith -# as_fn_error ERROR [LINENO LOG_FD]-# ---------------------------------+# as_fn_error STATUS ERROR [LINENO LOG_FD]+# ---------------------------------------- # Output "`basename $0`: error: ERROR" to stderr. If LINENO and LOG_FD are # provided, also output the error to LOG_FD, referencing LINENO. Then exit the-# script with status $?, using 1 if that was 0.+# script with STATUS, using 1 if that was 0. as_fn_error () {- as_status=$?; test $as_status -eq 0 && as_status=1- if test "$3"; then- as_lineno=${as_lineno-"$2"} as_lineno_stack=as_lineno_stack=$as_lineno_stack- $as_echo "$as_me:${as_lineno-$LINENO}: error: $1" >&$3+ as_status=$1; test $as_status -eq 0 && as_status=1+ if test "$4"; then+ as_lineno=${as_lineno-"$3"} as_lineno_stack=as_lineno_stack=$as_lineno_stack+ $as_echo "$as_me:${as_lineno-$LINENO}: error: $2" >&$4 fi- $as_echo "$as_me: error: $1" >&2+ $as_echo "$as_me: error: $2" >&2 as_fn_exit $as_status } # as_fn_error @@ -532,7 +540,7 @@ exec 6>&1 # Name of the host.-# hostname on some systems (SVR3.2, Linux) returns a bogus exit status,+# hostname on some systems (SVR3.2, old GNU/Linux) returns a bogus exit status, # so uname gets run too. ac_hostname=`(hostname || uname -n) 2>/dev/null | sed 1q` @@ -551,15 +559,18 @@ # Identity of this package. PACKAGE_NAME='accelerate-cuda' PACKAGE_TARNAME='accelerate-cuda'-PACKAGE_VERSION='0.9.0.0'-PACKAGE_STRING='accelerate-cuda 0.9.0.0'-PACKAGE_BUGREPORT='accelerate@projects.haskell.org'+PACKAGE_VERSION='0.13.0.0'+PACKAGE_STRING='accelerate-cuda 0.13.0.0'+PACKAGE_BUGREPORT='accelerate-haskell@googlegroups.com' PACKAGE_URL='' ac_unique_file="Data/Array/Accelerate/CUDA.hs" ac_subst_vars='LTLIBOBJS LIBOBJS-hs_int_type+cpp_flags+ghc_flags+type_hs_char+type_hs_int GHC OBJEXT EXEEXT@@ -681,8 +692,9 @@ fi case $ac_option in- *=*) ac_optarg=`expr "X$ac_option" : '[^=]*=\(.*\)'` ;;- *) ac_optarg=yes ;;+ *=?*) ac_optarg=`expr "X$ac_option" : '[^=]*=\(.*\)'` ;;+ *=) ac_optarg= ;;+ *) ac_optarg=yes ;; esac # Accept the important Cygnus configure options, so we can diagnose typos.@@ -727,7 +739,7 @@ ac_useropt=`expr "x$ac_option" : 'x-*disable-\(.*\)'` # Reject names that are not valid shell variable names. expr "x$ac_useropt" : ".*[^-+._$as_cr_alnum]" >/dev/null &&- as_fn_error "invalid feature name: $ac_useropt"+ as_fn_error $? "invalid feature name: $ac_useropt" ac_useropt_orig=$ac_useropt ac_useropt=`$as_echo "$ac_useropt" | sed 's/[-+.]/_/g'` case $ac_user_opts in@@ -753,7 +765,7 @@ ac_useropt=`expr "x$ac_option" : 'x-*enable-\([^=]*\)'` # Reject names that are not valid shell variable names. expr "x$ac_useropt" : ".*[^-+._$as_cr_alnum]" >/dev/null &&- as_fn_error "invalid feature name: $ac_useropt"+ as_fn_error $? "invalid feature name: $ac_useropt" ac_useropt_orig=$ac_useropt ac_useropt=`$as_echo "$ac_useropt" | sed 's/[-+.]/_/g'` case $ac_user_opts in@@ -957,7 +969,7 @@ ac_useropt=`expr "x$ac_option" : 'x-*with-\([^=]*\)'` # Reject names that are not valid shell variable names. expr "x$ac_useropt" : ".*[^-+._$as_cr_alnum]" >/dev/null &&- as_fn_error "invalid package name: $ac_useropt"+ as_fn_error $? "invalid package name: $ac_useropt" ac_useropt_orig=$ac_useropt ac_useropt=`$as_echo "$ac_useropt" | sed 's/[-+.]/_/g'` case $ac_user_opts in@@ -973,7 +985,7 @@ ac_useropt=`expr "x$ac_option" : 'x-*without-\(.*\)'` # Reject names that are not valid shell variable names. expr "x$ac_useropt" : ".*[^-+._$as_cr_alnum]" >/dev/null &&- as_fn_error "invalid package name: $ac_useropt"+ as_fn_error $? "invalid package name: $ac_useropt" ac_useropt_orig=$ac_useropt ac_useropt=`$as_echo "$ac_useropt" | sed 's/[-+.]/_/g'` case $ac_user_opts in@@ -1003,8 +1015,8 @@ | --x-librar=* | --x-libra=* | --x-libr=* | --x-lib=* | --x-li=* | --x-l=*) x_libraries=$ac_optarg ;; - -*) as_fn_error "unrecognized option: \`$ac_option'-Try \`$0 --help' for more information."+ -*) as_fn_error $? "unrecognized option: \`$ac_option'+Try \`$0 --help' for more information" ;; *=*)@@ -1012,7 +1024,7 @@ # Reject names that are not valid shell variable names. case $ac_envvar in #( '' | [0-9]* | *[!_$as_cr_alnum]* )- as_fn_error "invalid variable name: \`$ac_envvar'" ;;+ as_fn_error $? "invalid variable name: \`$ac_envvar'" ;; esac eval $ac_envvar=\$ac_optarg export $ac_envvar ;;@@ -1022,7 +1034,7 @@ $as_echo "$as_me: WARNING: you should use --build, --host, --target" >&2 expr "x$ac_option" : ".*[^-._$as_cr_alnum]" >/dev/null && $as_echo "$as_me: WARNING: invalid host type: $ac_option" >&2- : ${build_alias=$ac_option} ${host_alias=$ac_option} ${target_alias=$ac_option}+ : "${build_alias=$ac_option} ${host_alias=$ac_option} ${target_alias=$ac_option}" ;; esac@@ -1030,13 +1042,13 @@ if test -n "$ac_prev"; then ac_option=--`echo $ac_prev | sed 's/_/-/g'`- as_fn_error "missing argument to $ac_option"+ as_fn_error $? "missing argument to $ac_option" fi if test -n "$ac_unrecognized_opts"; then case $enable_option_checking in no) ;;- fatal) as_fn_error "unrecognized options: $ac_unrecognized_opts" ;;+ fatal) as_fn_error $? "unrecognized options: $ac_unrecognized_opts" ;; *) $as_echo "$as_me: WARNING: unrecognized options: $ac_unrecognized_opts" >&2 ;; esac fi@@ -1059,7 +1071,7 @@ [\\/$]* | ?:[\\/]* ) continue;; NONE | '' ) case $ac_var in *prefix ) continue;; esac;; esac- as_fn_error "expected an absolute directory name for --$ac_var: $ac_val"+ as_fn_error $? "expected an absolute directory name for --$ac_var: $ac_val" done # There might be people who depend on the old broken behavior: `$host'@@ -1073,8 +1085,8 @@ if test "x$host_alias" != x; then if test "x$build_alias" = x; then cross_compiling=maybe- $as_echo "$as_me: WARNING: If you wanted to set the --build type, don't use --host.- If a cross compiler is detected then cross compile mode will be used." >&2+ $as_echo "$as_me: WARNING: if you wanted to set the --build type, don't use --host.+ If a cross compiler is detected then cross compile mode will be used" >&2 elif test "x$build_alias" != "x$host_alias"; then cross_compiling=yes fi@@ -1089,9 +1101,9 @@ ac_pwd=`pwd` && test -n "$ac_pwd" && ac_ls_di=`ls -di .` && ac_pwd_ls_di=`cd "$ac_pwd" && ls -di .` ||- as_fn_error "working directory cannot be determined"+ as_fn_error $? "working directory cannot be determined" test "X$ac_ls_di" = "X$ac_pwd_ls_di" ||- as_fn_error "pwd does not report name of working directory"+ as_fn_error $? "pwd does not report name of working directory" # Find the source files, if location was not specified.@@ -1130,11 +1142,11 @@ fi if test ! -r "$srcdir/$ac_unique_file"; then test "$ac_srcdir_defaulted" = yes && srcdir="$ac_confdir or .."- as_fn_error "cannot find sources ($ac_unique_file) in $srcdir"+ as_fn_error $? "cannot find sources ($ac_unique_file) in $srcdir" fi ac_msg="sources are in $srcdir, but \`cd $srcdir' does not work" ac_abs_confdir=`(- cd "$srcdir" && test -r "./$ac_unique_file" || as_fn_error "$ac_msg"+ cd "$srcdir" && test -r "./$ac_unique_file" || as_fn_error $? "$ac_msg" pwd)` # When building in place, set srcdir=. if test "$ac_abs_confdir" = "$ac_pwd"; then@@ -1160,7 +1172,7 @@ # Omit some internal or obsolete options to make the list less imposing. # This message is too long to be a string in the A/UX 3.1 sh. cat <<_ACEOF-\`configure' configures accelerate-cuda 0.9.0.0 to adapt to many kinds of systems.+\`configure' configures accelerate-cuda 0.13.0.0 to adapt to many kinds of systems. Usage: $0 [OPTION]... [VAR=VALUE]... @@ -1174,7 +1186,7 @@ --help=short display options specific to this package --help=recursive display the short help of all the included packages -V, --version display version information and exit- -q, --quiet, --silent do not print \`checking...' messages+ -q, --quiet, --silent do not print \`checking ...' messages --cache-file=FILE cache test results in FILE [disabled] -C, --config-cache alias for \`--cache-file=config.cache' -n, --no-create do not create output files@@ -1221,7 +1233,7 @@ if test -n "$ac_init_help"; then case $ac_init_help in- short | recursive ) echo "Configuration of accelerate-cuda 0.9.0.0:";;+ short | recursive ) echo "Configuration of accelerate-cuda 0.13.0.0:";; esac cat <<\_ACEOF @@ -1237,7 +1249,7 @@ Use these variables to override the choices made by `configure' or to help it to find libraries and programs with nonstandard names/locations. -Report bugs to <accelerate@projects.haskell.org>.+Report bugs to <accelerate-haskell@googlegroups.com>. _ACEOF ac_status=$? fi@@ -1300,10 +1312,10 @@ test -n "$ac_init_help" && exit $ac_status if $ac_init_version; then cat <<\_ACEOF-accelerate-cuda configure 0.9.0.0-generated by GNU Autoconf 2.65+accelerate-cuda configure 0.13.0.0+generated by GNU Autoconf 2.68 -Copyright (C) 2009 Free Software Foundation, Inc.+Copyright (C) 2010 Free Software Foundation, Inc. This configure script is free software; the Free Software Foundation gives unlimited permission to copy, distribute and modify it. _ACEOF@@ -1347,7 +1359,7 @@ ac_retval=1 fi- eval $as_lineno_stack; test "x$as_lineno_stack" = x && { as_lineno=; unset as_lineno;}+ eval $as_lineno_stack; ${as_lineno_stack:+:} unset as_lineno as_fn_set_status $ac_retval } # ac_fn_cxx_try_compile@@ -1355,8 +1367,8 @@ This file contains any messages produced by compilers while running configure, to aid debugging if configure makes a mistake. -It was created by accelerate-cuda $as_me 0.9.0.0, which was-generated by GNU Autoconf 2.65. Invocation command line was+It was created by accelerate-cuda $as_me 0.13.0.0, which was+generated by GNU Autoconf 2.68. Invocation command line was $ $0 $@ @@ -1466,11 +1478,9 @@ { echo - cat <<\_ASBOX-## ---------------- ##+ $as_echo "## ---------------- ## ## Cache variables. ##-## ---------------- ##-_ASBOX+## ---------------- ##" echo # The following way of writing the cache mishandles newlines in values, (@@ -1504,11 +1514,9 @@ ) echo - cat <<\_ASBOX-## ----------------- ##+ $as_echo "## ----------------- ## ## Output variables. ##-## ----------------- ##-_ASBOX+## ----------------- ##" echo for ac_var in $ac_subst_vars do@@ -1521,11 +1529,9 @@ echo if test -n "$ac_subst_files"; then- cat <<\_ASBOX-## ------------------- ##+ $as_echo "## ------------------- ## ## File substitutions. ##-## ------------------- ##-_ASBOX+## ------------------- ##" echo for ac_var in $ac_subst_files do@@ -1539,11 +1545,9 @@ fi if test -s confdefs.h; then- cat <<\_ASBOX-## ----------- ##+ $as_echo "## ----------- ## ## confdefs.h. ##-## ----------- ##-_ASBOX+## ----------- ##" echo cat confdefs.h echo@@ -1598,7 +1602,12 @@ ac_site_file1=NONE ac_site_file2=NONE if test -n "$CONFIG_SITE"; then- ac_site_file1=$CONFIG_SITE+ # We do not want a PATH search for config.site.+ case $CONFIG_SITE in #((+ -*) ac_site_file1=./$CONFIG_SITE;;+ */*) ac_site_file1=$CONFIG_SITE;;+ *) ac_site_file1=./$CONFIG_SITE;;+ esac elif test "x$prefix" != xNONE; then ac_site_file1=$prefix/share/config.site ac_site_file2=$prefix/etc/config.site@@ -1613,7 +1622,11 @@ { $as_echo "$as_me:${as_lineno-$LINENO}: loading site script $ac_site_file" >&5 $as_echo "$as_me: loading site script $ac_site_file" >&6;} sed 's/^/| /' "$ac_site_file" >&5- . "$ac_site_file"+ . "$ac_site_file" \+ || { { $as_echo "$as_me:${as_lineno-$LINENO}: error: in \`$ac_pwd':" >&5+$as_echo "$as_me: error: in \`$ac_pwd':" >&2;}+as_fn_error $? "failed to load site script $ac_site_file+See \`config.log' for more details" "$LINENO" 5; } fi done @@ -1689,7 +1702,7 @@ $as_echo "$as_me: error: in \`$ac_pwd':" >&2;} { $as_echo "$as_me:${as_lineno-$LINENO}: error: changes in the environment can compromise the build" >&5 $as_echo "$as_me: error: changes in the environment can compromise the build" >&2;}- as_fn_error "run \`make distclean' and/or \`rm $cache_file' and start over" "$LINENO" 5+ as_fn_error $? "run \`make distclean' and/or \`rm $cache_file' and start over" "$LINENO" 5 fi ## -------------------- ## ## Main body of script. ##@@ -1703,7 +1716,7 @@ -ac_config_files="$ac_config_files cubits/accelerate_cuda_shape.h"+ac_config_files="$ac_config_files accelerate-cuda.buildinfo cubits/accelerate_cuda_shape.h" # This doesn't pick up a user specified '--with-compiler' flag@@ -1724,7 +1737,7 @@ set dummy $ac_tool_prefix$ac_prog; ac_word=$2 { $as_echo "$as_me:${as_lineno-$LINENO}: checking for $ac_word" >&5 $as_echo_n "checking for $ac_word... " >&6; }-if test "${ac_cv_prog_CXX+set}" = set; then :+if ${ac_cv_prog_CXX+:} false; then : $as_echo_n "(cached) " >&6 else if test -n "$CXX"; then@@ -1768,7 +1781,7 @@ set dummy $ac_prog; ac_word=$2 { $as_echo "$as_me:${as_lineno-$LINENO}: checking for $ac_word" >&5 $as_echo_n "checking for $ac_word... " >&6; }-if test "${ac_cv_prog_ac_ct_CXX+set}" = set; then :+if ${ac_cv_prog_ac_ct_CXX+:} false; then : $as_echo_n "(cached) " >&6 else if test -n "$ac_ct_CXX"; then@@ -1931,9 +1944,8 @@ { { $as_echo "$as_me:${as_lineno-$LINENO}: error: in \`$ac_pwd':" >&5 $as_echo "$as_me: error: in \`$ac_pwd':" >&2;}-{ as_fn_set_status 77-as_fn_error "C++ compiler cannot create executables-See \`config.log' for more details." "$LINENO" 5; }; }+as_fn_error 77 "C++ compiler cannot create executables+See \`config.log' for more details" "$LINENO" 5; } else { $as_echo "$as_me:${as_lineno-$LINENO}: result: yes" >&5 $as_echo "yes" >&6; }@@ -1975,8 +1987,8 @@ else { { $as_echo "$as_me:${as_lineno-$LINENO}: error: in \`$ac_pwd':" >&5 $as_echo "$as_me: error: in \`$ac_pwd':" >&2;}-as_fn_error "cannot compute suffix of executables: cannot compile and link-See \`config.log' for more details." "$LINENO" 5; }+as_fn_error $? "cannot compute suffix of executables: cannot compile and link+See \`config.log' for more details" "$LINENO" 5; } fi rm -f conftest conftest$ac_cv_exeext { $as_echo "$as_me:${as_lineno-$LINENO}: result: $ac_cv_exeext" >&5@@ -2033,9 +2045,9 @@ else { { $as_echo "$as_me:${as_lineno-$LINENO}: error: in \`$ac_pwd':" >&5 $as_echo "$as_me: error: in \`$ac_pwd':" >&2;}-as_fn_error "cannot run C++ compiled programs.+as_fn_error $? "cannot run C++ compiled programs. If you meant to cross compile, use \`--host'.-See \`config.log' for more details." "$LINENO" 5; }+See \`config.log' for more details" "$LINENO" 5; } fi fi fi@@ -2046,7 +2058,7 @@ ac_clean_files=$ac_clean_files_save { $as_echo "$as_me:${as_lineno-$LINENO}: checking for suffix of object files" >&5 $as_echo_n "checking for suffix of object files... " >&6; }-if test "${ac_cv_objext+set}" = set; then :+if ${ac_cv_objext+:} false; then : $as_echo_n "(cached) " >&6 else cat confdefs.h - <<_ACEOF >conftest.$ac_ext@@ -2086,8 +2098,8 @@ { { $as_echo "$as_me:${as_lineno-$LINENO}: error: in \`$ac_pwd':" >&5 $as_echo "$as_me: error: in \`$ac_pwd':" >&2;}-as_fn_error "cannot compute suffix of object files: cannot compile-See \`config.log' for more details." "$LINENO" 5; }+as_fn_error $? "cannot compute suffix of object files: cannot compile+See \`config.log' for more details" "$LINENO" 5; } fi rm -f conftest.$ac_cv_objext conftest.$ac_ext fi@@ -2097,7 +2109,7 @@ ac_objext=$OBJEXT { $as_echo "$as_me:${as_lineno-$LINENO}: checking whether we are using the GNU C++ compiler" >&5 $as_echo_n "checking whether we are using the GNU C++ compiler... " >&6; }-if test "${ac_cv_cxx_compiler_gnu+set}" = set; then :+if ${ac_cv_cxx_compiler_gnu+:} false; then : $as_echo_n "(cached) " >&6 else cat confdefs.h - <<_ACEOF >conftest.$ac_ext@@ -2134,7 +2146,7 @@ ac_save_CXXFLAGS=$CXXFLAGS { $as_echo "$as_me:${as_lineno-$LINENO}: checking whether $CXX accepts -g" >&5 $as_echo_n "checking whether $CXX accepts -g... " >&6; }-if test "${ac_cv_prog_cxx_g+set}" = set; then :+if ${ac_cv_prog_cxx_g+:} false; then : $as_echo_n "(cached) " >&6 else ac_save_cxx_werror_flag=$ac_cxx_werror_flag@@ -2220,7 +2232,7 @@ set dummy ghc; ac_word=$2 { $as_echo "$as_me:${as_lineno-$LINENO}: checking for $ac_word" >&5 $as_echo_n "checking for $ac_word... " >&6; }-if test "${ac_cv_path_GHC+set}" = set; then :+if ${ac_cv_path_GHC+:} false; then : $as_echo_n "(cached) " >&6 else case $GHC in@@ -2256,20 +2268,48 @@ fi -{ $as_echo "$as_me:${as_lineno-$LINENO}: checking size of Int" >&5+++++ { $as_echo "$as_me:${as_lineno-$LINENO}: checking size of Int" >&5 $as_echo_n "checking size of Int... " >&6; } -hs_int_size=`$GHC -w -ignore-dot-ghci -e "Foreign.sizeOf (undefined::Int)"`-{ $as_echo "$as_me:${as_lineno-$LINENO}: result: $hs_int_size" >&5-$as_echo "$hs_int_size" >&6; }+ sizeof_hs_Int=`$GHC -w -ignore-dot-ghci -e "Foreign.sizeOf (undefined::Int)"`+ { $as_echo "$as_me:${as_lineno-$LINENO}: result: $sizeof_hs_Int" >&5+$as_echo "$sizeof_hs_Int" >&6; } -case $hs_int_size in- 4) hs_int_type=int32_t ;;- 8) hs_int_type=int64_t ;;- *) as_fn_error "could not determine size of a Haskell Int" "$LINENO" 5-esac+ case $sizeof_hs_Int in+ 4) type_hs_int=Int32 ;;+ 8) type_hs_int=Int64 ;;+ esac + def="-DSIZEOF_HS$(echo Int | tr [:lower:] [:upper:])=$sizeof_hs_Int"+ cpp_flags="$cpp_flags $def"+ ghc_flags="$ghc_flags -optP$def" ++ { $as_echo "$as_me:${as_lineno-$LINENO}: checking size of Char" >&5+$as_echo_n "checking size of Char... " >&6; }++ sizeof_hs_Char=`$GHC -w -ignore-dot-ghci -e "Foreign.sizeOf (undefined::Char)"`+ { $as_echo "$as_me:${as_lineno-$LINENO}: result: $sizeof_hs_Char" >&5+$as_echo "$sizeof_hs_Char" >&6; }++ case $sizeof_hs_Char in+ 4) type_hs_char=Word32 ;;+ 8) type_hs_char=Word64 ;;+ esac++ def="-DSIZEOF_HS$(echo Char | tr [:lower:] [:upper:])=$sizeof_hs_Char"+ cpp_flags="$cpp_flags $def"+ ghc_flags="$ghc_flags -optP$def"++++++ cat >confcache <<\_ACEOF # This file is a shell script that caches the results of configure # tests run on this system so they can be shared between configure@@ -2334,10 +2374,21 @@ :end' >>confcache if diff "$cache_file" confcache >/dev/null 2>&1; then :; else if test -w "$cache_file"; then- test "x$cache_file" != "x/dev/null" &&+ if test "x$cache_file" != "x/dev/null"; then { $as_echo "$as_me:${as_lineno-$LINENO}: updating cache $cache_file" >&5 $as_echo "$as_me: updating cache $cache_file" >&6;}- cat confcache >$cache_file+ if test ! -f "$cache_file" || test -h "$cache_file"; then+ cat confcache >"$cache_file"+ else+ case $cache_file in #(+ */* | ?:*)+ mv -f confcache "$cache_file"$$ &&+ mv -f "$cache_file"$$ "$cache_file" ;; #(+ *)+ mv -f confcache "$cache_file" ;;+ esac+ fi+ fi else { $as_echo "$as_me:${as_lineno-$LINENO}: not updating unwritable cache $cache_file" >&5 $as_echo "$as_me: not updating unwritable cache $cache_file" >&6;}@@ -2389,6 +2440,7 @@ ac_libobjs= ac_ltlibobjs=+U= for ac_i in : $LIBOBJS; do test "x$ac_i" = x: && continue # 1. Remove the extension, and $U if already installed. ac_script='s/\$U\././;s/\.o$//;s/\.obj$//'@@ -2404,7 +2456,7 @@ -: ${CONFIG_STATUS=./config.status}+: "${CONFIG_STATUS=./config.status}" ac_write_fail=0 ac_clean_files_save=$ac_clean_files ac_clean_files="$ac_clean_files $CONFIG_STATUS"@@ -2505,6 +2557,7 @@ IFS=" "" $as_nl" # Find who we are. Look in the path if we contain no directory separator.+as_myself= case $0 in #(( *[\\/]* ) as_myself=$0 ;; *) as_save_IFS=$IFS; IFS=$PATH_SEPARATOR@@ -2550,19 +2603,19 @@ (unset CDPATH) >/dev/null 2>&1 && unset CDPATH -# as_fn_error ERROR [LINENO LOG_FD]-# ---------------------------------+# as_fn_error STATUS ERROR [LINENO LOG_FD]+# ---------------------------------------- # Output "`basename $0`: error: ERROR" to stderr. If LINENO and LOG_FD are # provided, also output the error to LOG_FD, referencing LINENO. Then exit the-# script with status $?, using 1 if that was 0.+# script with STATUS, using 1 if that was 0. as_fn_error () {- as_status=$?; test $as_status -eq 0 && as_status=1- if test "$3"; then- as_lineno=${as_lineno-"$2"} as_lineno_stack=as_lineno_stack=$as_lineno_stack- $as_echo "$as_me:${as_lineno-$LINENO}: error: $1" >&$3+ as_status=$1; test $as_status -eq 0 && as_status=1+ if test "$4"; then+ as_lineno=${as_lineno-"$3"} as_lineno_stack=as_lineno_stack=$as_lineno_stack+ $as_echo "$as_me:${as_lineno-$LINENO}: error: $2" >&$4 fi- $as_echo "$as_me: error: $1" >&2+ $as_echo "$as_me: error: $2" >&2 as_fn_exit $as_status } # as_fn_error @@ -2758,7 +2811,7 @@ test -d "$as_dir" && break done test -z "$as_dirs" || eval "mkdir $as_dirs"- } || test -d "$as_dir" || as_fn_error "cannot create directory $as_dir"+ } || test -d "$as_dir" || as_fn_error $? "cannot create directory $as_dir" } # as_fn_mkdir_p@@ -2811,8 +2864,8 @@ # report actual input values of CONFIG_FILES etc. instead of their # values after options handling. ac_log="-This file was extended by accelerate-cuda $as_me 0.9.0.0, which was-generated by GNU Autoconf 2.65. Invocation command line was+This file was extended by accelerate-cuda $as_me 0.13.0.0, which was+generated by GNU Autoconf 2.68. Invocation command line was CONFIG_FILES = $CONFIG_FILES CONFIG_HEADERS = $CONFIG_HEADERS@@ -2858,17 +2911,17 @@ Configuration files: $config_files -Report bugs to <accelerate@projects.haskell.org>."+Report bugs to <accelerate-haskell@googlegroups.com>." _ACEOF cat >>$CONFIG_STATUS <<_ACEOF || ac_write_fail=1 ac_cs_config="`$as_echo "$ac_configure_args" | sed 's/^ //; s/[\\""\`\$]/\\\\&/g'`" ac_cs_version="\\-accelerate-cuda config.status 0.9.0.0-configured by $0, generated by GNU Autoconf 2.65,+accelerate-cuda config.status 0.13.0.0+configured by $0, generated by GNU Autoconf 2.68, with options \\"\$ac_cs_config\\" -Copyright (C) 2009 Free Software Foundation, Inc.+Copyright (C) 2010 Free Software Foundation, Inc. This config.status script is free software; the Free Software Foundation gives unlimited permission to copy, distribute and modify it." @@ -2883,11 +2936,16 @@ while test $# != 0 do case $1 in- --*=*)+ --*=?*) ac_option=`expr "X$1" : 'X\([^=]*\)='` ac_optarg=`expr "X$1" : 'X[^=]*=\(.*\)'` ac_shift=: ;;+ --*=)+ ac_option=`expr "X$1" : 'X\([^=]*\)='`+ ac_optarg=+ ac_shift=:+ ;; *) ac_option=$1 ac_optarg=$2@@ -2909,6 +2967,7 @@ $ac_shift case $ac_optarg in *\'*) ac_optarg=`$as_echo "$ac_optarg" | sed "s/'/'\\\\\\\\''/g"` ;;+ '') as_fn_error $? "missing file argument" ;; esac as_fn_append CONFIG_FILES " '$ac_optarg'" ac_need_defaults=false;;@@ -2919,7 +2978,7 @@ ac_cs_silent=: ;; # This is an error.- -*) as_fn_error "unrecognized option: \`$1'+ -*) as_fn_error $? "unrecognized option: \`$1' Try \`$0 --help' for more information." ;; *) as_fn_append ac_config_targets " $1"@@ -2968,9 +3027,10 @@ for ac_config_target in $ac_config_targets do case $ac_config_target in+ "accelerate-cuda.buildinfo") CONFIG_FILES="$CONFIG_FILES accelerate-cuda.buildinfo" ;; "cubits/accelerate_cuda_shape.h") CONFIG_FILES="$CONFIG_FILES cubits/accelerate_cuda_shape.h" ;; - *) as_fn_error "invalid argument: \`$ac_config_target'" "$LINENO" 5;;+ *) as_fn_error $? "invalid argument: \`$ac_config_target'" "$LINENO" 5;; esac done @@ -2991,9 +3051,10 @@ # after its creation but before its name has been assigned to `$tmp'. $debug || {- tmp=+ tmp= ac_tmp= trap 'exit_status=$?- { test -z "$tmp" || test ! -d "$tmp" || rm -fr "$tmp"; } && exit $exit_status+ : "${ac_tmp:=$tmp}"+ { test ! -d "$ac_tmp" || rm -fr "$ac_tmp"; } && exit $exit_status ' 0 trap 'as_fn_exit 1' 1 2 13 15 }@@ -3001,12 +3062,13 @@ { tmp=`(umask 077 && mktemp -d "./confXXXXXX") 2>/dev/null` &&- test -n "$tmp" && test -d "$tmp"+ test -d "$tmp" } || { tmp=./conf$$-$RANDOM (umask 077 && mkdir "$tmp")-} || as_fn_error "cannot create a temporary directory in ." "$LINENO" 5+} || as_fn_error $? "cannot create a temporary directory in ." "$LINENO" 5+ac_tmp=$tmp # Set up the scripts for CONFIG_FILES section. # No need to generate them if there are no CONFIG_FILES.@@ -3023,12 +3085,12 @@ fi ac_cs_awk_cr=`$AWK 'BEGIN { print "a\rb" }' </dev/null 2>/dev/null` if test "$ac_cs_awk_cr" = "a${ac_cr}b"; then- ac_cs_awk_cr='\r'+ ac_cs_awk_cr='\\r' else ac_cs_awk_cr=$ac_cr fi -echo 'BEGIN {' >"$tmp/subs1.awk" &&+echo 'BEGIN {' >"$ac_tmp/subs1.awk" && _ACEOF @@ -3037,18 +3099,18 @@ echo "$ac_subst_vars" | sed 's/.*/&!$&$ac_delim/' && echo "_ACEOF" } >conf$$subs.sh ||- as_fn_error "could not make $CONFIG_STATUS" "$LINENO" 5-ac_delim_num=`echo "$ac_subst_vars" | grep -c '$'`+ as_fn_error $? "could not make $CONFIG_STATUS" "$LINENO" 5+ac_delim_num=`echo "$ac_subst_vars" | grep -c '^'` ac_delim='%!_!# ' for ac_last_try in false false false false false :; do . ./conf$$subs.sh ||- as_fn_error "could not make $CONFIG_STATUS" "$LINENO" 5+ as_fn_error $? "could not make $CONFIG_STATUS" "$LINENO" 5 ac_delim_n=`sed -n "s/.*$ac_delim\$/X/p" conf$$subs.awk | grep -c X` if test $ac_delim_n = $ac_delim_num; then break elif $ac_last_try; then- as_fn_error "could not make $CONFIG_STATUS" "$LINENO" 5+ as_fn_error $? "could not make $CONFIG_STATUS" "$LINENO" 5 else ac_delim="$ac_delim!$ac_delim _$ac_delim!! " fi@@ -3056,7 +3118,7 @@ rm -f conf$$subs.sh cat >>$CONFIG_STATUS <<_ACEOF || ac_write_fail=1-cat >>"\$tmp/subs1.awk" <<\\_ACAWK &&+cat >>"\$ac_tmp/subs1.awk" <<\\_ACAWK && _ACEOF sed -n ' h@@ -3104,7 +3166,7 @@ rm -f conf$$subs.awk cat >>$CONFIG_STATUS <<_ACEOF || ac_write_fail=1 _ACAWK-cat >>"\$tmp/subs1.awk" <<_ACAWK &&+cat >>"\$ac_tmp/subs1.awk" <<_ACAWK && for (key in S) S_is_set[key] = 1 FS = "" @@ -3136,21 +3198,29 @@ sed "s/$ac_cr\$//; s/$ac_cr/$ac_cs_awk_cr/g" else cat-fi < "$tmp/subs1.awk" > "$tmp/subs.awk" \- || as_fn_error "could not setup config files machinery" "$LINENO" 5+fi < "$ac_tmp/subs1.awk" > "$ac_tmp/subs.awk" \+ || as_fn_error $? "could not setup config files machinery" "$LINENO" 5 _ACEOF -# VPATH may cause trouble with some makes, so we remove $(srcdir),-# ${srcdir} and @srcdir@ from VPATH if srcdir is ".", strip leading and+# VPATH may cause trouble with some makes, so we remove sole $(srcdir),+# ${srcdir} and @srcdir@ entries from VPATH if srcdir is ".", strip leading and # trailing colons and then remove the whole line if VPATH becomes empty # (actually we leave an empty line to preserve line numbers). if test "x$srcdir" = x.; then- ac_vpsub='/^[ ]*VPATH[ ]*=/{-s/:*\$(srcdir):*/:/-s/:*\${srcdir}:*/:/-s/:*@srcdir@:*/:/-s/^\([^=]*=[ ]*\):*/\1/+ ac_vpsub='/^[ ]*VPATH[ ]*=[ ]*/{+h+s///+s/^/:/+s/[ ]*$/:/+s/:\$(srcdir):/:/g+s/:\${srcdir}:/:/g+s/:@srcdir@:/:/g+s/^:*// s/:*$//+x+s/\(=[ ]*\).*/\1/+G+s/\n// s/^[^=]*=[ ]*$// }' fi@@ -3168,7 +3238,7 @@ esac case $ac_mode$ac_tag in :[FHL]*:*);;- :L* | :C*:*) as_fn_error "invalid tag \`$ac_tag'" "$LINENO" 5;;+ :L* | :C*:*) as_fn_error $? "invalid tag \`$ac_tag'" "$LINENO" 5;; :[FH]-) ac_tag=-:-;; :[FH]*) ac_tag=$ac_tag:$ac_tag.in;; esac@@ -3187,7 +3257,7 @@ for ac_f do case $ac_f in- -) ac_f="$tmp/stdin";;+ -) ac_f="$ac_tmp/stdin";; *) # Look for the file first in the build tree, then in the source tree # (if the path is not absolute). The absolute path cannot be DOS-style, # because $ac_f cannot contain `:'.@@ -3196,7 +3266,7 @@ [\\/$]*) false;; *) test -f "$srcdir/$ac_f" && ac_f="$srcdir/$ac_f";; esac ||- as_fn_error "cannot find input file: \`$ac_f'" "$LINENO" 5;;+ as_fn_error 1 "cannot find input file: \`$ac_f'" "$LINENO" 5;; esac case $ac_f in *\'*) ac_f=`$as_echo "$ac_f" | sed "s/'/'\\\\\\\\''/g"`;; esac as_fn_append ac_file_inputs " '$ac_f'"@@ -3222,8 +3292,8 @@ esac case $ac_tag in- *:-:* | *:-) cat >"$tmp/stdin" \- || as_fn_error "could not create $ac_file" "$LINENO" 5 ;;+ *:-:* | *:-) cat >"$ac_tmp/stdin" \+ || as_fn_error $? "could not create $ac_file" "$LINENO" 5 ;; esac ;; esac@@ -3348,23 +3418,24 @@ s&@abs_top_builddir@&$ac_abs_top_builddir&;t t $ac_datarootdir_hack "-eval sed \"\$ac_sed_extra\" "$ac_file_inputs" | $AWK -f "$tmp/subs.awk" >$tmp/out \- || as_fn_error "could not create $ac_file" "$LINENO" 5+eval sed \"\$ac_sed_extra\" "$ac_file_inputs" | $AWK -f "$ac_tmp/subs.awk" \+ >$ac_tmp/out || as_fn_error $? "could not create $ac_file" "$LINENO" 5 test -z "$ac_datarootdir_hack$ac_datarootdir_seen" &&- { ac_out=`sed -n '/\${datarootdir}/p' "$tmp/out"`; test -n "$ac_out"; } &&- { ac_out=`sed -n '/^[ ]*datarootdir[ ]*:*=/p' "$tmp/out"`; test -z "$ac_out"; } &&+ { ac_out=`sed -n '/\${datarootdir}/p' "$ac_tmp/out"`; test -n "$ac_out"; } &&+ { ac_out=`sed -n '/^[ ]*datarootdir[ ]*:*=/p' \+ "$ac_tmp/out"`; test -z "$ac_out"; } && { $as_echo "$as_me:${as_lineno-$LINENO}: WARNING: $ac_file contains a reference to the variable \`datarootdir'-which seems to be undefined. Please make sure it is defined." >&5+which seems to be undefined. Please make sure it is defined" >&5 $as_echo "$as_me: WARNING: $ac_file contains a reference to the variable \`datarootdir'-which seems to be undefined. Please make sure it is defined." >&2;}+which seems to be undefined. Please make sure it is defined" >&2;} - rm -f "$tmp/stdin"+ rm -f "$ac_tmp/stdin" case $ac_file in- -) cat "$tmp/out" && rm -f "$tmp/out";;- *) rm -f "$ac_file" && mv "$tmp/out" "$ac_file";;+ -) cat "$ac_tmp/out" && rm -f "$ac_tmp/out";;+ *) rm -f "$ac_file" && mv "$ac_tmp/out" "$ac_file";; esac \- || as_fn_error "could not create $ac_file" "$LINENO" 5+ || as_fn_error $? "could not create $ac_file" "$LINENO" 5 ;; @@ -3379,7 +3450,7 @@ ac_clean_files=$ac_clean_files_save test $ac_write_fail = 0 ||- as_fn_error "write failure creating $CONFIG_STATUS" "$LINENO" 5+ as_fn_error $? "write failure creating $CONFIG_STATUS" "$LINENO" 5 # configure is writing to config.log, and then calls config.status.@@ -3400,7 +3471,7 @@ exec 5>>config.log # Use ||, not &&, to avoid exiting from the if with $? = 1, which # would make configure fail if this is the last instruction.- $ac_cs_success || as_fn_exit $?+ $ac_cs_success || as_fn_exit 1 fi if test -n "$ac_unrecognized_opts" && test "$enable_option_checking" != no; then { $as_echo "$as_me:${as_lineno-$LINENO}: WARNING: unrecognized options: $ac_unrecognized_opts" >&5
cubits/accelerate_cuda_extras.h view
@@ -17,6 +17,7 @@ #include "accelerate_cuda_shape.h" #include "accelerate_cuda_stencil.h" #include "accelerate_cuda_texture.h"+#include "accelerate_cuda_type.h" #endif
cubits/accelerate_cuda_function.h view
@@ -13,8 +13,8 @@ #ifndef __ACCELERATE_CUDA_FUNCTION_H__ #define __ACCELERATE_CUDA_FUNCTION_H__ -#include <stdint.h> #include <cuda_runtime.h>+#include "accelerate_cuda_type.h" #ifdef __cplusplus @@ -26,16 +26,16 @@ * Left/Right bitwise rotation */ template <typename T>-static __inline__ __device__ T rotateL(const T x, const int32_t i)+static __inline__ __device__ T rotateL(const T x, const Int32 i) {- const int32_t i8 = i & 8 * sizeof(x) - 1;+ const Int32 i8 = i & 8 * sizeof(x) - 1; return i8 == 0 ? x : x << i8 | x >> 8 * sizeof(x) - i8; } template <typename T>-static __inline__ __device__ T rotateR(const T x, const int32_t i)+static __inline__ __device__ T rotateR(const T x, const Int32 i) {- const int32_t i8 = i & 8 * sizeof(x) - 1;+ const Int32 i8 = i & 8 * sizeof(x) - 1; return i8 == 0 ? x : x >> i8 | x << 8 * sizeof(x) - i8; } @@ -63,24 +63,24 @@ * Type coercion */ template <typename T>-static __inline__ __device__ uint32_t reinterpret32(const T x)+static __inline__ __device__ Word32 reinterpret32(const T x) {- union { T a; uint32_t b; } u;+ union { T a; Word32 b; } u; u.a = x; return u.b; } template <>-static __inline__ __device__ uint32_t reinterpret32(const float x)+static __inline__ __device__ Word32 reinterpret32(const float x) { return __float_as_int(x); } template <typename T>-static __inline__ __device__ uint64_t reinterpret64(const T x)+static __inline__ __device__ Word64 reinterpret64(const T x) {- union { T a; uint64_t b; } u;+ union { T a; Word64 b; } u; u.a = x; return u.b;@@ -88,7 +88,7 @@ #if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 130 template <>-static __inline__ __device__ uint64_t reinterpret64(const double x)+static __inline__ __device__ Word64 reinterpret64(const double x) { return __double_as_longlong(x); }@@ -101,36 +101,38 @@ template <typename T> static __inline__ __device__ T atomicCAS32(T* address, T compare, T val) {- union { T a; uint32_t b; } u;+ union { T a; Word32 b; } u; - u.b = atomicCAS((uint32_t*) address, reinterpret32<T>(compare), reinterpret32<T>(val));+ u.b = atomicCAS((Word32*) address, reinterpret32<T>(compare), reinterpret32<T>(val)); return u.a; } +#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 110 template <>-static __inline__ __device__ int32_t atomicCAS32(int32_t* address, int32_t compare, int32_t val)+static __inline__ __device__ Int32 atomicCAS32(Int32* address, Int32 compare, Int32 val) { return atomicCAS(address, compare, val); } template <>-static __inline__ __device__ uint32_t atomicCAS32(uint32_t* address, uint32_t compare, uint32_t val)+static __inline__ __device__ Word32 atomicCAS32(Word32* address, Word32 compare, Word32 val) { return atomicCAS(address, compare, val); }+#endif #if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 120 template <typename T> static __inline__ __device__ T atomicCAS64(T* address, T compare, T val) {- union { T a; unsigned long long int b; } u;+ union { T a; Word64 b; } u; - u.b = atomicCAS((unsigned long long int*) address, reinterpret64<T>(compare), reinterpret64<T>(val));+ u.b = atomicCAS((Word64*) address, reinterpret64<T>(compare), reinterpret64<T>(val)); return u.a; } template <>-static __inline__ __device__ unsigned long long int atomicCAS64(unsigned long long int* address, unsigned long long int compare, unsigned long long int val)+static __inline__ __device__ Word64 atomicCAS64(Word64* address, Word64 compare, Word64 val) { return atomicCAS(address, compare, val); }
cubits/accelerate_cuda_shape.h view
@@ -13,8 +13,8 @@ #ifndef __ACCELERATE_CUDA_SHAPE_H__ #define __ACCELERATE_CUDA_SHAPE_H__ -#include <stdint.h> #include <cuda_runtime.h>+#include "accelerate_cuda_type.h" /* * For the time being, the D.A.A.CUDA.Execute converts shape components from@@ -23,9 +23,9 @@ * future hardware gains better 64-bit support and/or we need to access very * large arrays. *- * typedef int64_t Ix;+ * typedef Int64 Ix; */-typedef int32_t Ix;+typedef Int32 Ix; typedef void* DIM0; typedef Ix DIM1; typedef struct { Ix a1,a0; } DIM2;
cubits/accelerate_cuda_shape.h.in view
@@ -13,8 +13,8 @@ #ifndef __ACCELERATE_CUDA_SHAPE_H__ #define __ACCELERATE_CUDA_SHAPE_H__ -#include <stdint.h> #include <cuda_runtime.h>+#include "accelerate_cuda_type.h" /* * For the time being, the D.A.A.CUDA.Execute converts shape components from@@ -23,9 +23,9 @@ * future hardware gains better 64-bit support and/or we need to access very * large arrays. *- * typedef @hs_int_type@ Ix;+ * typedef @type_hs_int@ Ix; */-typedef int32_t Ix;+typedef Int32 Ix; typedef void* DIM0; typedef Ix DIM1; typedef struct { Ix a1,a0; } DIM2;
cubits/accelerate_cuda_stencil.h view
@@ -53,7 +53,7 @@ static __inline__ __device__ DIM1 clamp(const DIM1 sz, const DIM1 i) { // CUDA-4.0 does not include 64-bit min/max functions for Fermi (?)- return max(0, min((int) i, (int) sz-1));+ return max(0, min(i, sz-1)); } @@ -88,6 +88,6 @@ else return i; } -#endif+#endif // __cplusplus+#endif // __ACCELERATE_CUDA_STENCIL_H__ -#endif
cubits/accelerate_cuda_texture.h view
@@ -9,8 +9,7 @@ * Stability : experimental * * CUDA texture definitions and access functions are defined in terms of- * templates, and hence only available through the C++ interface. Expose some- * dummy wrappers to enable parsing with language-c.+ * templates, and hence only available through the C++ interface. * * We don't have a definition for `Int' or `Word', since the bitwidth of the * Haskell and C types may be different.@@ -25,38 +24,38 @@ #ifndef __ACCELERATE_CUDA_TEXTURE_H__ #define __ACCELERATE_CUDA_TEXTURE_H__ -#include <stdint.h> #include <cuda_runtime.h>+#include "accelerate_cuda_type.h" #if defined(__cplusplus) && defined(__CUDACC__) -typedef texture<uint2, 1> TexWord64;-typedef texture<uint32_t, 1> TexWord32;-typedef texture<uint16_t, 1> TexWord16;-typedef texture<uint8_t, 1> TexWord8;-typedef texture<int2, 1> TexInt64;-typedef texture<int32_t, 1> TexInt32;-typedef texture<int16_t, 1> TexInt16;-typedef texture<int8_t, 1> TexInt8;-typedef texture<float, 1> TexFloat;-typedef texture<char, 1> TexCChar;+typedef texture<uint2, 1> TexWord64;+typedef texture<Word32, 1> TexWord32;+typedef texture<Word16, 1> TexWord16;+typedef texture<Word8, 1> TexWord8;+typedef texture<int2, 1> TexInt64;+typedef texture<Int32, 1> TexInt32;+typedef texture<Int16, 1> TexInt16;+typedef texture<Int8, 1> TexInt8;+typedef texture<float, 1> TexFloat;+typedef texture<char, 1> TexCChar; -static __inline__ __device__ uint8_t indexArray(TexWord8 t, const int x) { return tex1Dfetch(t,x); }-static __inline__ __device__ uint16_t indexArray(TexWord16 t, const int x) { return tex1Dfetch(t,x); }-static __inline__ __device__ uint32_t indexArray(TexWord32 t, const int x) { return tex1Dfetch(t,x); }-static __inline__ __device__ uint64_t indexArray(TexWord64 t, const int x)+static __inline__ __device__ Word8 indexArray(TexWord8 t, const int x) { return tex1Dfetch(t,x); }+static __inline__ __device__ Word16 indexArray(TexWord16 t, const int x) { return tex1Dfetch(t,x); }+static __inline__ __device__ Word32 indexArray(TexWord32 t, const int x) { return tex1Dfetch(t,x); }+static __inline__ __device__ Word64 indexArray(TexWord64 t, const int x) {- union { uint2 x; uint64_t y; } v;+ union { uint2 x; Word64 y; } v; v.x = tex1Dfetch(t,x); return v.y; } -static __inline__ __device__ int8_t indexArray(TexInt8 t, const int x) { return tex1Dfetch(t,x); }-static __inline__ __device__ int16_t indexArray(TexInt16 t, const int x) { return tex1Dfetch(t,x); }-static __inline__ __device__ int32_t indexArray(TexInt32 t, const int x) { return tex1Dfetch(t,x); }-static __inline__ __device__ int64_t indexArray(TexInt64 t, const int x)+static __inline__ __device__ Int8 indexArray(TexInt8 t, const int x) { return tex1Dfetch(t,x); }+static __inline__ __device__ Int16 indexArray(TexInt16 t, const int x) { return tex1Dfetch(t,x); }+static __inline__ __device__ Int32 indexArray(TexInt32 t, const int x) { return tex1Dfetch(t,x); }+static __inline__ __device__ Int64 indexArray(TexInt64 t, const int x) {- union { int2 x; int64_t y; } v;+ union { int2 x; Int64 y; } v; v.x = tex1Dfetch(t,x); return v.y; }@@ -98,35 +97,6 @@ return __hiloint2double(v.y,v.x); } #endif--#else--typedef void* TexWord64;-typedef void* TexWord32;-typedef void* TexWord16;-typedef void* TexWord8;-typedef void* TexInt64;-typedef void* TexInt32;-typedef void* TexInt16;-typedef void* TexInt8;-typedef void* TexCShort;-typedef void* TexCUShort;-typedef void* TexCInt;-typedef void* TexCUInt;-typedef void* TexCLong;-typedef void* TexCULong;-typedef void* TexCLLong;-typedef void* TexCULLong;-typedef void* TexFloat;-typedef void* TexDouble;-typedef void* TexCFloat;-typedef void* TexCDouble;-typedef void* TexCChar;-typedef void* TexCSChar;-typedef void* TexCUChar;--void* indexArray(const void*, const int);-void* indexDArray(const void*, const int); #endif // defined(__cplusplus) && defined(__CUDACC__) #endif // __ACCELERATE_CUDA_TEXTURE_H__
+ cubits/accelerate_cuda_type.h view
@@ -0,0 +1,32 @@+/* -----------------------------------------------------------------------------+ *+ * Module : Type+ * Copyright : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee+ * [2009..2012] Manuel M T Chakravarty, Gabriele Keller, Trevor L. McDonell+ * License : BSD3+ *+ * Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+ * Stability : experimental+ *+ * ---------------------------------------------------------------------------*/++#ifndef __ACCELERATE_CUDA_TYPE_H__+#define __ACCELERATE_CUDA_TYPE_H__++/*+ * The word size on CUDA devices is always 32-bits. If the host is a 64-bit+ * architecture, the system stdint.h implementation can define these incorrectly+ * for device code.+ *+ */+typedef signed char Int8;+typedef short Int16;+typedef int Int32;+typedef long long Int64;++typedef unsigned char Word8;+typedef unsigned short Word16;+typedef unsigned int Word32;+typedef unsigned long long Word64;++#endif