accelerate-llvm-ptx 1.0.0.1 → 1.1.0.0
raw patch · 33 files changed
+2150/−1158 lines, 33 filesdep +deepseqdep +file-embeddep +processdep ~acceleratedep ~accelerate-llvmdep ~base
Dependencies added: deepseq, file-embed, process, template-haskell
Dependency ranges changed: accelerate, accelerate-llvm, base, bytestring, cuda, llvm-hs, llvm-hs-pure
Files
- CHANGELOG.md +37/−0
- Data/Array/Accelerate/LLVM/PTX.hs +251/−39
- Data/Array/Accelerate/LLVM/PTX/Analysis/Device.hs +1/−1
- Data/Array/Accelerate/LLVM/PTX/Analysis/Launch.hs +23/−20
- Data/Array/Accelerate/LLVM/PTX/Array/Prim.hs +5/−0
- Data/Array/Accelerate/LLVM/PTX/Array/Remote.hs +7/−4
- Data/Array/Accelerate/LLVM/PTX/CodeGen.hs +14/−13
- Data/Array/Accelerate/LLVM/PTX/CodeGen/Base.hs +8/−6
- Data/Array/Accelerate/LLVM/PTX/CodeGen/Fold.hs +34/−31
- Data/Array/Accelerate/LLVM/PTX/CodeGen/FoldSeg.hs +30/−27
- Data/Array/Accelerate/LLVM/PTX/CodeGen/Intrinsic.hs +359/−0
- Data/Array/Accelerate/LLVM/PTX/CodeGen/Queue.hs +2/−2
- Data/Array/Accelerate/LLVM/PTX/CodeGen/Scan.hs +26/−23
- Data/Array/Accelerate/LLVM/PTX/Compile.hs +191/−161
- Data/Array/Accelerate/LLVM/PTX/Compile/Cache.hs +40/−0
- Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice.hs +116/−514
- Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice/Load.hs +127/−0
- Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice/TH.hs +184/−0
- Data/Array/Accelerate/LLVM/PTX/Compile/Link.hs +0/−183
- Data/Array/Accelerate/LLVM/PTX/Embed.hs +86/−0
- Data/Array/Accelerate/LLVM/PTX/Execute.hs +73/−81
- Data/Array/Accelerate/LLVM/PTX/Execute/Async.hs +1/−2
- Data/Array/Accelerate/LLVM/PTX/Execute/Event.hs +2/−2
- Data/Array/Accelerate/LLVM/PTX/Execute/Marshal.hs +0/−1
- Data/Array/Accelerate/LLVM/PTX/Execute/Stream/Reservoir.hs +1/−1
- Data/Array/Accelerate/LLVM/PTX/Foreign.hs +4/−2
- Data/Array/Accelerate/LLVM/PTX/Link.hs +153/−0
- Data/Array/Accelerate/LLVM/PTX/Link/Cache.hs +22/−0
- Data/Array/Accelerate/LLVM/PTX/Link/Object.hs +41/−0
- Data/Array/Accelerate/LLVM/PTX/State.hs +8/−5
- Data/Array/Accelerate/LLVM/PTX/Target.hs +29/−24
- README.md +188/−0
- accelerate-llvm-ptx.cabal +87/−16
+ CHANGELOG.md view
@@ -0,0 +1,37 @@+# Change Log++Notable changes to the project will be documented in this file.++The format is based on [Keep a Changelog](http://keepachangelog.com/) and the+project adheres to the [Haskell Package Versioning+Policy (PVP)](https://pvp.haskell.org)++## [1.1.0.0] - 2017-09-21+### Added+ * support for GHC-8.2+ * caching of compilation results ([accelerate-llvm#17])+ * support for ahead-of-time compilation (`runQ` and `runQAsync`)++### Changed+ * generalise `run1*` to polyvariadic `runN*`++### Fixed+ * Fixed synchronisation bug in multidimensional reduction+ ++## [1.0.0.1] - 2017-05-25+### Fixed+ * [#386] (partial fix)++## [1.0.0.0] - 2017-03-31+ * initial release+++[1.1.0.0]: https://github.com/AccelerateHS/accelerate-llvm/compare/1.0.0.0...HEAD+[1.0.0.1]: https://github.com/AccelerateHS/accelerate-llvm/compare/1.0.0.0...1.0.0.1+[1.0.0.0]: https://github.com/AccelerateHS/accelerate-llvm/compare/be7f91295f77434b2103c70aa1cabb6a4f2b09a8...1.0.0.0++[#386]: https://github.com/AccelerateHS/accelerate/issues/386++[accelerate-llvm#17]: https://github.com/AccelerateHS/accelerate-llvm/issues/17+
Data/Array/Accelerate/LLVM/PTX.hs view
@@ -3,6 +3,7 @@ {-# LANGUAGE FlexibleInstances #-} {-# LANGUAGE GADTs #-} {-# LANGUAGE TemplateHaskell #-}+{-# LANGUAGE TypeFamilies #-} {-# LANGUAGE TypeSynonymInstances #-} -- | -- Module : Data.Array.Accelerate.LLVM.PTX@@ -26,6 +27,7 @@ -- * Synchronous execution run, runWith, run1, run1With,+ runN, runNWith, stream, streamWith, -- * Asynchronous execution@@ -34,7 +36,12 @@ runAsync, runAsyncWith, run1Async, run1AsyncWith,+ runNAsync, runNAsyncWith, + -- * Ahead-of-time compilation+ runQ, runQWith,+ runQAsync, runQAsyncWith,+ -- * Execution targets PTX, createTargetForDevice, createTargetFromContext, @@ -44,27 +51,38 @@ ) where -- accelerate-import Data.Array.Accelerate.Array.Sugar ( Arrays )+import Data.Array.Accelerate.AST ( PreOpenAfun(..) )+import Data.Array.Accelerate.Array.Sugar ( Arrays ) import Data.Array.Accelerate.Async-import Data.Array.Accelerate.Debug as Debug+import Data.Array.Accelerate.Debug as Debug import Data.Array.Accelerate.Error-import Data.Array.Accelerate.Smart ( Acc )+import Data.Array.Accelerate.Smart ( Acc ) import Data.Array.Accelerate.Trafo +import Data.Array.Accelerate.LLVM.Execute.Async ( AsyncR(..) )+import Data.Array.Accelerate.LLVM.Execute.Environment ( AvalR(..) ) import Data.Array.Accelerate.LLVM.PTX.Compile+import Data.Array.Accelerate.LLVM.PTX.Embed ( embedOpenAcc ) import Data.Array.Accelerate.LLVM.PTX.Execute+import Data.Array.Accelerate.LLVM.PTX.Execute.Environment ( Aval )+import Data.Array.Accelerate.LLVM.PTX.Link import Data.Array.Accelerate.LLVM.PTX.State import Data.Array.Accelerate.LLVM.PTX.Target-import qualified Data.Array.Accelerate.LLVM.PTX.Context as CT-import qualified Data.Array.Accelerate.LLVM.PTX.Array.Data as AD+import Data.Array.Accelerate.LLVM.State+import qualified Data.Array.Accelerate.LLVM.PTX.Array.Data as AD+import qualified Data.Array.Accelerate.LLVM.PTX.Context as CT+import qualified Data.Array.Accelerate.LLVM.PTX.Execute.Async as E -import Foreign.CUDA.Driver as CUDA ( CUDAException, mallocHostForeignPtr )+import Foreign.CUDA.Driver as CUDA ( CUDAException, mallocHostForeignPtr ) -- standard library+import Data.Typeable import Control.Exception import Control.Monad.Trans import System.IO.Unsafe import Text.Printf+import qualified Language.Haskell.TH as TH+import qualified Language.Haskell.TH.Syntax as TH -- Accelerate: LLVM backend for NVIDIA GPUs@@ -72,12 +90,17 @@ -- | Compile and run a complete embedded array program. ----- Note that it is recommended that you use 'run1' whenever possible.+-- The result is copied back to the host only once the arrays are demanded (or+-- the result is forced to normal form). For results consisting of multiple+-- components (a tuple of arrays or array of tuples) this applies per primitive+-- array. Evaluating the result of 'run' to WHNF will initiate the computation,+-- but does not copy the results back from the device. --+-- /NOTE:/ it is recommended to use 'runN' or 'runQ' whenever possible.+-- run :: Arrays a => Acc a -> a run = runWith defaultTarget - -- | As 'run', but execute using the specified target rather than using the -- default, automatically selected device. --@@ -101,7 +124,6 @@ runAsync :: Arrays a => Acc a -> IO (Async a) runAsync = runAsyncWith defaultTarget - -- | As 'runWith', but execute asynchronously. Be sure not to destroy the context, -- or attempt to attach it to a different host thread, before all outstanding -- operations have completed.@@ -114,28 +136,44 @@ dumpGraph acc evalPTX target $ do acc `seq` dumpSimplStats- exec <- phase "compile" (compileAcc acc)- res <- phase "execute" (executeAcc exec >>= AD.copyToHostLazy)+ build <- phase "compile" (compileAcc acc)+ exec <- phase "link" (linkAcc build)+ res <- phase "execute" (executeAcc exec >>= AD.copyToHostLazy) return res --- | Prepare and execute an embedded array program of one argument.+-- | This is 'runN', specialised to an array program of one argument. --+run1 :: (Arrays a, Arrays b) => (Acc a -> Acc b) -> a -> b+run1 = run1With defaultTarget++-- | As 'run1', but execute using the specified target rather than using the+-- default, automatically selected device.+--+run1With :: (Arrays a, Arrays b) => PTX -> (Acc a -> Acc b) -> a -> b+run1With = runNWith+++-- | Prepare and execute an embedded array program.+-- -- This function can be used to improve performance in cases where the array -- program is constant between invocations, because it enables us to bypass -- front-end conversion stages and move directly to the execution phase. If you -- have a computation applied repeatedly to different input data, use this,--- specifying any changing aspects of the computation via the input parameter.+-- specifying any changing aspects of the computation via the input parameters. -- If the function is only evaluated once, this is equivalent to 'run'. ----- To use 'run1' effectively you must express your program as a function of one--- argument. If your program takes more than one argument, you can use--- 'Data.Array.Accelerate.lift' and 'Data.Array.Accelerate.unlift' to tuple up--- the arguments.+-- In order to use 'runN' you must express your Accelerate program as a function+-- of array terms: ----- At an example, once your program is expressed as a function of one argument,--- instead of the usual:+-- > f :: (Arrays a, Arrays b, ... Arrays c) => Acc a -> Acc b -> ... -> Acc c --+-- This function then returns the compiled version of 'f':+--+-- > runN f :: (Arrays a, Arrays b, ... Arrays c) => a -> b -> ... -> c+--+-- At an example, rather than:+-- -- > step :: Acc (Vector a) -> Acc (Vector b) -- > step = ... -- >@@ -144,59 +182,233 @@ -- -- Instead write: ----- > simulate xs = run1 step xs+-- > simulate = runN step -- -- You can use the debugging options to check whether this is working--- successfully by, for example, observing no output from the @-ddump-cc@ flag--- at the second and subsequent invocations.+-- successfully. For example, running with the @-ddump-phases@ flag should show+-- that the compilation steps only happen once, not on the second and subsequent+-- invocations of 'simulate'. Note that this typically relies on GHC knowing+-- that it can lift out the function returned by 'runN' and reuse it. --+-- As with 'run', the resulting array(s) are only copied back to the host once+-- they are actually demanded (forced to normal form). Thus, splitting a program+-- into multiple 'runN' steps does not imply transferring intermediate+-- computations back and forth between host and device. However note that+-- Accelerate is not able to optimise (fuse) across separate 'runN' invocations.+-- -- See the programs in the 'accelerate-examples' package for examples. ---run1 :: (Arrays a, Arrays b) => (Acc a -> Acc b) -> a -> b-run1 = run1With defaultTarget-+-- See also 'runQ', which compiles the Accelerate program at _Haskell_ compile+-- time, thus eliminating the runtime overhead altogether.+--+runN :: Afunction f => f -> AfunctionR f+runN = runNWith defaultTarget --- | As 'run1', but execute using the specified target rather than using the--- default, automatically selected device.+-- | As 'runN', but execute using the specified target device. ---run1With :: (Arrays a, Arrays b) => PTX -> (Acc a -> Acc b) -> a -> b-run1With = run1' unsafePerformIO+runNWith :: Afunction f => PTX -> f -> AfunctionR f+runNWith target f = exec+ where+ !acc = convertAfunWith config f+ !afun = unsafePerformIO $ do+ dumpGraph acc+ evalPTX target $ do+ build <- phase "compile" (compileAfun acc) >>= dumpStats+ link <- phase "link" (linkAfun build)+ return link+ !exec = go afun (return Aempty) + go :: ExecOpenAfun PTX aenv t -> LLVM PTX (Aval aenv) -> t+ go (Alam l) k = \arrs ->+ let k' = do aenv <- k+ AsyncR _ a <- E.async (AD.useRemoteAsync arrs)+ return (aenv `Apush` a)+ in go l k'+ go (Abody b) k = unsafePerformIO . phase "execute" . evalPTX target $ do+ aenv <- k+ r <- E.async (executeOpenAcc b aenv)+ AD.copyToHostLazy =<< E.get r + -- | As 'run1', but the computation is executed asynchronously. -- run1Async :: (Arrays a, Arrays b) => (Acc a -> Acc b) -> a -> IO (Async b) run1Async = run1AsyncWith defaultTarget - -- | As 'run1With', but execute asynchronously. -- run1AsyncWith :: (Arrays a, Arrays b) => PTX -> (Acc a -> Acc b) -> a -> IO (Async b)-run1AsyncWith = run1' asyncBound+run1AsyncWith = runNAsyncWith -run1' :: (Arrays a, Arrays b) => (IO b -> c) -> PTX -> (Acc a -> Acc b) -> a -> c-run1' using target f = \a -> using (execute a)++-- | As 'runN', but execute asynchronously.+--+runNAsync :: (Afunction f, RunAsync r, AfunctionR f ~ RunAsyncR r) => f -> r+runNAsync = runNAsyncWith defaultTarget++-- | As 'runNWith', but execute asynchronously.+--+runNAsyncWith :: (Afunction f, RunAsync r, AfunctionR f ~ RunAsyncR r) => PTX -> f -> r+runNAsyncWith target f = runAsync' target afun (return Aempty) where- !acc = convertAfunWith config f- !afun = unsafePerformIO $ do- dumpGraph acc- phase "compile" (evalPTX target (compileAfun acc)) >>= dumpStats- execute a = phase "execute" (evalPTX target (executeAfun1 afun a >>= AD.copyToHostLazy))+ !acc = convertAfunWith config f+ !afun = unsafePerformIO $ do+ dumpGraph acc+ evalPTX target $ do+ build <- phase "compile" (compileAfun acc) >>= dumpStats+ exec <- phase "link" (linkAfun build)+ return exec +class RunAsync f where+ type RunAsyncR f+ runAsync' :: PTX -> ExecOpenAfun PTX aenv (RunAsyncR f) -> LLVM PTX (Aval aenv) -> f +instance RunAsync b => RunAsync (a -> b) where+ type RunAsyncR (a -> b) = a -> RunAsyncR b+ runAsync' _ Abody{} _ _ = error "runAsync: function oversaturated"+ runAsync' target (Alam l) k arrs =+ let k' = do aenv <- k+ AsyncR _ a <- E.async (AD.useRemoteAsync arrs)+ return (aenv `Apush` a)+ in runAsync' target l k'++instance RunAsync (IO (Async b)) where+ type RunAsyncR (IO (Async b)) = b+ runAsync' _ Alam{} _ = error "runAsync: function not fully applied"+ runAsync' target (Abody b) k = asyncBound . phase "execute" . evalPTX target $ do+ aenv <- k+ r <- E.async (executeOpenAcc b aenv)+ AD.copyToHostLazy =<< E.get r++ -- | Stream a lazily read list of input arrays through the given program, -- collecting results as we go. -- stream :: (Arrays a, Arrays b) => (Acc a -> Acc b) -> [a] -> [b] stream = streamWith defaultTarget - -- | As 'stream', but execute using the specified target. -- streamWith :: (Arrays a, Arrays b) => PTX -> (Acc a -> Acc b) -> [a] -> [b] streamWith target f arrs = map go arrs where !go = run1With target f+++-- | Ahead-of-time compilation for an embedded array program.+--+-- This function will generate, compile, and link into the final executable,+-- code to execute the given Accelerate computation /at Haskell compile time/.+-- This eliminates any runtime overhead associated with the other @run*@+-- operations. The generated code will be compiled for the current (default) GPU+-- architecture.+--+-- Since the Accelerate program will be generated at Haskell compile time,+-- construction of the Accelerate program, in particular via meta-programming,+-- will be limited to operations available to that phase. Also note that any+-- arrays which are embedded into the program via 'Data.Array.Accelerate.use'+-- will be stored as part of the final executable.+--+-- Usage of this function in your program is similar to that of 'runN'. First,+-- express your Accelerate program as a function of array terms:+--+-- > f :: (Arrays a, Arrays b, ... Arrays c) => Acc a -> Acc b -> ... -> Acc c+--+-- This function then returns a compiled version of @f@ as a Template Haskell+-- splice, to be added into your program at Haskell compile time:+--+-- > {-# LANGUAGE TemplateHaskell #-}+-- >+-- > f' :: a -> b -> ... -> c+-- > f' = $( runQ f )+--+-- Note that at the splice point the usage of @f@ must monomorphic; i.e. the+-- types @a@, @b@ and @c@ must be at some known concrete type.+--+-- See the <https://github.com/tmcdonell/lulesh-accelerate lulesh-accelerate>+-- project for an example.+--+-- [/Note:/]+--+-- Due to <https://ghc.haskell.org/trac/ghc/ticket/13587 GHC#13587>, this+-- currently must be as an /untyped/ splice.+--+-- The correct type of this function is similar to that of 'runN':+--+-- > runQ :: Afunction f => f -> Q (TExp (AfunctionR f))+--+-- @since 1.1.0.0+--+runQ :: Afunction f => f -> TH.ExpQ+runQ = runQ' [| unsafePerformIO |] [| defaultTarget |]++-- | Ahead-of-time analogue of 'runNWith'. See 'runQ' for more information.+--+-- /NOTE:/ The supplied (at runtime) target must be compatible with the+-- architecture that this function was compiled for (the 'defaultTarget' of the+-- compiling machine). Running on a device with the same compute capability is+-- best, but this should also be forward compatible to newer architectures.+--+-- The correct type of this function is:+--+-- > runQWith :: Afunction f => f -> Q (TExp (PTX -> AfunctionR f))+--+-- @since 1.1.0.0+--+runQWith :: Afunction f => f -> TH.ExpQ+runQWith f = do+ target <- TH.newName "target"+ TH.lamE [TH.varP target] (runQ' [| unsafePerformIO |] (TH.varE target) f)+++-- | Ahead-of-time analogue of 'runNAsync'. See 'runQ' for more information.+--+-- The correct type of this function is:+--+-- > runQAsync :: (Afunction f, RunAsync r, AfunctionR f ~ RunAsyncR r) => f -> Q (TExp r)+--+-- @since 1.1.0.0+--+runQAsync :: Afunction f => f -> TH.ExpQ+runQAsync = runQ' [| async |] [| defaultTarget |]++-- | Ahead-of-time analogue of 'runNAsyncWith'. See 'runQWith' for more information.+--+-- The correct type of this function is:+--+-- > runQAsyncWith :: (Afunction f, RunAsync r, AfunctionR f ~ RunAsyncR r) => f -> Q (TExp (PTX -> r))+--+-- @since 1.1.0.0+--+runQAsyncWith :: Afunction f => f -> TH.ExpQ+runQAsyncWith f = do+ target <- TH.newName "target"+ TH.lamE [TH.varP target] (runQ' [| async |] (TH.varE target) f)+++runQ' :: Afunction f => TH.ExpQ -> TH.ExpQ -> f -> TH.ExpQ+runQ' using target f = do+ afun <- let acc = convertAfunWith config f+ in TH.runIO $ do+ dumpGraph acc+ evalPTX defaultTarget $+ phase "compile" (compileAfun acc) >>= dumpStats+ let+ go :: Typeable aenv => CompiledOpenAfun PTX aenv t -> [TH.PatQ] -> [TH.ExpQ] -> [TH.StmtQ] -> TH.ExpQ+ go (Alam lam) xs as stmts = do+ x <- TH.newName "x" -- lambda bound variable+ a <- TH.newName "a" -- local array name+ s <- TH.bindS (TH.conP 'AsyncR [TH.wildP, TH.varP a]) [| E.async (AD.useRemoteAsync $(TH.varE x)) |]+ go lam (TH.varP x : xs) (TH.varE a : as) (return s : stmts)++ go (Abody body) xs as stmts =+ let aenv = foldr (\a gamma -> [| $gamma `Apush` $a |] ) [| Aempty |] as+ eval = TH.noBindS [| AD.copyToHostLazy =<< E.get =<< E.async (executeOpenAcc $(TH.unTypeQ (embedOpenAcc defaultTarget body)) $aenv) |]+ in+ TH.lamE (reverse xs) [| $using . phase "execute" . evalPTX $target $+ $(TH.doE (reverse (eval : stmts))) |]+ --+ go afun [] [] [] -- How the Accelerate program should be evaluated.
Data/Array/Accelerate/LLVM/PTX/Analysis/Device.hs view
@@ -32,7 +32,7 @@ selectBestDevice = do dev <- mapM CUDA.device . enumFromTo 0 . subtract 1 =<< CUDA.count prop <- mapM CUDA.props dev- return . head . sortBy (flip cmp `on` snd) $ zip dev prop+ return . minimumBy (flip cmp `on` snd) $ zip dev prop where compute = computeCapability flops d = multiProcessorCount d * coresPerMultiProcessor d * clockRate d
Data/Array/Accelerate/LLVM/PTX/Analysis/Launch.hs view
@@ -1,4 +1,5 @@-{-# LANGUAGE RecordWildCards #-}+{-# LANGUAGE QuasiQuotes #-}+{-# LANGUAGE TemplateHaskell #-} -- | -- Module : Data.Array.Accelerate.LLVM.PTX.Analysis.Launch -- Copyright : [2008..2017] Manuel M T Chakravarty, Gabriele Keller@@ -14,29 +15,26 @@ DeviceProperties, Occupancy, LaunchConfig, simpleLaunchConfig, launchConfig,- multipleOf,+ multipleOf, multipleOfQ, ) where --- library import Foreign.CUDA.Analysis as CUDA+import Language.Haskell.TH --- Kernel annotation for the PTX backend consists of the launch configuration------ data instance KernelAnn PTX = ANN_PTX LaunchConfig- -- | Given information about the resource usage of the compiled kernel, -- determine the optimum launch parameters. -- type LaunchConfig- = Int -- maximum #threads per block- -> Int -- #registers per thread- -> Int -- #bytes of static shared memory+ = Int -- maximum #threads per block+ -> Int -- #registers per thread+ -> Int -- #bytes of static shared memory -> ( Occupancy- , Int -- thread block size- , Int -> Int -- grid size required to process the given input size- , Int -- #bytes dynamic shared memory+ , Int -- thread block size+ , Int -> Int -- grid size required to process the given input size+ , Int -- #bytes dynamic shared memory+ , Q (TExp (Int -> Int)) ) -- | Analytics for a simple kernel which requires no additional shared memory or@@ -44,29 +42,34 @@ -- size, in increments of a single warp, with the highest occupancy is used. -- simpleLaunchConfig :: DeviceProperties -> LaunchConfig-simpleLaunchConfig dev = launchConfig dev (decWarp dev) (const 0) multipleOf+simpleLaunchConfig dev = launchConfig dev (decWarp dev) (const 0) multipleOf multipleOfQ -- | Determine the optimal kernel launch configuration for a kernel. -- launchConfig- :: DeviceProperties -- ^ Device architecture to optimise for- -> [Int] -- ^ Thread block sizes to consider- -> (Int -> Int) -- ^ Shared memory (#bytes) as a function of thread block size- -> (Int -> Int -> Int) -- ^ Determine grid size for input size 'n' (first arg) over thread blocks of size 'm' (second arg)+ :: DeviceProperties -- ^ Device architecture to optimise for+ -> [Int] -- ^ Thread block sizes to consider+ -> (Int -> Int) -- ^ Shared memory (#bytes) as a function of thread block size+ -> (Int -> Int -> Int) -- ^ Determine grid size for input size 'n' (first arg) over thread blocks of size 'm' (second arg)+ -> Q (TExp (Int -> Int -> Int)) -> LaunchConfig-launchConfig dev candidates dynamic_smem grid_size maxThreads registers static_smem =+launchConfig dev candidates dynamic_smem grid_size grid_sizeQ maxThreads registers static_smem = let (cta, occ) = optimalBlockSizeOf dev (filter (<= maxThreads) candidates) (const registers) smem maxGrid = multiProcessorCount dev * activeThreadBlocks occ grid n = maxGrid `min` grid_size n cta smem n = static_smem + dynamic_smem n+ gridQ = [|| \n -> (maxGrid::Int) `min` $$grid_sizeQ (n::Int) (cta::Int) ||] in- ( occ, cta, grid, dynamic_smem cta )+ ( occ, cta, grid, dynamic_smem cta, gridQ ) -- | The next highest multiple of 'y' from 'x'. -- multipleOf :: Int -> Int -> Int multipleOf x y = ((x + y - 1) `quot` y)++multipleOfQ :: Q (TExp (Int -> Int -> Int))+multipleOfQ = [|| multipleOf ||]
Data/Array/Accelerate/LLVM/PTX/Array/Prim.hs view
@@ -131,6 +131,7 @@ nonblocking stream $ transfer "pokeArray" bytes (Just st) $ CUDA.pokeArrayAsync n src dst (Just st) liftIO (touchLifetime stream)+ liftIO (Debug.didCopyBytesToRemote (fromIntegral bytes)) {-# INLINEABLE pokeArrayR #-}@@ -163,6 +164,7 @@ transfer "pokeArray" bytes (Just st) $ CUDA.pokeArrayAsync n (src `CUDA.plusHostPtr` offset) (dst `CUDA.plusDevPtr` offset) (Just st) liftIO (touchLifetime stream)+ liftIO (Debug.didCopyBytesToRemote (fromIntegral bytes)) -- | Read a single element from an array at a given row-major index@@ -179,6 +181,7 @@ bracket (CUDA.mallocHostArray [] 1) CUDA.freeHost $ \dst -> do let !st = unsafeGetValue stream message $ "indexArray: " ++ showBytes (sizeOf (undefined::a))+ Debug.didCopyBytesFromRemote (fromIntegral (sizeOf (undefined::a))) CUDA.peekArrayAsync 1 (src `CUDA.advanceDevPtr` i) dst (Just st) CUDA.block st touchLifetime stream@@ -213,6 +216,7 @@ nonblocking stream $ transfer "peekArray" bytes (Just st) $ CUDA.peekArrayAsync n src dst (Just st) liftIO (touchLifetime stream)+ liftIO (Debug.didCopyBytesFromRemote (fromIntegral bytes)) {-# INLINEABLE peekArrayR #-} peekArrayR@@ -244,6 +248,7 @@ transfer "peekArray" bytes (Just st) $ CUDA.peekArrayAsync n (src `CUDA.plusDevPtr` offset) (dst `CUDA.plusHostPtr` offset) (Just st) liftIO (touchLifetime stream)+ liftIO (Debug.didCopyBytesFromRemote (fromIntegral bytes)) -- | Copy data between arrays in the same context
Data/Array/Accelerate/LLVM/PTX/Array/Remote.hs view
@@ -57,8 +57,9 @@ | otherwise = liftIO $ do ep <- try (CUDA.mallocArray n) case ep of- Right p -> return (Just p)- Left (ExitCode OutOfMemory) -> return Nothing+ Right p -> do liftIO (Debug.didAllocateBytesRemote (fromIntegral n))+ return (Just p)+ Left (ExitCode OutOfMemory) -> do return Nothing Left e -> do message ("malloc failed with error: " ++ show e) throwIO e @@ -67,7 +68,8 @@ dst = CUDA.HostPtr (ptrsOfArrayData ad) in blocking $ \stream ->- withLifetime stream $ \st ->+ withLifetime stream $ \st -> do+ Debug.didCopyBytesFromRemote (fromIntegral bytes) transfer "peekRemote" bytes (Just st) $ CUDA.peekArrayAsync n src dst (Just st) pokeRemote n dst ad =@@ -75,7 +77,8 @@ src = CUDA.HostPtr (ptrsOfArrayData ad) in blocking $ \stream ->- withLifetime stream $ \st ->+ withLifetime stream $ \st -> do+ Debug.didCopyBytesToRemote (fromIntegral bytes) transfer "pokeRemote" bytes (Just st) $ CUDA.pokeArrayAsync n src dst (Just st) castRemotePtr _ = CUDA.castDevPtr
Data/Array/Accelerate/LLVM/PTX/CodeGen.hs view
@@ -23,6 +23,7 @@ import Data.Array.Accelerate.LLVM.PTX.CodeGen.Fold import Data.Array.Accelerate.LLVM.PTX.CodeGen.FoldSeg import Data.Array.Accelerate.LLVM.PTX.CodeGen.Generate+import Data.Array.Accelerate.LLVM.PTX.CodeGen.Intrinsic () import Data.Array.Accelerate.LLVM.PTX.CodeGen.Map import Data.Array.Accelerate.LLVM.PTX.CodeGen.Permute import Data.Array.Accelerate.LLVM.PTX.CodeGen.Scan@@ -30,17 +31,17 @@ instance Skeleton PTX where- map = mkMap- generate = mkGenerate- fold = mkFold- fold1 = mkFold1- foldSeg = mkFoldSeg- fold1Seg = mkFold1Seg- scanl = mkScanl- scanl1 = mkScanl1- scanl' = mkScanl'- scanr = mkScanr- scanr1 = mkScanr1- scanr' = mkScanr'- permute = mkPermute+ map ptx _ = mkMap ptx+ generate ptx _ = mkGenerate ptx+ fold ptx _ = mkFold ptx+ fold1 ptx _ = mkFold1 ptx+ foldSeg ptx _ = mkFoldSeg ptx+ fold1Seg ptx _ = mkFold1Seg ptx+ scanl ptx _ = mkScanl ptx+ scanl1 ptx _ = mkScanl1 ptx+ scanl' ptx _ = mkScanl' ptx+ scanr ptx _ = mkScanr ptx+ scanr1 ptx _ = mkScanr1 ptx+ scanr' ptx _ = mkScanr' ptx+ permute ptx _ = mkPermute ptx
Data/Array/Accelerate/LLVM/PTX/CodeGen/Base.hs view
@@ -55,6 +55,7 @@ import LLVM.AST.Type.Operand import LLVM.AST.Type.Representation import qualified LLVM.AST.Global as LLVM+import qualified LLVM.AST.Constant as LLVM hiding ( type' ) import qualified LLVM.AST.Linkage as LLVM import qualified LLVM.AST.Name as LLVM import qualified LLVM.AST.Type as LLVM@@ -82,6 +83,7 @@ -- standard library import Control.Applicative import Control.Monad ( void )+import Data.String import Text.Printf import Prelude as P @@ -242,7 +244,7 @@ _ -> $internalError "atomicAdd" "unexpected operand type" t_ret = PrimType (ScalarPrimType t_val)- fun = Label $ printf "llvm.nvvm.atomic.load.add.f%d.p%df%d" width addrspace width+ fun = fromString $ printf "llvm.nvvm.atomic.load.add.f%d.p%df%d" width addrspace width in void $ call (Lam t_addr addr (Lam (ScalarPrimType t_val) val (Body t_ret fun))) [NoUnwind] @@ -283,7 +285,7 @@ declare $ LLVM.globalVariableDefaults { LLVM.addrSpace = sharedMemAddrSpace , LLVM.type' = LLVM.ArrayType n (downcast t)- , LLVM.linkage = LLVM.Internal+ , LLVM.linkage = LLVM.External , LLVM.name = downcast nm , LLVM.alignment = 4 `P.max` P.fromIntegral (sizeOf tt) }@@ -312,7 +314,7 @@ , LLVM.name = LLVM.Name "__shared__" , LLVM.alignment = 4 }- return $ ConstantOperand $ GlobalReference type' "__shared__"+ return $ ConstantOperand $ GlobalReference (PrimType (PtrPrimType (ArrayType 0 scalarType) sharedMemAddrSpace)) "__shared__" -- Declared a new dynamically allocated array in the __shared__ memory space@@ -390,9 +392,9 @@ _ <- kernel code <- createBlocks addMetadata "nvvm.annotations"- [ Just . MetadataOperand $ ConstantOperand (GlobalReference VoidType (Name l))- , Just . MetadataStringOperand $ "kernel"- , Just . MetadataOperand $ scalar scalarType (1::Int)+ [ Just . MetadataConstantOperand $ LLVM.GlobalReference (LLVM.PointerType (LLVM.FunctionType LLVM.VoidType [ t | LLVM.Parameter t _ _ <- param ] False) (AddrSpace 0)) (LLVM.Name l)+ , Just . MetadataStringOperand $ "kernel"+ , Just . MetadataConstantOperand $ LLVM.Int 32 1 ] return $ Kernel { kernelMetadata = KM_PTX config
Data/Array/Accelerate/LLVM/PTX/CodeGen/Fold.hs view
@@ -3,6 +3,7 @@ {-# LANGUAGE RebindableSyntax #-} {-# LANGUAGE RecordWildCards #-} {-# LANGUAGE ScopedTypeVariables #-}+{-# LANGUAGE TemplateHaskell #-} {-# LANGUAGE TypeOperators #-} {-# LANGUAGE ViewPatterns #-} -- |@@ -150,12 +151,12 @@ (arrOut, paramOut) = mutableArray ("out" :: Name (Scalar e)) paramEnv = envParam aenv --- config = launchConfig dev (CUDA.incWarp dev) smem multipleOf+ config = launchConfig dev (CUDA.incWarp dev) smem multipleOf multipleOfQ smem n = warps * (1 + per_warp) * bytes where ws = CUDA.warpSize dev- warps = n `div` ws- per_warp = ws + ws `div` 2+ warps = n `P.quot` ws+ per_warp = ws + ws `P.quot` 2 bytes = sizeOf (eltType (undefined :: e)) in makeOpenAccWith config "foldAllS" (paramGang ++ paramOut ++ paramEnv) $ do@@ -201,12 +202,12 @@ (arrTmp, paramTmp) = mutableArray ("tmp" :: Name (Vector e)) paramEnv = envParam aenv --- config = launchConfig dev (CUDA.incWarp dev) smem const+ config = launchConfig dev (CUDA.incWarp dev) smem const [|| const ||] smem n = warps * (1 + per_warp) * bytes where ws = CUDA.warpSize dev- warps = n `div` ws- per_warp = ws + ws `div` 2+ warps = n `P.quot` ws+ per_warp = ws + ws `P.quot` 2 bytes = sizeOf (eltType (undefined :: e)) in makeOpenAccWith config "foldAllM1" (paramGang ++ paramTmp ++ paramEnv) $ do@@ -255,12 +256,12 @@ (arrOut, paramOut) = mutableArray ("out" :: Name (Vector e)) paramEnv = envParam aenv --- config = launchConfig dev (CUDA.incWarp dev) smem const+ config = launchConfig dev (CUDA.incWarp dev) smem const [|| const ||] smem n = warps * (1 + per_warp) * bytes where ws = CUDA.warpSize dev- warps = n `div` ws- per_warp = ws + ws `div` 2+ warps = n `P.quot` ws+ per_warp = ws + ws `P.quot` 2 bytes = sizeOf (eltType (undefined :: e)) in makeOpenAccWith config "foldAllM2" (paramGang ++ paramTmp ++ paramOut ++ paramEnv) $ do@@ -320,12 +321,12 @@ (arrOut, paramOut) = mutableArray ("out" :: Name (Array sh e)) paramEnv = envParam aenv --- config = launchConfig dev (CUDA.incWarp dev) smem const+ config = launchConfig dev (CUDA.incWarp dev) smem const [|| const ||] smem n = warps * (1 + per_warp) * bytes where ws = CUDA.warpSize dev- warps = n `div` ws- per_warp = ws + ws `div` 2+ warps = n `P.quot` ws+ per_warp = ws + ws `P.quot` 2 bytes = sizeOf (eltType (undefined :: e)) in makeOpenAccWith config "fold" (paramGang ++ paramOut ++ paramEnv) $ do@@ -365,25 +366,27 @@ __syncthreads -- Threads cooperatively reduce this stripe of the input- i <- A.add numType offset tid- i' <- A.fromIntegral integralType numType i- valid <- A.sub numType to offset- r' <- if A.gte scalarType valid bd- -- All threads of the block are valid, so we can avoid- -- bounds checks.- then do- x <- app1 delayedLinearIndex i'- reduceBlockSMem dev combine Nothing x+ i <- A.add numType offset tid+ v' <- A.sub numType to offset+ r' <- if A.gte scalarType v' bd+ -- All threads of the block are valid, so we can avoid+ -- bounds checks.+ then do+ x <- app1 delayedLinearIndex =<< A.fromIntegral integralType numType i+ y <- reduceBlockSMem dev combine Nothing x+ return y - -- Otherwise we require bounds checks when reading the- -- input and during the reduction.- else- if A.lt scalarType i to- then do- x <- app1 delayedLinearIndex i'- reduceBlockSMem dev combine (Just valid) x- else- return r+ -- Otherwise, we require bounds checks when reading the input+ -- and during the reduction. Note that even though only the+ -- valid threads will contribute useful work in the+ -- reduction, we must still have all threads enter the+ -- reduction procedure to avoid synchronisation divergence.+ else do+ x <- if A.lt scalarType i to+ then app1 delayedLinearIndex =<< A.fromIntegral integralType numType i+ else return r+ y <- reduceBlockSMem dev combine (Just v') x+ return y if A.eq scalarType tid (lift 0) then app2 combine r r'@@ -436,7 +439,7 @@ -- Temporary storage required for each warp bytes = sizeOf (eltType (undefined::e))- warp_smem_elems = CUDA.warpSize dev + (CUDA.warpSize dev `div` 2)+ warp_smem_elems = CUDA.warpSize dev + (CUDA.warpSize dev `P.quot` 2) -- Step 1: Reduction in every warp --
Data/Array/Accelerate/LLVM/PTX/CodeGen/FoldSeg.hs view
@@ -1,8 +1,8 @@ {-# LANGUAGE GADTs #-} {-# LANGUAGE OverloadedStrings #-} {-# LANGUAGE RebindableSyntax #-}-{-# LANGUAGE RecordWildCards #-} {-# LANGUAGE ScopedTypeVariables #-}+{-# LANGUAGE TemplateHaskell #-} {-# LANGUAGE TypeOperators #-} {-# LANGUAGE ViewPatterns #-} -- |@@ -107,12 +107,12 @@ (arrOut, paramOut) = mutableArray ("out" :: Name (Array (sh :. Int) e)) paramEnv = envParam aenv --- config = launchConfig dev (CUDA.decWarp dev) dsmem const+ config = launchConfig dev (CUDA.decWarp dev) dsmem const [|| const ||] dsmem n = warps * (1 + per_warp) * bytes where ws = CUDA.warpSize dev- warps = n `div` ws- per_warp = ws + ws `div` 2+ warps = n `P.quot` ws+ per_warp = ws + ws `P.quot` 2 bytes = sizeOf (eltType (undefined :: e)) in makeOpenAccWith config "foldSeg_block" (paramGang ++ paramOut ++ paramEnv) $ do@@ -172,7 +172,8 @@ 0 -> return (A.pair u v) _ -> do q <- A.quot integralType s ss a <- A.mul numType q sz- A.pair <$> A.add numType u a <*> A.add numType v a+ A.pair <$> A.add numType u a+ <*> A.add numType v a void $ if A.eq scalarType inf sup@@ -196,7 +197,7 @@ -- those threads die and will not participate in the computation of -- _any_ further segment. I'm not sure if this is a CUDA oddity -- (e.g. we must have all threads convergent on __syncthreads) or- -- a bug in NVPTX.+ -- a bug in NVPTX / ptxas. -- bd <- blockDim i0 <- A.add numType inf tid@@ -228,17 +229,20 @@ -- All threads in the block are in bounds, so we -- can avoid bounds checks. then do- x' <- app1 (delayedLinearIndex arr) =<< A.fromIntegral integralType numType i'- reduceBlockSMem dev combine Nothing x'+ x <- app1 (delayedLinearIndex arr) =<< A.fromIntegral integralType numType i'+ y <- reduceBlockSMem dev combine Nothing x+ return y - -- Not all threads are valid.- else- if A.lt scalarType i' sup- then do- x' <- app1 (delayedLinearIndex arr) =<< A.fromIntegral integralType numType i'- reduceBlockSMem dev combine (Just v') x'- else- return r+ -- Not all threads are valid. Note that we still+ -- have all threads enter the reduction procedure+ -- to avoid thread divergence on synchronisation+ -- points, similar to the above NOTE.+ else do+ x <- if A.lt scalarType i' sup+ then app1 (delayedLinearIndex arr) =<< A.fromIntegral integralType numType i'+ else return r+ y <- reduceBlockSMem dev combine (Just v') x+ return y -- first thread incorporates the result from the previous -- iteration@@ -282,15 +286,16 @@ (arrOut, paramOut) = mutableArray ("out" :: Name (Array (sh :. Int) e)) paramEnv = envParam aenv --- config = launchConfig dev (CUDA.decWarp dev) dsmem grid+ config = launchConfig dev (CUDA.decWarp dev) dsmem grid gridQ dsmem n = warps * (2 + per_warp_elems) * bytes where- warps = n `div` ws+ warps = n `P.quot` ws --- grid n m = multipleOf n (m `div` ws)+ grid n m = multipleOf n (m `P.quot` ws)+ gridQ = [|| \n m -> $$multipleOfQ n (m `P.quot` ws) ||] -- per_warp_bytes = per_warp_elems * bytes- per_warp_elems = ws + (ws `div` 2)+ per_warp_elems = ws + (ws `P.quot` 2) ws = CUDA.warpSize dev bytes = sizeOf (eltType (undefined :: e)) @@ -428,13 +433,11 @@ return y else do- if A.lt scalarType i' sup- then do- x <- app1 (delayedLinearIndex arr) =<< A.fromIntegral integralType numType i'- y <- reduceWarpSMem dev combine smem (Just v') x- return y- else- return r+ x <- if A.lt scalarType i' sup+ then app1 (delayedLinearIndex arr) =<< A.fromIntegral integralType numType i'+ else return r+ y <- reduceWarpSMem dev combine smem (Just v') x+ return y -- The first lane incorporates the result from the previous -- iteration
+ Data/Array/Accelerate/LLVM/PTX/CodeGen/Intrinsic.hs view
@@ -0,0 +1,359 @@+{-# LANGUAGE OverloadedStrings #-}+{-# OPTIONS_GHC -fno-warn-orphans #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.CodeGen.Intrinsic+-- Copyright : [2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.CodeGen.Intrinsic ( )+ where++import LLVM.AST.Type.Name ( Label(..) )++import Data.Array.Accelerate.LLVM.CodeGen.Intrinsic+import Data.Array.Accelerate.LLVM.PTX.Target++import Data.ByteString.Short ( ShortByteString )+import Data.HashMap.Strict ( HashMap )+import Data.Monoid+import qualified Data.HashMap.Strict as HashMap+++instance Intrinsic PTX where+ intrinsicForTarget _ = libdeviceIndex++-- The list of functions implemented by libdevice. These are all more-or-less+-- named consistently based on the standard mathematical functions they+-- implement, with the "__nv_" prefix stripped.+--+libdeviceIndex :: HashMap ShortByteString Label+libdeviceIndex =+ let nv base = (base, Label $ "__nv_" <> base)+ in+ HashMap.fromList $ map nv+ [ "abs"+ , "acos"+ , "acosf"+ , "acosh"+ , "acoshf"+ , "asin"+ , "asinf"+ , "asinh"+ , "asinhf"+ , "atan"+ , "atan2"+ , "atan2f"+ , "atanf"+ , "atanh"+ , "atanhf"+ , "brev"+ , "brevll"+ , "byte_perm"+ , "cbrt"+ , "cbrtf"+ , "ceil"+ , "ceilf"+ , "clz"+ , "clzll"+ , "copysign"+ , "copysignf"+ , "cos"+ , "cosf"+ , "cosh"+ , "coshf"+ , "cospi"+ , "cospif"+ , "dadd_rd"+ , "dadd_rn"+ , "dadd_ru"+ , "dadd_rz"+ , "ddiv_rd"+ , "ddiv_rn"+ , "ddiv_ru"+ , "ddiv_rz"+ , "dmul_rd"+ , "dmul_rn"+ , "dmul_ru"+ , "dmul_rz"+ , "double2float_rd"+ , "double2float_rn"+ , "double2float_ru"+ , "double2float_rz"+ , "double2hiint"+ , "double2int_rd"+ , "double2int_rn"+ , "double2int_ru"+ , "double2int_rz"+ , "double2ll_rd"+ , "double2ll_rn"+ , "double2ll_ru"+ , "double2ll_rz"+ , "double2loint"+ , "double2uint_rd"+ , "double2uint_rn"+ , "double2uint_ru"+ , "double2uint_rz"+ , "double2ull_rd"+ , "double2ull_rn"+ , "double2ull_ru"+ , "double2ull_rz"+ , "double_as_longlong"+ , "drcp_rd"+ , "drcp_rn"+ , "drcp_ru"+ , "drcp_rz"+ , "dsqrt_rd"+ , "dsqrt_rn"+ , "dsqrt_ru"+ , "dsqrt_rz"+ , "erf"+ , "erfc"+ , "erfcf"+ , "erfcinv"+ , "erfcinvf"+ , "erfcx"+ , "erfcxf"+ , "erff"+ , "erfinv"+ , "erfinvf"+ , "exp"+ , "exp10"+ , "exp10f"+ , "exp2"+ , "exp2f"+ , "expf"+ , "expm1"+ , "expm1f"+ , "fabs"+ , "fabsf"+ , "fadd_rd"+ , "fadd_rn"+ , "fadd_ru"+ , "fadd_rz"+ , "fast_cosf"+ , "fast_exp10f"+ , "fast_expf"+ , "fast_fdividef"+ , "fast_log10f"+ , "fast_log2f"+ , "fast_logf"+ , "fast_powf"+ , "fast_sincosf"+ , "fast_sinf"+ , "fast_tanf"+ , "fdim"+ , "fdimf"+ , "fdiv_rd"+ , "fdiv_rn"+ , "fdiv_ru"+ , "fdiv_rz"+ , "ffs"+ , "ffsll"+ , "finitef"+ , "float2half_rn"+ , "float2int_rd"+ , "float2int_rn"+ , "float2int_ru"+ , "float2int_rz"+ , "float2ll_rd"+ , "float2ll_rn"+ , "float2ll_ru"+ , "float2ll_rz"+ , "float2uint_rd"+ , "float2uint_rn"+ , "float2uint_ru"+ , "float2uint_rz"+ , "float2ull_rd"+ , "float2ull_rn"+ , "float2ull_ru"+ , "float2ull_rz"+ , "float_as_int"+ , "floor"+ , "floorf"+ , "fma"+ , "fma_rd"+ , "fma_rn"+ , "fma_ru"+ , "fma_rz"+ , "fmaf"+ , "fmaf_rd"+ , "fmaf_rn"+ , "fmaf_ru"+ , "fmaf_rz"+ , "fmax"+ , "fmaxf"+ , "fmin"+ , "fminf"+ , "fmod"+ , "fmodf"+ , "fmul_rd"+ , "fmul_rn"+ , "fmul_ru"+ , "fmul_rz"+ , "frcp_rd"+ , "frcp_rn"+ , "frcp_ru"+ , "frcp_rz"+ , "frexp"+ , "frexpf"+ , "frsqrt_rn"+ , "fsqrt_rd"+ , "fsqrt_rn"+ , "fsqrt_ru"+ , "fsqrt_rz"+ , "fsub_rd"+ , "fsub_rn"+ , "fsub_ru"+ , "fsub_rz"+ , "hadd"+ , "half2float"+ , "hiloint2double"+ , "hypot"+ , "hypotf"+ , "ilogb"+ , "ilogbf"+ , "int2double_rn"+ , "int2float_rd"+ , "int2float_rn"+ , "int2float_ru"+ , "int2float_rz"+ , "int_as_float"+ , "isfinited"+ , "isinfd"+ , "isinff"+ , "isnand"+ , "isnanf"+ , "j0"+ , "j0f"+ , "j1"+ , "j1f"+ , "jn"+ , "jnf"+ , "ldexp"+ , "ldexpf"+ , "lgamma"+ , "lgammaf"+ , "ll2double_rd"+ , "ll2double_rn"+ , "ll2double_ru"+ , "ll2double_rz"+ , "ll2float_rd"+ , "ll2float_rn"+ , "ll2float_ru"+ , "ll2float_rz"+ , "llabs"+ , "llmax"+ , "llmin"+ , "llrint"+ , "llrintf"+ , "llround"+ , "llroundf"+ , "log"+ , "log10"+ , "log10f"+ , "log1p"+ , "log1pf"+ , "log2"+ , "log2f"+ , "logb"+ , "logbf"+ , "logf"+ , "longlong_as_double"+ , "max"+ , "min"+ , "modf"+ , "modff"+ , "mul24"+ , "mul64hi"+ , "mulhi"+ , "nan"+ , "nanf"+ , "nearbyint"+ , "nearbyintf"+ , "nextafter"+ , "nextafterf"+ , "normcdf"+ , "normcdff"+ , "normcdfinv"+ , "normcdfinvf"+ , "popc"+ , "popcll"+ , "pow"+ , "powf"+ , "powi"+ , "powif"+ , "rcbrt"+ , "rcbrtf"+ , "remainder"+ , "remainderf"+ , "remquo"+ , "remquof"+ , "rhadd"+ , "rint"+ , "rintf"+ , "round"+ , "roundf"+ , "rsqrt"+ , "rsqrtf"+ , "sad"+ , "saturatef"+ , "scalbn"+ , "scalbnf"+ , "signbitd"+ , "signbitf"+ , "sin"+ , "sincos"+ , "sincosf"+ , "sincospi"+ , "sincospif"+ , "sinf"+ , "sinh"+ , "sinhf"+ , "sinpi"+ , "sinpif"+ , "sqrt"+ , "sqrtf"+ , "tan"+ , "tanf"+ , "tanh"+ , "tanhf"+ , "tgamma"+ , "tgammaf"+ , "trunc"+ , "truncf"+ , "uhadd"+ , "uint2double_rn"+ , "uint2float_rd"+ , "uint2float_rn"+ , "uint2float_ru"+ , "uint2float_rz"+ , "ull2double_rd"+ , "ull2double_rn"+ , "ull2double_ru"+ , "ull2double_rz"+ , "ull2float_rd"+ , "ull2float_rn"+ , "ull2float_ru"+ , "ull2float_rz"+ , "ullmax"+ , "ullmin"+ , "umax"+ , "umin"+ , "umul24"+ , "umul64hi"+ , "umulhi"+ , "urhadd"+ , "usad"+ , "y0"+ , "y0f"+ , "y1"+ , "y1f"+ , "yn"+ , "ynf"+ ]+
Data/Array/Accelerate/LLVM/PTX/CodeGen/Queue.hs view
@@ -1,6 +1,6 @@ {-# LANGUAGE OverloadedStrings #-} {-# LANGUAGE ScopedTypeVariables #-}-{-# LANGUAGE ViewPatterns #-}+{-# LANGUAGE TemplateHaskell #-} -- | -- Module : Data.Array.Accelerate.LLVM.PTX.CodeGen.Queue -- Copyright : [2014..2017] Trevor L. McDonell@@ -109,7 +109,7 @@ mkQueueInit dev = let (start, _end, paramGang) = gangParam- config = launchConfig dev [1] (\_ -> 0) (\_ _ -> 1)+ config = launchConfig dev [1] (\_ -> 0) (\_ _ -> 1) [|| \_ _ -> 1 ||] in makeOpenAccWith config "qinit" paramGang $ do (queue,_) <- globalWorkQueue
Data/Array/Accelerate/LLVM/PTX/CodeGen/Scan.hs view
@@ -4,6 +4,7 @@ {-# LANGUAGE RebindableSyntax #-} {-# LANGUAGE RecordWildCards #-} {-# LANGUAGE ScopedTypeVariables #-}+{-# LANGUAGE TemplateHaskell #-} {-# LANGUAGE TypeOperators #-} {-# LANGUAGE ViewPatterns #-} -- |@@ -256,12 +257,12 @@ (arrTmp, paramTmp) = mutableArray ("tmp" :: Name (Vector e)) paramEnv = envParam aenv --- config = launchConfig dev (CUDA.incWarp dev) smem const+ config = launchConfig dev (CUDA.incWarp dev) smem const [|| const ||] smem n = warps * (1 + per_warp) * bytes where ws = CUDA.warpSize dev- warps = n `div` ws- per_warp = ws + ws `div` 2+ warps = n `P.quot` ws+ per_warp = ws + ws `P.quot` 2 bytes = sizeOf (eltType (undefined :: e)) in makeOpenAccWith config "scanP1" (paramGang ++ paramTmp ++ paramOut ++ paramEnv) $ do@@ -366,13 +367,14 @@ (arrTmp, paramTmp) = mutableArray ("tmp" :: Name (Vector e)) paramEnv = envParam aenv --- config = launchConfig dev (CUDA.incWarp dev) smem grid+ config = launchConfig dev (CUDA.incWarp dev) smem grid gridQ grid _ _ = 1+ gridQ = [|| \_ _ -> 1 ||] smem n = warps * (1 + per_warp) * bytes where ws = CUDA.warpSize dev- warps = n `div` ws- per_warp = ws + ws `div` 2+ warps = n `P.quot` ws+ per_warp = ws + ws `P.quot` 2 bytes = sizeOf (eltType (undefined :: e)) in makeOpenAccWith config "scanP2" (paramGang ++ paramTmp ++ paramEnv) $ do@@ -455,7 +457,7 @@ stride = local scalarType ("ix.stride" :: Name Int32) paramStride = scalarParameter scalarType ("ix.stride" :: Name Int32) --- config = launchConfig dev (CUDA.incWarp dev) (const 0) const+ config = launchConfig dev (CUDA.incWarp dev) (const 0) const [|| const ||] in makeOpenAccWith config "scanP3" (paramGang ++ paramTmp ++ paramOut ++ paramStride : paramEnv) $ do @@ -548,12 +550,12 @@ (arrTmp, paramTmp) = mutableArray ("tmp" :: Name (Vector e)) paramEnv = envParam aenv --- config = launchConfig dev (CUDA.incWarp dev) smem const+ config = launchConfig dev (CUDA.incWarp dev) smem const [|| const ||] smem n = warps * (1 + per_warp) * bytes where ws = CUDA.warpSize dev- warps = n `div` ws- per_warp = ws + ws `div` 2+ warps = n `P.quot` ws+ per_warp = ws + ws `P.quot` 2 bytes = sizeOf (eltType (undefined :: e)) in makeOpenAccWith config "scanP1" (paramGang ++ paramTmp ++ paramOut ++ paramEnv) $ do@@ -655,13 +657,14 @@ (arrSum, paramSum) = mutableArray ("sum" :: Name (Scalar e)) paramEnv = envParam aenv --- config = launchConfig dev (CUDA.incWarp dev) smem grid+ config = launchConfig dev (CUDA.incWarp dev) smem grid gridQ grid _ _ = 1+ gridQ = [|| \_ _ -> 1 ||] smem n = warps * (1 + per_warp) * bytes where ws = CUDA.warpSize dev- warps = n `div` ws- per_warp = ws + ws `div` 2+ warps = n `P.quot` ws+ per_warp = ws + ws `P.quot` 2 bytes = sizeOf (eltType (undefined :: e)) in makeOpenAccWith config "scanP2" (paramGang ++ paramTmp ++ paramSum ++ paramEnv) $ do@@ -749,7 +752,7 @@ stride = local scalarType ("ix.stride" :: Name Int32) paramStride = scalarParameter scalarType ("ix.stride" :: Name Int32) --- config = launchConfig dev (CUDA.incWarp dev) (const 0) const+ config = launchConfig dev (CUDA.incWarp dev) (const 0) const [|| const ||] in makeOpenAccWith config "scanP3" (paramGang ++ paramTmp ++ paramOut ++ paramStride : paramEnv) $ do @@ -830,12 +833,12 @@ (arrOut, paramOut) = mutableArray ("out" :: Name (Array (sh:.Int) e)) paramEnv = envParam aenv --- config = launchConfig dev (CUDA.incWarp dev) smem const+ config = launchConfig dev (CUDA.incWarp dev) smem const [|| const ||] smem n = warps * (1 + per_warp) * bytes where ws = CUDA.warpSize dev- warps = n `div` ws- per_warp = ws + ws `div` 2+ warps = n `P.quot` ws+ per_warp = ws + ws `P.quot` 2 bytes = sizeOf (eltType (undefined :: e)) in makeOpenAccWith config "scan" (paramGang ++ paramOut ++ paramEnv) $ do@@ -1020,12 +1023,12 @@ (arrSum, paramSum) = mutableArray ("sum" :: Name (Array sh e)) paramEnv = envParam aenv --- config = launchConfig dev (CUDA.incWarp dev) smem const+ config = launchConfig dev (CUDA.incWarp dev) smem const [|| const ||] smem n = warps * (1 + per_warp) * bytes where ws = CUDA.warpSize dev- warps = n `div` ws- per_warp = ws + ws `div` 2+ warps = n `P.quot` ws+ per_warp = ws + ws `P.quot` 2 bytes = sizeOf (eltType (undefined :: e)) in makeOpenAccWith config "scan" (paramGang ++ paramOut ++ paramSum ++ paramEnv) $ do@@ -1216,11 +1219,11 @@ -> CodeGen (IR e) scanBlockSMem dir dev combine nelem = warpScan >=> warpPrefix where- int32 :: Integral a => a -> IR (Int32)+ int32 :: Integral a => a -> IR Int32 int32 = lift . P.fromIntegral -- Temporary storage required for each warp- warp_smem_elems = CUDA.warpSize dev + (CUDA.warpSize dev `div` 2)+ warp_smem_elems = CUDA.warpSize dev + (CUDA.warpSize dev `P.quot` 2) warp_smem_bytes = warp_smem_elems * sizeOf (eltType (undefined::e)) -- Step 1: Scan in every warp@@ -1304,7 +1307,7 @@ -- Number of steps required to scan warp steps = P.floor (log2 (P.fromIntegral (CUDA.warpSize dev)))- halfWarp = P.fromIntegral (CUDA.warpSize dev `div` 2)+ halfWarp = P.fromIntegral (CUDA.warpSize dev `P.quot` 2) -- Unfold the scan as a recursive code generation function scan :: Int -> IR e -> CodeGen (IR e)
Data/Array/Accelerate/LLVM/PTX/Compile.hs view
@@ -18,22 +18,25 @@ module Data.Array.Accelerate.LLVM.PTX.Compile ( module Data.Array.Accelerate.LLVM.Compile,- ExecutableR(..), Kernel(..), ObjectCode,+ ObjectR(..), ) where -- llvm-hs-import LLVM.AST hiding ( Module ) import qualified LLVM.AST as AST import qualified LLVM.AST.Name as LLVM-import qualified LLVM.Analysis as LLVM import qualified LLVM.Context as LLVM import qualified LLVM.Module as LLVM import qualified LLVM.PassManager as LLVM+import qualified LLVM.Target as LLVM+import qualified LLVM.Internal.Module as LLVM.Internal+import qualified LLVM.Internal.FFI.LLVMCTypes as LLVM.Internal.FFI+#ifdef ACCELERATE_INTERNAL_CHECKS+import qualified LLVM.Analysis as LLVM+#endif -- accelerate import Data.Array.Accelerate.Error ( internalError )-import Data.Array.Accelerate.Lifetime import Data.Array.Accelerate.Trafo ( DelayedOpenAcc ) import Data.Array.Accelerate.LLVM.CodeGen@@ -41,105 +44,197 @@ import Data.Array.Accelerate.LLVM.CodeGen.Module ( Module(..) ) import Data.Array.Accelerate.LLVM.Compile import Data.Array.Accelerate.LLVM.State-#ifdef ACCELERATE_USE_NVVM import Data.Array.Accelerate.LLVM.Util-#endif import Data.Array.Accelerate.LLVM.PTX.Analysis.Launch import Data.Array.Accelerate.LLVM.PTX.CodeGen-import Data.Array.Accelerate.LLVM.PTX.Compile.Link-import Data.Array.Accelerate.LLVM.PTX.Context+import Data.Array.Accelerate.LLVM.PTX.Compile.Cache+import Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice import Data.Array.Accelerate.LLVM.PTX.Foreign ( ) import Data.Array.Accelerate.LLVM.PTX.Target -import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug+import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug -- cuda+import Foreign.CUDA.Path import qualified Foreign.CUDA.Analysis as CUDA-import qualified Foreign.CUDA.Driver as CUDA-#ifdef ACCELERATE_USE_NVVM import qualified Foreign.NVVM as NVVM-#endif -- standard library+import Control.Concurrent+import Control.DeepSeq+import Control.Exception import Control.Monad.Except import Control.Monad.State import Data.ByteString ( ByteString )-import Data.List ( intercalate )+import Data.ByteString.Short ( ShortByteString )+import Data.Maybe+import Data.Word+import Foreign.C+import Foreign.ForeignPtr+import Foreign.Ptr+import Foreign.Storable+import GHC.IO.Exception ( IOErrorType(..), IOException(..) )+import System.Directory+import System.Exit+import System.FilePath+import System.IO+import System.IO.Unsafe+import System.Process import Text.Printf ( printf )-import qualified Data.ByteString.Char8 as B+import qualified Data.ByteString as B+import qualified Data.ByteString.Char8 as B8+import qualified Data.ByteString.Internal as B import qualified Data.Map as Map import Prelude as P instance Compile PTX where- data ExecutableR PTX = PTXR { ptxKernel :: ![Kernel]- , ptxModule :: {-# UNPACK #-} !ObjectCode- }- compileForTarget = compileForPTX+ data ObjectR PTX = ObjectR { objId :: {-# UNPACK #-} !UID+ , ptxConfig :: ![(ShortByteString, LaunchConfig)]+ , objData :: {- LAZY -} ByteString+ }+ compileForTarget = compile -data Kernel = Kernel {- kernelFun :: {-# UNPACK #-} !CUDA.Fun- , kernelOccupancy :: {-# UNPACK #-} !CUDA.Occupancy- , kernelSharedMemBytes :: {-# UNPACK #-} !Int- , kernelThreadBlockSize :: {-# UNPACK #-} !Int- , kernelThreadBlocks :: (Int -> Int)- , kernelName :: String- }+-- | Compile an Accelerate expression to object code.+--+-- This generates the target code together with a list of each kernel function+-- defined in the module paired with its occupancy information.+--+compile :: DelayedOpenAcc aenv a -> Gamma aenv -> LLVM PTX (ObjectR PTX)+compile acc aenv = do+ target <- gets llvmTarget+ (uid, cacheFile) <- cacheOfOpenAcc acc -type ObjectCode = Lifetime CUDA.Module+ -- Generate code for this Acc operation+ --+ let Module ast md = llvmOfOpenAcc target uid acc aenv+ dev = ptxDeviceProperties target+ config = [ (f,x) | (LLVM.Name f, KM_PTX x) <- Map.toList md ] + -- Lower the generated LLVM into a CUBIN object code.+ --+ -- The 'objData' field is lazily evaluated since the object code might have+ -- already been loaded into the current context from a different function, in+ -- which case it will be found by the linker cache.+ --+ cubin <- liftIO . unsafeInterleaveIO $ do+ exists <- doesFileExist cacheFile+ recomp <- Debug.queryFlag Debug.force_recomp+ if exists && not (fromMaybe False recomp)+ then do+ Debug.traceIO Debug.dump_cc (printf "cc: found cached object code %016x" uid)+ B.readFile cacheFile --- | Compile a given module for the NVPTX backend. This produces a CUDA module--- as well as a list of the kernel functions in the module, together with some--- occupancy information.+ else+ LLVM.withContext $ \ctx -> do+ ptx <- compilePTX dev ctx ast+ cubin <- compileCUBIN dev cacheFile ptx+ return cubin++ return $! ObjectR uid config cubin+++-- | Compile the LLVM module to PTX assembly. This is done either by the+-- closed-source libNVVM library, or via the standard NVPTX backend (which is+-- the default). ---compileForPTX- :: DelayedOpenAcc aenv a- -> Gamma aenv- -> LLVM PTX (ExecutableR PTX)-compileForPTX acc aenv = do- target <- gets llvmTarget- let- Module ast md = llvmOfOpenAcc target acc aenv- dev = ptxDeviceProperties target+compilePTX :: CUDA.DeviceProperties -> LLVM.Context -> AST.Module -> IO ByteString+compilePTX dev ctx ast = do+#ifdef ACCELERATE_USE_NVVM+ ptx <- withLibdeviceNVVM dev ctx ast (compileModuleNVVM dev (AST.moduleName ast))+#else+ ptx <- withLibdeviceNVPTX dev ctx ast (compileModuleNVPTX dev)+#endif+ Debug.when Debug.dump_asm $ Debug.traceIO Debug.verbose (B8.unpack ptx)+ return ptx+++-- | Compile the given PTX assembly to a CUBIN file (SASS object code). The+-- compiled code will be stored at the given FilePath.+--+compileCUBIN :: CUDA.DeviceProperties -> FilePath -> ByteString -> IO ByteString+compileCUBIN dev sass ptx = do+ _verbose <- Debug.queryFlag Debug.verbose+ _debug <- Debug.queryFlag Debug.debug_cc --- liftIO . LLVM.withContext $ \ctx -> do- ptx <- compileModule dev ctx ast- funs <- sequence [ linkFunction ptx f x | (LLVM.Name f, KM_PTX x) <- Map.toList md ]- ptx' <- newLifetime ptx- addFinalizer ptx' $ do- Debug.traceIO Debug.dump_gc- $ printf "gc: unload module: %s"- $ intercalate "," (P.map kernelName funs)- withContext (ptxContext target) (CUDA.unload ptx)- return $! PTXR funs ptx'+ let verboseFlag = if _verbose then [ "-v" ] else []+ debugFlag = if _debug then [ "-g", "-lineinfo" ] else []+ arch = printf "-arch=sm_%d%d" m n+ CUDA.Compute m n = CUDA.computeCapability dev+ flags = "-" : "-o" : sass : arch : verboseFlag ++ debugFlag+ --+ cp = (proc (cudaBinPath </> "ptxas") flags)+ { std_in = CreatePipe+ , std_out = NoStream+ , std_err = CreatePipe+ } + -- Invoke the 'ptxas' executable (which must be on the PATH) to compile the+ -- PTX into SASS. The output is written directly to the final cache location.+ --+ withCreateProcess cp $ \(Just inh) Nothing (Just errh) ph -> do --- | Compile the LLVM module to produce a CUDA module.+ -- fork off a thread to start consuming stderr+ info <- hGetContents errh+ withForkWait (evaluate (rnf info)) $ \waitErr -> do++ -- write the PTX to the input handle+ -- closing the handle performs an implicit flush, thus may trigger SIGPIPE+ ignoreSIGPIPE $ B.hPut inh ptx+ ignoreSIGPIPE $ hClose inh++ -- wait on the output+ waitErr+ hClose errh++ -- wait on the process+ ex <- waitForProcess ph+ case ex of+ ExitFailure r -> $internalError "compile" (printf "ptxas %s (exit %d)\n%s" (unwords flags) r info)+ ExitSuccess -> return ()++ when _verbose $+ unless (null info) $+ Debug.traceIO Debug.dump_cc (printf "ptx: compiled entry function(s)\n%s" info)++ -- Read back the results+ B.readFile sass+++-- | Fork a thread while doing something else, but kill it if there's an+-- exception. ----- * If we are using NVVM, this includes all LLVM optimisations plus some--- sekrit optimisations.+-- This is important because we want to kill the thread that is holding the+-- Handle lock, because when we clean up the process we try to close that+-- handle, which could otherwise deadlock. ----- * If we are just using the llvm ptx backend, we still need to run the--- standard optimisations.+-- Stolen from the 'process' package. ---compileModule :: CUDA.DeviceProperties -> LLVM.Context -> AST.Module -> IO CUDA.Module-compileModule dev ctx ast =- let name = moduleName ast in-#ifdef ACCELERATE_USE_NVVM- withLibdeviceNVVM dev ctx ast (compileModuleNVVM dev name)-#else- withLibdeviceNVPTX dev ctx ast (compileModuleNVPTX dev name)-#endif+withForkWait :: IO () -> (IO () -> IO a) -> IO a+withForkWait async body = do+ waitVar <- newEmptyMVar :: IO (MVar (Either SomeException ()))+ mask $ \restore -> do+ tid <- forkIO $ try (restore async) >>= putMVar waitVar+ let wait = takeMVar waitVar >>= either throwIO return+ restore (body wait) `onException` killThread tid +ignoreSIGPIPE :: IO () -> IO ()+ignoreSIGPIPE =+ handle $ \e ->+ case e of+ IOError{..} | ResourceVanished <- ioe_type+ , Just ioe <- ioe_errno+ , Errno ioe == ePIPE+ -> return ()+ _ -> throwIO e -#ifdef ACCELERATE_USE_NVVM+ -- Compile and optimise the module to PTX using the (closed source) NVVM--- library. This may produce faster object code than the LLVM NVPTX compiler.+-- library. This _may_ produce faster object code than the LLVM NVPTX compiler. ---compileModuleNVVM :: CUDA.DeviceProperties -> String -> [(String, ByteString)] -> LLVM.Module -> IO CUDA.Module+compileModuleNVVM :: CUDA.DeviceProperties -> String -> [(String, ByteString)] -> LLVM.Module -> IO ByteString compileModuleNVVM dev name libdevice mdl = do _debug <- Debug.queryFlag Debug.debug_cc --@@ -165,7 +260,7 @@ Debug.when Debug.dump_cc $ do Debug.when Debug.verbose $ do ll <- LLVM.moduleLLVMAssembly mdl -- TLM: unfortunate to do the lowering twice in debug mode- Debug.traceIO Debug.verbose ll+ Debug.traceIO Debug.verbose (B8.unpack ll) -- Lower the generated module to bitcode, then compile and link together with -- the shim header and libdevice library (if necessary)@@ -173,120 +268,55 @@ ptx <- NVVM.compileModules (("",header) : (name,bc) : libdevice) flags unless (B.null (NVVM.compileLog ptx)) $ do- Debug.traceIO Debug.dump_cc $ "llvm: " ++ B.unpack (NVVM.compileLog ptx)+ Debug.traceIO Debug.dump_cc $ "llvm: " ++ B8.unpack (NVVM.compileLog ptx) - -- Link into a new CUDA module in the current context- linkPTX name (NVVM.compileResult ptx)+ -- Return the generated binary code+ return (NVVM.compileResult ptx) -#else+ -- Compiling with the NVPTX backend uses LLVM-3.3 and above ---compileModuleNVPTX :: CUDA.DeviceProperties -> String -> LLVM.Module -> IO CUDA.Module-compileModuleNVPTX dev name mdl =+compileModuleNVPTX :: CUDA.DeviceProperties -> LLVM.Module -> IO ByteString+compileModuleNVPTX dev mdl = withPTXTargetMachine dev $ \nvptx -> do -- Run the standard optimisation pass --- let pss = LLVM.defaultCuratedPassSetSpec { LLVM.optLevel = Just 3 }- runError e = either ($internalError "compileModuleNVPTX") id `fmap` runExceptT e-+ let pss = LLVM.defaultCuratedPassSetSpec { LLVM.optLevel = Just 3 } LLVM.withPassManager pss $ \pm -> do #ifdef ACCELERATE_INTERNAL_CHECKS- runError $ LLVM.verify mdl+ LLVM.verify mdl #endif b1 <- LLVM.runPassManager pm mdl -- debug printout Debug.when Debug.dump_cc $ do Debug.traceIO Debug.dump_cc $ printf "llvm: optimisation did work? %s" (show b1)- Debug.traceIO Debug.verbose =<< LLVM.moduleLLVMAssembly mdl+ Debug.traceIO Debug.verbose . B8.unpack =<< LLVM.moduleLLVMAssembly mdl -- Lower the LLVM module into target assembly (PTX)- ptx <- runError (LLVM.moduleTargetAssembly nvptx mdl)-- -- Link into a new CUDA module in the current context- linkPTX name (B.pack ptx)-#endif---- | Load the given CUDA PTX into a new module that is linked into the current--- context.----linkPTX :: String -> ByteString -> IO CUDA.Module-linkPTX name ptx = do- _verbose <- Debug.queryFlag Debug.verbose- _debug <- Debug.queryFlag Debug.debug_cc- --- let v = if _verbose then [ CUDA.Verbose ] else []- d = if _debug then [ CUDA.GenerateDebugInfo, CUDA.GenerateLineInfo ] else []- flags = concat [v,d]- --- Debug.when (Debug.dump_asm) $- Debug.traceIO Debug.verbose (B.unpack ptx)-- jit <- CUDA.loadDataEx ptx flags-- Debug.traceIO Debug.dump_asm $- printf "ptx: compiled entry function \"%s\" in %s\n%s"- name- (Debug.showFFloatSIBase (Just 2) 1000 (CUDA.jitTime jit / 1000) "s")- (B.unpack (CUDA.jitInfoLog jit))-- return $! CUDA.jitModule jit----- | Extract the named function from the module and package into a Kernel--- object, which includes meta-information on resource usage.------ If we are in debug mode, print statistics on kernel resource usage, etc.----linkFunction- :: CUDA.Module -- the compiled module- -> String -- __global__ entry function name- -> LaunchConfig -- launch configuration for this global function- -> IO Kernel-linkFunction mdl name configure = do- f <- CUDA.getFun mdl name- regs <- CUDA.requires f CUDA.NumRegs- ssmem <- CUDA.requires f CUDA.SharedSizeBytes- cmem <- CUDA.requires f CUDA.ConstSizeBytes- lmem <- CUDA.requires f CUDA.LocalSizeBytes- maxt <- CUDA.requires f CUDA.MaxKernelThreadsPerBlock-- let- (occ, cta, grid, dsmem) = configure maxt regs ssmem-- msg1, msg2 :: String- msg1 = printf "kernel function '%s' used %d registers, %d bytes smem, %d bytes lmem, %d bytes cmem"- name regs (ssmem + dsmem) lmem cmem-- msg2 = printf "multiprocessor occupancy %.1f %% : %d threads over %d warps in %d blocks"- (CUDA.occupancy100 occ)- (CUDA.activeThreads occ)- (CUDA.activeWarps occ)- (CUDA.activeThreadBlocks occ)-- Debug.traceIO Debug.dump_cc (printf "cc: %s\n ... %s" msg1 msg2)- return $ Kernel f occ dsmem cta grid name+ moduleTargetAssembly nvptx mdl -{----- | Extract the names of the function definitions from the module.------ Note: [Extracting global function names]------ It is important to run this on the module given to us by code generation.--- After combining modules with 'libdevice', extra function definitions,--- corresponding to basic maths operations, will be added to the module. These--- functions will not be callable as __global__ functions.------ The list of names will be exported in the order that they appear in the--- module.+-- | Produce target specific assembly as a 'ByteString'. ---globalFunctions :: [Definition] -> [String]-globalFunctions defs =- [ n | GlobalDefinition Function{..} <- defs- , not (null basicBlocks)- , let Name n = name- ]---}+moduleTargetAssembly :: LLVM.TargetMachine -> LLVM.Module -> IO ByteString+moduleTargetAssembly tm m = unsafe0 =<< LLVM.Internal.emitToByteString LLVM.Internal.FFI.codeGenFileTypeAssembly tm m+ where+ -- Ensure that the ByteString is NULL-terminated, so that it can be passed+ -- directly to C. This will unsafely mutate the underlying ForeignPtr if the+ -- string is not NULL-terminated but the last character is a whitespace+ -- character (there are usually a few blank lines at the end).+ --+ unsafe0 :: ByteString -> IO ByteString+ unsafe0 bs@(B.PS fp s l) =+ liftIO . withForeignPtr fp $ \p -> do+ let p' :: Ptr Word8+ p' = p `plusPtr` (s+l-1)+ --+ x <- peek p'+ case x of+ 0 -> return bs+ _ | B.isSpaceWord8 x -> poke p' 0 >> return bs+ _ -> return (B.snoc bs 0)
+ Data/Array/Accelerate/LLVM/PTX/Compile/Cache.hs view
@@ -0,0 +1,40 @@+{-# OPTIONS_GHC -fno-warn-orphans #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Compile.Cache+-- Copyright : [2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Compile.Cache (++ module Data.Array.Accelerate.LLVM.Compile.Cache++) where++import Data.Array.Accelerate.LLVM.Compile.Cache+import Data.Array.Accelerate.LLVM.PTX.Target++import Control.Monad.State+import Data.Version+import Foreign.CUDA.Analysis+import System.FilePath+import qualified Data.ByteString.Char8 as B8+import qualified Data.ByteString.Short.Char8 as S8++import Paths_accelerate_llvm_ptx+++instance Persistent PTX where+ targetCacheTemplate = do+ dev <- gets ptxDeviceProperties+ let Compute m n = computeCapability dev+ --+ return $ "accelerate-llvm-ptx-" ++ showVersion version+ </> S8.unpack ptxTargetTriple+ </> B8.unpack (ptxISAVersion m n)+ </> "morp.sass"+
Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice.hs view
@@ -1,9 +1,7 @@-{-# LANGUAGE FlexibleInstances #-} {-# LANGUAGE OverloadedStrings #-} {-# LANGUAGE RecordWildCards #-} {-# LANGUAGE TemplateHaskell #-}-{-# LANGUAGE TupleSections #-}-{-# OPTIONS_GHC -fno-warn-orphans #-}+{-# LANGUAGE ViewPatterns #-} -- | -- Module : Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice -- Copyright : [2014..2017] Trevor L. McDonell@@ -17,559 +15,163 @@ module Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice ( - nvvmReflect, libdevice,+ withLibdeviceNVVM,+ withLibdeviceNVPTX, ) where -- llvm-hs import LLVM.Context-import LLVM.Module as LLVM-import LLVM.AST as AST ( Module(..), Definition(..) )-import LLVM.AST.Attribute+import qualified LLVM.Module as LLVM++import LLVM.AST as AST import LLVM.AST.Global as G-import qualified LLVM.AST.Name as AST+import LLVM.AST.Linkage -- accelerate-import LLVM.AST.Type.Name ( Label(..) )-import LLVM.AST.Type.Representation--import Data.Array.Accelerate.Error-import Data.Array.Accelerate.LLVM.CodeGen.Base-import Data.Array.Accelerate.LLVM.CodeGen.Downcast-import Data.Array.Accelerate.LLVM.CodeGen.Intrinsic-import Data.Array.Accelerate.LLVM.PTX.Target+import Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice.Load+import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug -- cuda import Foreign.CUDA.Analysis -- standard library-import Control.Monad.Except+import Control.Monad import Data.ByteString ( ByteString )-import Data.HashMap.Strict ( HashMap )+import Data.ByteString.Short.Char8 ( ShortByteString )+import Data.HashSet ( HashSet ) import Data.List import Data.Maybe-import System.Directory-import System.FilePath-import System.IO.Unsafe import Text.Printf-import qualified Data.ByteString as B-import qualified Data.ByteString.Char8 as B8-import qualified Data.HashMap.Strict as HashMap----- NVVM Reflect--- --------------class NVVMReflect a where- nvvmReflect :: a--instance NVVMReflect AST.Module where- nvvmReflect = nvvmReflectPass_mdl--instance NVVMReflect (String, ByteString) where- nvvmReflect = nvvmReflectPass_bc+import qualified Data.ByteString.Short.Char8 as S8+import qualified Data.ByteString.Short.Extra as BS+import qualified Data.HashSet as Set --- This is a hacky module that can be linked against in order to provide the--- same functionality as running the NVVMReflect pass.------ Note: [NVVM Reflect Pass]------ To accommodate various math-related compiler flags that can affect code--- generation of libdevice code, the library code depends on a special LLVM IR--- pass (NVVMReflect) to handle conditional compilation within LLVM IR. This--- pass looks for calls to the @__nvvm_reflect function and replaces them with--- constants based on the defined reflection parameters.------ libdevice currently uses the following reflection parameters to control code--- generation:------ * __CUDA_FTZ={0,1} fast math that flushes denormals to zero+-- | Lower an LLVM AST to C++ objects and link it against the libdevice module,+-- iff any libdevice functions are referenced from the base module. ----- Since this is currently the only reflection parameter supported, and that we--- prefer correct results over pure speed, we do not flush denormals to zero. If--- the list of supported parameters ever changes, we may need to re-evaluate--- this implementation.+-- Note: [Linking with libdevice] ---nvvmReflectPass_mdl :: AST.Module-nvvmReflectPass_mdl =- AST.Module- { moduleName = "nvvm-reflect"- , moduleSourceFileName = []- , moduleDataLayout = targetDataLayout (undefined::PTX)- , moduleTargetTriple = targetTriple (undefined::PTX)- , moduleDefinitions = [GlobalDefinition $ functionDefaults- { name = AST.Name "__nvvm_reflect"- , returnType = downcast (integralType :: IntegralType Int32)- , parameters = ( [ptrParameter scalarType (UnName 0 :: Name (Ptr Int8))], False )- , G.functionAttributes = map Right [NoUnwind, ReadNone, AlwaysInline]- , basicBlocks = []- }]- }--{-# NOINLINE nvvmReflectPass_bc #-}-nvvmReflectPass_bc :: (String, ByteString)-nvvmReflectPass_bc = (name,) . unsafePerformIO $ do- withContext $ \ctx -> do- runError $ withModuleFromAST ctx nvvmReflectPass_mdl (return . B8.pack <=< moduleLLVMAssembly)- where- name = "__nvvm_reflect"- runError = either ($internalError "nvvmReflectPass") return <=< runExceptT----- libdevice--- ------------- Compatible version of libdevice for a given compute capability should be--- listed here:+-- The CUDA toolkit comes with an LLVM bitcode library called 'libdevice' that+-- implements many common mathematical functions. The library can be used as a+-- high performance math library for targets of the LLVM NVPTX backend, such as+-- this one. To link a module 'foo' with libdevice, the following compilation+-- pipeline is recommended: ----- https://github.com/llvm-mirror/llvm/blob/master/lib/Target/NVPTX/NVPTX.td#L72+-- 1. Save all external functions in module 'foo' ---class Libdevice a where- libdevice :: Compute -> a--instance Libdevice AST.Module where- libdevice (Compute n m) =- case (n,m) of- (2,_) -> libdevice_20_mdl -- 2.0, 2.1- (3,x) | x < 5 -> libdevice_30_mdl -- 3.0, 3.2- | otherwise -> libdevice_35_mdl -- 3.5, 3.7- (5,_) -> libdevice_50_mdl -- 5.x- (6,_) -> libdevice_50_mdl -- 6.x- _ -> $internalError "libdevice" "no binary for this architecture"--instance Libdevice (String, ByteString) where- libdevice (Compute n m) =- case (n,m) of- (2,_) -> libdevice_20_bc -- 2.0, 2.1- (3,x) | x < 5 -> libdevice_30_bc -- 3.0, 3.2- | otherwise -> libdevice_35_bc -- 3.5, 3.7- (5,_) -> libdevice_50_bc -- 5.x- (6,_) -> libdevice_50_bc -- 6.x- _ -> $internalError "libdevice" "no binary for this architecture"----- Load the libdevice bitcode files as an LLVM AST module. The top-level--- unsafePerformIO ensures that the data is only read from disk once per program--- execution.+-- 2. Link module 'foo' with the appropriate 'libdevice_compute_XX.YY.bc' ---{-# NOINLINE libdevice_20_mdl #-}-{-# NOINLINE libdevice_30_mdl #-}-{-# NOINLINE libdevice_35_mdl #-}-{-# NOINLINE libdevice_50_mdl #-}-libdevice_20_mdl, libdevice_30_mdl, libdevice_35_mdl, libdevice_50_mdl :: AST.Module-libdevice_20_mdl = unsafePerformIO $ libdeviceModule (Compute 2 0)-libdevice_30_mdl = unsafePerformIO $ libdeviceModule (Compute 3 0)-libdevice_35_mdl = unsafePerformIO $ libdeviceModule (Compute 3 5)-libdevice_50_mdl = unsafePerformIO $ libdeviceModule (Compute 5 0)---- Load the libdevice bitcode files as raw binary data. The top-level--- unsafePerformIO ensures that the data is read only once per program--- execution.+-- 3. Internalise all functions not in the list from (1) ---{-# NOINLINE libdevice_20_bc #-}-{-# NOINLINE libdevice_30_bc #-}-{-# NOINLINE libdevice_35_bc #-}-{-# NOINLINE libdevice_50_bc #-}-libdevice_20_bc, libdevice_30_bc, libdevice_35_bc, libdevice_50_bc :: (String,ByteString)-libdevice_20_bc = unsafePerformIO $ libdeviceBitcode (Compute 2 0)-libdevice_30_bc = unsafePerformIO $ libdeviceBitcode (Compute 3 0)-libdevice_35_bc = unsafePerformIO $ libdeviceBitcode (Compute 3 5)-libdevice_50_bc = unsafePerformIO $ libdeviceBitcode (Compute 5 0)----- Load the libdevice bitcode file for the given compute architecture, and raise--- it to a Haskell AST that can be kept for future use. The name of the bitcode--- files follows:+-- 4. Eliminate all unused internal functions ----- libdevice.compute_XX.YY.bc+-- 5. Run the NVVMReflect pass (see note: [NVVM Reflect Pass]) ----- Where XX represents the compute capability, and YY represents a version(?) We--- search the libdevice PATH for all files of the appropriate compute capability--- and load the most recent.+-- 6. Run the standard optimisation pipeline ---libdeviceModule :: Compute -> IO AST.Module-libdeviceModule arch = do- let bc :: (String, ByteString)- bc = libdevice arch+withLibdeviceNVPTX+ :: DeviceProperties+ -> Context+ -> Module+ -> (LLVM.Module -> IO a)+ -> IO a+withLibdeviceNVPTX dev ctx ast next =+ case Set.null externs of+ True -> LLVM.withModuleFromAST ctx ast next+ False ->+ LLVM.withModuleFromAST ctx ast $ \mdl ->+ LLVM.withModuleFromAST ctx nvvmReflect $ \refl ->+ LLVM.withModuleFromAST ctx (internalise externs libdev) $ \libd -> do+ LLVM.linkModules mdl refl+ LLVM.linkModules mdl libd+ Debug.traceIO Debug.dump_cc msg+ next mdl+ where+ -- Replace the target triple and datalayout from the libdevice.bc module+ -- with those of the generated code. This avoids warnings such as "linking+ -- two modules of different target triples..."+ libdev = (libdevice arch) { moduleTargetTriple = moduleTargetTriple ast+ , moduleDataLayout = moduleDataLayout ast+ }+ externs = analyse ast+ arch = computeCapability dev - -- TLM: we have called 'withContext' again here, although the LLVM state- -- already carries a version of the context. We do this so that we can- -- fully apply this function that can be lifted out to a CAF and only- -- executed once per program execution.- --- withContext $ \ctx ->- either ($internalError "libdeviceModule") id `fmap`- runExceptT (withModuleFromBitcode ctx bc moduleAST)+ msg = printf "cc: linking with libdevice: %s"+ $ intercalate ", "+ $ map S8.unpack+ $ Set.toList externs --- Load the libdevice bitcode file for the given compute architecture. The name--- of the bitcode files follows the format:+-- | Lower an LLVM AST to C++ objects and prepare it for linking against+-- libdevice using the nvvm bindings, iff any libdevice functions are referenced+-- from the base module. ----- libdevice.compute_XX.YY.bc+-- Rather than internalise and strip any unused functions ourselves, allow the+-- nvvm library to do so when linking the two modules together. ----- Where XX represents the compute capability, and YY represents a version(?) We--- search the libdevice PATH for all files of the appropriate compute capability--- and load the "most recent" (by sort order).+-- TLM: This really should work with the above method, however for some reason+-- we get a "CUDA Exception: function named symbol not found" error, even though+-- the function is clearly visible in the generated code. hmm... ---libdeviceBitcode :: Compute -> IO (String, ByteString)-libdeviceBitcode (Compute m n) = do- let arch = printf "libdevice.compute_%d%d" m n- err = $internalError "libdevice" (printf "not found: %s.YY.bc" arch)- best f = arch `isPrefixOf` f && takeExtension f == ".bc"+withLibdeviceNVVM+ :: DeviceProperties+ -> Context+ -> Module+ -> ([(String, ByteString)] -> LLVM.Module -> IO a)+ -> IO a+withLibdeviceNVVM dev ctx ast next =+ LLVM.withModuleFromAST ctx ast $ \mdl -> do+ when withlib $ Debug.traceIO Debug.dump_cc msg+ next lib mdl+ where+ externs = analyse ast+ withlib = not (Set.null externs)+ lib | withlib = [ nvvmReflect, libdevice arch ]+ | otherwise = [] - path <- libdevicePath- files <- getDirectoryContents path- name <- maybe err return . listToMaybe . sortBy (flip compare) $ filter best files- bc <- B.readFile (path </> name)+ arch = computeCapability dev - return (name, bc)+ msg = printf "cc: linking with libdevice: %s"+ $ intercalate ", "+ $ map S8.unpack+ $ Set.toList externs --- Determine the location of the libdevice bitcode libraries. We search for the--- location of the 'nvcc' executable in the PATH. From that, we assume the--- location of the libdevice bitcode files.+-- | Analyse the LLVM AST module and determine if any of the external+-- declarations are intrinsics implemented by libdevice. The set of such+-- functions is returned, and will be used when determining which functions from+-- libdevice to internalise. ---libdevicePath :: IO FilePath-libdevicePath = do- nvcc <- fromMaybe (error "could not find 'nvcc' in PATH") `fmap` findExecutable "nvcc"-- let ccvn = reverse (splitPath nvcc)- dir = "libdevice" : "nvvm" : drop 2 ccvn+analyse :: Module -> HashSet ShortByteString+analyse Module{..} =+ let intrinsic (GlobalDefinition Function{..})+ | null basicBlocks+ , Name n <- name+ , "__nv_" <- BS.take 5 n+ = Just n - return (joinPath (reverse dir))+ intrinsic _+ = Nothing+ in+ Set.fromList (mapMaybe intrinsic moduleDefinitions) -instance Intrinsic PTX where- intrinsicForTarget _ = libdeviceIndex---- The list of functions implemented by libdevice. These are all more-or-less--- named consistently based on the standard mathematical functions they--- implement, with the "__nv_" prefix stripped.+-- | Mark all definitions in the module as internal linkage. This means that+-- unused definitions can be removed as dead code. Be careful to leave any+-- declarations as external. ---libdeviceIndex :: HashMap String Label-libdeviceIndex =- let nv base = (base, Label $ "__nv_" ++ base)+internalise :: HashSet ShortByteString -> Module -> Module+internalise externals Module{..} =+ let internal (GlobalDefinition Function{..})+ | Name n <- name+ , not (Set.member n externals) -- we don't call this function directly; and+ , not (null basicBlocks) -- it is not an external declaration+ = GlobalDefinition Function { linkage=Internal, .. }++ internal x+ = x in- HashMap.fromList $ map nv- [ "abs"- , "acos"- , "acosf"- , "acosh"- , "acoshf"- , "asin"- , "asinf"- , "asinh"- , "asinhf"- , "atan"- , "atan2"- , "atan2f"- , "atanf"- , "atanh"- , "atanhf"- , "brev"- , "brevll"- , "byte_perm"- , "cbrt"- , "cbrtf"- , "ceil"- , "ceilf"- , "clz"- , "clzll"- , "copysign"- , "copysignf"- , "cos"- , "cosf"- , "cosh"- , "coshf"- , "cospi"- , "cospif"- , "dadd_rd"- , "dadd_rn"- , "dadd_ru"- , "dadd_rz"- , "ddiv_rd"- , "ddiv_rn"- , "ddiv_ru"- , "ddiv_rz"- , "dmul_rd"- , "dmul_rn"- , "dmul_ru"- , "dmul_rz"- , "double2float_rd"- , "double2float_rn"- , "double2float_ru"- , "double2float_rz"- , "double2hiint"- , "double2int_rd"- , "double2int_rn"- , "double2int_ru"- , "double2int_rz"- , "double2ll_rd"- , "double2ll_rn"- , "double2ll_ru"- , "double2ll_rz"- , "double2loint"- , "double2uint_rd"- , "double2uint_rn"- , "double2uint_ru"- , "double2uint_rz"- , "double2ull_rd"- , "double2ull_rn"- , "double2ull_ru"- , "double2ull_rz"- , "double_as_longlong"- , "drcp_rd"- , "drcp_rn"- , "drcp_ru"- , "drcp_rz"- , "dsqrt_rd"- , "dsqrt_rn"- , "dsqrt_ru"- , "dsqrt_rz"- , "erf"- , "erfc"- , "erfcf"- , "erfcinv"- , "erfcinvf"- , "erfcx"- , "erfcxf"- , "erff"- , "erfinv"- , "erfinvf"- , "exp"- , "exp10"- , "exp10f"- , "exp2"- , "exp2f"- , "expf"- , "expm1"- , "expm1f"- , "fabs"- , "fabsf"- , "fadd_rd"- , "fadd_rn"- , "fadd_ru"- , "fadd_rz"- , "fast_cosf"- , "fast_exp10f"- , "fast_expf"- , "fast_fdividef"- , "fast_log10f"- , "fast_log2f"- , "fast_logf"- , "fast_powf"- , "fast_sincosf"- , "fast_sinf"- , "fast_tanf"- , "fdim"- , "fdimf"- , "fdiv_rd"- , "fdiv_rn"- , "fdiv_ru"- , "fdiv_rz"- , "ffs"- , "ffsll"- , "finitef"- , "float2half_rn"- , "float2int_rd"- , "float2int_rn"- , "float2int_ru"- , "float2int_rz"- , "float2ll_rd"- , "float2ll_rn"- , "float2ll_ru"- , "float2ll_rz"- , "float2uint_rd"- , "float2uint_rn"- , "float2uint_ru"- , "float2uint_rz"- , "float2ull_rd"- , "float2ull_rn"- , "float2ull_ru"- , "float2ull_rz"- , "float_as_int"- , "floor"- , "floorf"- , "fma"- , "fma_rd"- , "fma_rn"- , "fma_ru"- , "fma_rz"- , "fmaf"- , "fmaf_rd"- , "fmaf_rn"- , "fmaf_ru"- , "fmaf_rz"- , "fmax"- , "fmaxf"- , "fmin"- , "fminf"- , "fmod"- , "fmodf"- , "fmul_rd"- , "fmul_rn"- , "fmul_ru"- , "fmul_rz"- , "frcp_rd"- , "frcp_rn"- , "frcp_ru"- , "frcp_rz"- , "frexp"- , "frexpf"- , "frsqrt_rn"- , "fsqrt_rd"- , "fsqrt_rn"- , "fsqrt_ru"- , "fsqrt_rz"- , "fsub_rd"- , "fsub_rn"- , "fsub_ru"- , "fsub_rz"- , "hadd"- , "half2float"- , "hiloint2double"- , "hypot"- , "hypotf"- , "ilogb"- , "ilogbf"- , "int2double_rn"- , "int2float_rd"- , "int2float_rn"- , "int2float_ru"- , "int2float_rz"- , "int_as_float"- , "isfinited"- , "isinfd"- , "isinff"- , "isnand"- , "isnanf"- , "j0"- , "j0f"- , "j1"- , "j1f"- , "jn"- , "jnf"- , "ldexp"- , "ldexpf"- , "lgamma"- , "lgammaf"- , "ll2double_rd"- , "ll2double_rn"- , "ll2double_ru"- , "ll2double_rz"- , "ll2float_rd"- , "ll2float_rn"- , "ll2float_ru"- , "ll2float_rz"- , "llabs"- , "llmax"- , "llmin"- , "llrint"- , "llrintf"- , "llround"- , "llroundf"- , "log"- , "log10"- , "log10f"- , "log1p"- , "log1pf"- , "log2"- , "log2f"- , "logb"- , "logbf"- , "logf"- , "longlong_as_double"- , "max"- , "min"- , "modf"- , "modff"- , "mul24"- , "mul64hi"- , "mulhi"- , "nan"- , "nanf"- , "nearbyint"- , "nearbyintf"- , "nextafter"- , "nextafterf"- , "normcdf"- , "normcdff"- , "normcdfinv"- , "normcdfinvf"- , "popc"- , "popcll"- , "pow"- , "powf"- , "powi"- , "powif"- , "rcbrt"- , "rcbrtf"- , "remainder"- , "remainderf"- , "remquo"- , "remquof"- , "rhadd"- , "rint"- , "rintf"- , "round"- , "roundf"- , "rsqrt"- , "rsqrtf"- , "sad"- , "saturatef"- , "scalbn"- , "scalbnf"- , "signbitd"- , "signbitf"- , "sin"- , "sincos"- , "sincosf"- , "sincospi"- , "sincospif"- , "sinf"- , "sinh"- , "sinhf"- , "sinpi"- , "sinpif"- , "sqrt"- , "sqrtf"- , "tan"- , "tanf"- , "tanh"- , "tanhf"- , "tgamma"- , "tgammaf"- , "trunc"- , "truncf"- , "uhadd"- , "uint2double_rn"- , "uint2float_rd"- , "uint2float_rn"- , "uint2float_ru"- , "uint2float_rz"- , "ull2double_rd"- , "ull2double_rn"- , "ull2double_ru"- , "ull2double_rz"- , "ull2float_rd"- , "ull2float_rn"- , "ull2float_ru"- , "ull2float_rz"- , "ullmax"- , "ullmin"- , "umax"- , "umin"- , "umul24"- , "umul64hi"- , "umulhi"- , "urhadd"- , "usad"- , "y0"- , "y0f"- , "y1"- , "y1f"- , "yn"- , "ynf"- ]+ Module { moduleDefinitions = map internal moduleDefinitions, .. }
+ Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice/Load.hs view
@@ -0,0 +1,127 @@+{-# LANGUAGE FlexibleInstances #-}+{-# LANGUAGE OverloadedStrings #-}+{-# LANGUAGE RecordWildCards #-}+{-# LANGUAGE TemplateHaskell #-}+{-# LANGUAGE TupleSections #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice.Load+-- Copyright : [2014..2017] Trevor L. McDonell+-- [2014..2014] Vinod Grover (NVIDIA Corporation)+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice.Load (++ nvvmReflect, libdevice,++) where++-- llvm-hs+import LLVM.Context+import LLVM.Module as LLVM+import LLVM.AST as AST ( Module(..) )++-- accelerate+import Data.Array.Accelerate.Error+import Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice.TH+import Data.Array.Accelerate.LLVM.PTX.Execute.Event ( ) -- GHC#1012+import Data.Array.Accelerate.LLVM.PTX.Execute.Stream ( ) -- GHC#1012++-- cuda+import Foreign.CUDA.Analysis++-- standard library+import Data.ByteString ( ByteString )+import System.IO.Unsafe+++-- NVVM Reflect+-- ------------++class NVVMReflect a where+ nvvmReflect :: a++instance NVVMReflect AST.Module where+ nvvmReflect = nvvmReflectModule++instance NVVMReflect (String, ByteString) where+ nvvmReflect = $$( nvvmReflectBitcode nvvmReflectModule )+++-- libdevice+-- ---------++-- Compatible version of libdevice for a given compute capability should be+-- listed here:+--+-- https://github.com/llvm-mirror/llvm/blob/master/lib/Target/NVPTX/NVPTX.td#L72+--+class Libdevice a where+ libdevice :: Compute -> a++instance Libdevice AST.Module where+ libdevice (Compute n m) =+ case (n,m) of+ (2,_) -> libdevice_20_mdl -- 2.0, 2.1+ (3,x) | x < 5 -> libdevice_30_mdl -- 3.0, 3.2+ | otherwise -> libdevice_35_mdl -- 3.5, 3.7+ (5,_) -> libdevice_50_mdl -- 5.x+ (6,_) -> libdevice_50_mdl -- 6.x+ _ -> $internalError "libdevice" "no binary for this architecture"++instance Libdevice (String, ByteString) where+ libdevice (Compute n m) =+ case (n,m) of+ (2,_) -> libdevice_20_bc -- 2.0, 2.1+ (3,x) | x < 5 -> libdevice_30_bc -- 3.0, 3.2+ | otherwise -> libdevice_35_bc -- 3.5, 3.7+ (5,_) -> libdevice_50_bc -- 5.x+ (6,_) -> libdevice_50_bc -- 6.x+ _ -> $internalError "libdevice" "no binary for this architecture"+++-- Load the libdevice bitcode files as an LLVM AST module. The top-level+-- unsafePerformIO ensures that the data is only read from disk once per program+-- execution.+--+{-# NOINLINE libdevice_20_mdl #-}+{-# NOINLINE libdevice_30_mdl #-}+{-# NOINLINE libdevice_35_mdl #-}+{-# NOINLINE libdevice_50_mdl #-}+libdevice_20_mdl, libdevice_30_mdl, libdevice_35_mdl, libdevice_50_mdl :: AST.Module+libdevice_20_mdl = unsafePerformIO $ libdeviceModule (Compute 2 0)+libdevice_30_mdl = unsafePerformIO $ libdeviceModule (Compute 3 0)+libdevice_35_mdl = unsafePerformIO $ libdeviceModule (Compute 3 5)+libdevice_50_mdl = unsafePerformIO $ libdeviceModule (Compute 5 0)++-- Load the libdevice bitcode files as raw binary data.+--+libdevice_20_bc, libdevice_30_bc, libdevice_35_bc, libdevice_50_bc :: (String,ByteString)+libdevice_20_bc = $$( libdeviceBitcode (Compute 2 0) )+libdevice_30_bc = $$( libdeviceBitcode (Compute 3 0) )+libdevice_35_bc = $$( libdeviceBitcode (Compute 3 5) )+libdevice_50_bc = $$( libdeviceBitcode (Compute 5 0) )+++-- Load the libdevice bitcode file for the given compute architecture, and raise+-- it to a Haskell AST that can be kept for future use. The name of the bitcode+-- files follows:+--+-- libdevice.compute_XX.YY.bc+--+-- Where XX represents the compute capability, and YY represents a version(?) We+-- search the libdevice PATH for all files of the appropriate compute capability+-- and load the most recent.+--+libdeviceModule :: Compute -> IO AST.Module+libdeviceModule arch = do+ let bc :: (String, ByteString)+ bc = libdevice arch+ --+ withContext $ \ctx ->+ withModuleFromBitcode ctx bc moduleAST+
+ Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice/TH.hs view
@@ -0,0 +1,184 @@+{-# LANGUAGE OverloadedStrings #-}+{-# LANGUAGE TemplateHaskell #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice.TH+-- Copyright : [2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice.TH (++ nvvmReflectModule, nvvmReflectBitcode,+ libdeviceBitcode,++) where++import qualified LLVM.AST as AST+import qualified LLVM.AST.Attribute as AST+import qualified LLVM.AST.Global as AST.G+import qualified LLVM.Context as LLVM+import qualified LLVM.Module as LLVM++import LLVM.AST.Type.Representation++import Data.Array.Accelerate.Error+import Data.Array.Accelerate.LLVM.CodeGen.Base+import Data.Array.Accelerate.LLVM.CodeGen.Downcast+import Data.Array.Accelerate.LLVM.PTX.Target++import Foreign.CUDA.Path+import Foreign.CUDA.Analysis++import Data.ByteString ( ByteString )+import Data.FileEmbed+import Data.List+import Data.Maybe+import Language.Haskell.TH.Syntax hiding ( Name )+import System.Directory+import System.FilePath+import Text.Printf+import qualified Data.ByteString.Short as BS+++-- This is a hacky module that can be linked against in order to provide the+-- same functionality as running the NVVMReflect pass.+--+-- Note: [NVVM Reflect Pass]+--+-- To accommodate various math-related compiler flags that can affect code+-- generation of libdevice code, the library code depends on a special LLVM IR+-- pass (NVVMReflect) to handle conditional compilation within LLVM IR. This+-- pass looks for calls to the @__nvvm_reflect function and replaces them with+-- constants based on the defined reflection parameters.+--+-- libdevice currently uses the following reflection parameters to control code+-- generation:+--+-- * __CUDA_FTZ={0,1} fast math that flushes denormals to zero+--+-- Since this is currently the only reflection parameter supported, and that we+-- prefer correct results over pure speed, we do not flush denormals to zero. If+-- the list of supported parameters ever changes, we may need to re-evaluate+-- this implementation.+--+nvvmReflectModule :: AST.Module+nvvmReflectModule =+ AST.Module+ { AST.moduleName = "nvvm-reflect"+ , AST.moduleSourceFileName = BS.empty+ , AST.moduleDataLayout = targetDataLayout (undefined::PTX)+ , AST.moduleTargetTriple = targetTriple (undefined::PTX)+ , AST.moduleDefinitions = [AST.GlobalDefinition $ AST.G.functionDefaults+ { AST.G.name = AST.Name "__nvvm_reflect"+ , AST.G.returnType = downcast (integralType :: IntegralType Int32)+ , AST.G.parameters = ( [ptrParameter scalarType (UnName 0 :: Name (Ptr Int8))], False )+ , AST.G.functionAttributes = map Right [AST.NoUnwind, AST.ReadNone, AST.AlwaysInline]+ , AST.G.basicBlocks = []+ }]+ }+++-- Lower the given NVVM Reflect module into bitcode.+--+nvvmReflectBitcode :: AST.Module -> Q (TExp (String, ByteString))+nvvmReflectBitcode mdl = do+ let name = "__nvvm_reflect"+ --+ bs <- runIO $ LLVM.withContext $ \ctx -> do+ LLVM.withModuleFromAST ctx mdl LLVM.moduleLLVMAssembly+ be <- bsToExp bs+ return . TExp $ TupE [ LitE (StringL name), be ]+++-- Load the libdevice bitcode file for the given compute architecture. The name+-- of the bitcode files follows the format:+--+-- libdevice.compute_XX.YY.bc+--+-- Where XX represents the compute capability, and YY represents a version(?) We+-- search the libdevice PATH for all files of the appropriate compute capability+-- and load the "most recent" (by sort order).+--+libdeviceBitcode :: Compute -> Q (TExp (String, ByteString))+libdeviceBitcode (Compute m n) = do+ let arch = printf "libdevice.compute_%d%d" m n+ err = $internalError "libdevice" (printf "not found: %s.YY.bc" arch)+ best f = arch `isPrefixOf` f && takeExtension f == ".bc"+ base = cudaInstallPath </> "nvvm" </> "libdevice"+ --+ files <- runIO $ getDirectoryContents base+ --+ let name = fromMaybe err . listToMaybe . sortBy (flip compare) $ filter best files+ path = base </> name+ --+ bc <- embedFile path+ return . TExp $ TupE [ LitE (StringL name), bc ]+++-- Determine the location of the libdevice bitcode libraries. We search for the+-- location of the 'nvcc' executable in the PATH. From that, we assume the+-- location of the libdevice bitcode files.+--+-- libdevicePath :: IO FilePath+-- libdevicepath = do+-- nvcc <- fromMaybe (error "could not find 'nvcc' in PATH") `fmap` findExecutable "nvcc"+-- --+-- let ccvn = reverse (splitPath nvcc)+-- dir = "libdevice" : "nvvm" : drop 2 ccvn+-- --+-- return (joinPath (reverse dir))+++-- With these instances it is possible to also write TH function to raise the+-- libNVVM modules to an AST. However, generating those large ASTs results in+-- awful compile times.+--+-- $( deriveLift ''AST.AddrSpace )+-- $( deriveLift ''AST.AlignType )+-- $( deriveLift ''AST.AlignmentInfo )+-- $( deriveLift ''AST.BasicBlock )+-- $( deriveLift ''AST.CallingConvention )+-- $( deriveLift ''AST.Constant )+-- $( deriveLift ''AST.DataLayout )+-- $( deriveLift ''AST.Definition )+-- $( deriveLift ''AST.Dialect )+-- $( deriveLift ''AST.Endianness )+-- $( deriveLift ''AST.FastMathFlags )+-- $( deriveLift ''AST.FloatingPointFormat )+-- $( deriveLift ''AST.FloatingPointPredicate )+-- $( deriveLift ''AST.FunctionAttribute )+-- $( deriveLift ''AST.Global )+-- $( deriveLift ''AST.GroupID )+-- $( deriveLift ''AST.InlineAssembly )+-- $( deriveLift ''AST.Instruction )+-- $( deriveLift ''AST.IntegerPredicate )+-- $( deriveLift ''AST.LandingPadClause )+-- $( deriveLift ''AST.Linkage )+-- $( deriveLift ''AST.Mangling )+-- $( deriveLift ''AST.MemoryOrdering )+-- $( deriveLift ''AST.Metadata )+-- $( deriveLift ''AST.MetadataNode )+-- $( deriveLift ''AST.MetadataNodeID )+-- $( deriveLift ''AST.Model )+-- $( deriveLift ''AST.Module )+-- $( deriveLift ''AST.Name )+-- $( deriveLift ''AST.Named )+-- $( deriveLift ''AST.Operand )+-- $( deriveLift ''AST.Parameter )+-- $( deriveLift ''AST.ParameterAttribute )+-- $( deriveLift ''AST.RMWOperation )+-- $( deriveLift ''AST.SelectionKind )+-- $( deriveLift ''AST.SomeFloat )+-- $( deriveLift ''AST.StorageClass )+-- $( deriveLift ''AST.SynchronizationScope )+-- $( deriveLift ''AST.TailCallKind )+-- $( deriveLift ''AST.Terminator )+-- $( deriveLift ''AST.Type )+-- $( deriveLift ''AST.UnnamedAddr )+-- $( deriveLift ''AST.Visibility )+-- $( deriveLift ''NonEmpty )+
− Data/Array/Accelerate/LLVM/PTX/Compile/Link.hs
@@ -1,183 +0,0 @@-{-# LANGUAGE OverloadedStrings #-}-{-# LANGUAGE RecordWildCards #-}-{-# LANGUAGE TemplateHaskell #-}-{-# LANGUAGE ViewPatterns #-}--- |--- Module : Data.Array.Accelerate.LLVM.PTX.Compile.Link--- Copyright : [2014..2017] Trevor L. McDonell--- [2014..2014] Vinod Grover (NVIDIA Corporation)--- License : BSD3------ Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>--- Stability : experimental--- Portability : non-portable (GHC extensions)-----module Data.Array.Accelerate.LLVM.PTX.Compile.Link (-- withLibdeviceNVVM,- withLibdeviceNVPTX,--) where---- llvm-hs-import LLVM.Context-import qualified LLVM.Module as LLVM--import LLVM.AST as AST-import LLVM.AST.Global as G-import LLVM.AST.Linkage---- accelerate-import Data.Array.Accelerate.Error--import Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice-import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug---- cuda-import Foreign.CUDA.Analysis---- standard library-import Control.Monad.Except-import Data.ByteString ( ByteString )-import Data.HashSet ( HashSet )-import Data.List-import Data.Maybe-import Text.Printf-import qualified Data.HashSet as Set----- | Lower an LLVM AST to C++ objects and link it against the libdevice module,--- iff any libdevice functions are referenced from the base module.------ Note: [Linking with libdevice]------ The CUDA toolkit comes with an LLVM bitcode library called 'libdevice' that--- implements many common mathematical functions. The library can be used as a--- high performance math library for targets of the LLVM NVPTX backend, such as--- this one. To link a module 'foo' with libdevice, the following compilation--- pipeline is recommended:------ 1. Save all external functions in module 'foo'------ 2. Link module 'foo' with the appropriate 'libdevice_compute_XX.YY.bc'------ 3. Internalise all functions not in the list from (1)------ 4. Eliminate all unused internal functions------ 5. Run the NVVMReflect pass (see note: [NVVM Reflect Pass])------ 6. Run the standard optimisation pipeline----withLibdeviceNVPTX- :: DeviceProperties- -> Context- -> Module- -> (LLVM.Module -> IO a)- -> IO a-withLibdeviceNVPTX dev ctx ast next =- case Set.null externs of- True -> runError $ LLVM.withModuleFromAST ctx ast next- False ->- runError $ LLVM.withModuleFromAST ctx ast $ \mdl ->- runError $ LLVM.withModuleFromAST ctx nvvmReflect $ \refl ->- runError $ LLVM.withModuleFromAST ctx (internalise externs libdev) $ \libd -> do- runError $ linkModules mdl refl- runError $ linkModules mdl libd- Debug.traceIO Debug.dump_cc msg- next mdl- where- -- Replace the target triple and datalayout from the libdevice.bc module- -- with those of the generated code. This avoids warnings such as "linking- -- two modules of different target triples..."- libdev = (libdevice arch) { moduleTargetTriple = moduleTargetTriple ast- , moduleDataLayout = moduleDataLayout ast- }- externs = analyse ast- arch = computeCapability dev- runError = either ($internalError "withLibdeviceNVPTX") return <=< runExceptT-- msg = printf "cc: linking with libdevice: %s"- $ intercalate ", " (Set.toList externs)----- | Link LLVM modules by copying parts of the second argument into the first.----linkModules- :: LLVM.Module -- module into which to link (destination: contains all symbols)- -> LLVM.Module -- module to copy into the other (this is destroyed in the process)- -> ExceptT String IO ()-linkModules = LLVM.linkModules----- | Lower an LLVM AST to C++ objects and prepare it for linking against--- libdevice using the nvvm bindings, iff any libdevice functions are referenced--- from the base module.------ Rather than internalise and strip any unused functions ourselves, allow the--- nvvm library to do so when linking the two modules together.------ TLM: This really should work with the above method, however for some reason--- we get a "CUDA Exception: function named symbol not found" error, even though--- the function is clearly visible in the generated code. hmm...----withLibdeviceNVVM- :: DeviceProperties- -> Context- -> Module- -> ([(String, ByteString)] -> LLVM.Module -> IO a)- -> IO a-withLibdeviceNVVM dev ctx ast next =- runError $ LLVM.withModuleFromAST ctx ast $ \mdl -> do- when withlib $ Debug.traceIO Debug.dump_cc msg- next lib mdl- where- externs = analyse ast- withlib = not (Set.null externs)- lib | withlib = [ nvvmReflect, libdevice arch ]- | otherwise = []-- arch = computeCapability dev- runError = either ($internalError "withLibdeviceNVPTX") return <=< runExceptT-- msg = printf "cc: linking with libdevice: %s"- $ intercalate ", " (Set.toList externs)----- | Analyse the LLVM AST module and determine if any of the external--- declarations are intrinsics implemented by libdevice. The set of such--- functions is returned, and will be used when determining which functions from--- libdevice to internalise.----analyse :: Module -> HashSet String-analyse Module{..} =- let intrinsic (GlobalDefinition Function{..})- | null basicBlocks- , Name n <- name- , "__nv_" <- take 5 n- = Just n-- intrinsic _- = Nothing- in- Set.fromList (mapMaybe intrinsic moduleDefinitions)----- | Mark all definitions in the module as internal linkage. This means that--- unused definitions can be removed as dead code. Be careful to leave any--- declarations as external.----internalise :: HashSet String -> Module -> Module-internalise externals Module{..} =- let internal (GlobalDefinition Function{..})- | Name n <- name- , not (Set.member n externals) -- we don't call this function directly; and- , not (null basicBlocks) -- it is not an external declaration- = GlobalDefinition (Function { linkage=Internal, .. })-- internal x- = x- in- Module { moduleDefinitions = map internal moduleDefinitions, .. }-
+ Data/Array/Accelerate/LLVM/PTX/Embed.hs view
@@ -0,0 +1,86 @@+{-# LANGUAGE QuasiQuotes #-}+{-# LANGUAGE TemplateHaskell #-}+{-# OPTIONS_GHC -fno-warn-orphans #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Embed+-- Copyright : [2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Embed (++ module Data.Array.Accelerate.LLVM.Embed,++) where++import Data.ByteString.Short.Char8 as S8+import Data.ByteString.Short.Internal as BS++import Data.Array.Accelerate.Lifetime++import Data.Array.Accelerate.LLVM.Compile+import Data.Array.Accelerate.LLVM.Embed++import Data.Array.Accelerate.LLVM.PTX.Compile+import Data.Array.Accelerate.LLVM.PTX.Link+import Data.Array.Accelerate.LLVM.PTX.Target+import Data.Array.Accelerate.LLVM.PTX.Context++-- import qualified Foreign.CUDA.Analysis as CUDA+import qualified Foreign.CUDA.Driver as CUDA++import Foreign.Ptr+import GHC.Ptr ( Ptr(..) )+import Language.Haskell.TH ( Q, TExp )+import System.IO.Unsafe+import qualified Data.ByteString as B+import qualified Data.ByteString.Unsafe as B+import qualified Language.Haskell.TH as TH+import qualified Language.Haskell.TH.Syntax as TH+++instance Embed PTX where+ embedForTarget = embed++-- Embed the given object code and set up to be reloaded at execution time.+--+embed :: PTX -> ObjectR PTX -> Q (TExp (ExecutableR PTX))+embed target (ObjectR _ cfg obj) = do+ -- Load the module to recover information such as number of registers and+ -- bytes of shared memory. It may be possible to do this without requiring an+ -- active CUDA context.+ kmd <- TH.runIO $ withContext (ptxContext target) $ do+ jit <- B.unsafeUseAsCString obj $ \p -> CUDA.loadDataFromPtrEx (castPtr p) []+ ks <- mapM (uncurry (linkFunctionQ (CUDA.jitModule jit))) cfg+ CUDA.unload (CUDA.jitModule jit)+ return ks++ -- Generate the embedded kernel executable. This will load the embedded object+ -- code into the current (at execution time) context.+ [|| unsafePerformIO $ do+ jit <- CUDA.loadDataFromPtrEx $$( TH.unsafeTExpCoerce [| Ptr $(TH.litE (TH.StringPrimL (B.unpack obj))) |] ) []+ fun <- newLifetime (FunctionTable $$(listE (map (linkQ 'jit) kmd)))+ return $ PTXR fun+ ||]+ where+ linkQ :: TH.Name -> (Kernel, Q (TExp (Int -> Int))) -> Q (TExp Kernel)+ linkQ jit (Kernel name _ dsmem cta _, grid) =+ [|| unsafePerformIO $ do+ f <- CUDA.getFun (CUDA.jitModule $$(TH.unsafeTExpCoerce (TH.varE jit))) $$(TH.unsafeTExpCoerce (TH.lift (S8.unpack name)))+ return $ Kernel $$(liftSBS name) f dsmem cta $$grid+ ||]++ listE :: [Q (TExp a)] -> Q (TExp [a])+ listE xs = TH.unsafeTExpCoerce (TH.listE (map TH.unTypeQ xs))++ liftSBS :: ShortByteString -> Q (TExp ShortByteString)+ liftSBS bs =+ let bytes = BS.unpack bs+ len = BS.length bs+ in+ [|| unsafePerformIO $ BS.createFromPtr $$( TH.unsafeTExpCoerce [| Ptr $(TH.litE (TH.StringPrimL bytes)) |]) len ||]+
Data/Array/Accelerate/LLVM/PTX/Execute.hs view
@@ -1,6 +1,7 @@ {-# LANGUAGE CPP #-} {-# LANGUAGE FlexibleContexts #-} {-# LANGUAGE GADTs #-}+{-# LANGUAGE OverloadedStrings #-} {-# LANGUAGE RecordWildCards #-} {-# LANGUAGE ScopedTypeVariables #-} {-# LANGUAGE TemplateHaskell #-}@@ -19,7 +20,8 @@ module Data.Array.Accelerate.LLVM.PTX.Execute ( - executeAcc, executeAfun1,+ executeAcc, executeAfun,+ executeOpenAcc, ) where @@ -36,10 +38,10 @@ import Data.Array.Accelerate.LLVM.PTX.Analysis.Launch ( multipleOf ) import Data.Array.Accelerate.LLVM.PTX.Array.Data import Data.Array.Accelerate.LLVM.PTX.Array.Prim ( memsetArrayAsync )-import Data.Array.Accelerate.LLVM.PTX.Compile import Data.Array.Accelerate.LLVM.PTX.Execute.Async import Data.Array.Accelerate.LLVM.PTX.Execute.Environment import Data.Array.Accelerate.LLVM.PTX.Execute.Marshal+import Data.Array.Accelerate.LLVM.PTX.Link import Data.Array.Accelerate.LLVM.PTX.Target import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug @@ -52,6 +54,7 @@ -- library import Control.Monad ( when ) import Control.Monad.State ( gets, liftIO )+import Data.ByteString.Short.Char8 ( ShortByteString, unpack ) import Data.Int ( Int32 ) import Data.List ( find ) import Data.Maybe ( fromMaybe )@@ -110,32 +113,29 @@ -> Stream -> sh -> LLVM PTX (Array sh e)-simpleOp PTXR{..} gamma aenv stream sh = do- let kernel = case ptxKernel of+simpleOp exe gamma aenv stream sh = withExecutable exe $ \ptxExecutable -> do+ let kernel = case functionTable ptxExecutable of k:_ -> k _ -> $internalError "simpleOp" "no kernels found" -- out <- allocateRemote sh ptx <- gets llvmTarget- liftIO $ executeOp ptx kernel ptxModule gamma aenv stream (IE 0 (size sh)) out+ liftIO $ executeOp ptx kernel gamma aenv stream (IE 0 (size sh)) out return out simpleNamed :: (Shape sh, Elt e)- => String+ => ShortByteString -> ExecutableR PTX -> Gamma aenv -> Aval aenv -> Stream -> sh -> LLVM PTX (Array sh e)-simpleNamed fun exe@PTXR{..} gamma aenv stream sh = do- let kernel = fromMaybe ($internalError "simpleNamed" ("not found: " ++ fun))- $ lookupKernel fun exe- --+simpleNamed fun exe gamma aenv stream sh = withExecutable exe $ \ptxExecutable -> do out <- allocateRemote sh ptx <- gets llvmTarget- liftIO $ executeOp ptx kernel ptxModule gamma aenv stream (IE 0 (size sh)) out+ liftIO $ executeOp ptx (ptxExecutable !# fun) gamma aenv stream (IE 0 (size sh)) out return out @@ -204,19 +204,18 @@ -> Stream -> DIM1 -> LLVM PTX (Scalar e)-foldAllOp exe@PTXR{..} gamma aenv stream (Z :. n) = do+foldAllOp exe gamma aenv stream (Z :. n) = withExecutable exe $ \ptxExecutable -> do ptx <- gets llvmTarget let- err = $internalError "foldAll" "kernel not found"- ks = fromMaybe err (lookupKernel "foldAllS" exe)- km1 = fromMaybe err (lookupKernel "foldAllM1" exe)- km2 = fromMaybe err (lookupKernel "foldAllM2" exe)+ ks = ptxExecutable !# "foldAllS"+ km1 = ptxExecutable !# "foldAllM1"+ km2 = ptxExecutable !# "foldAllM2" -- if kernelThreadBlocks ks n == 1 then do -- The array is small enough that we can compute it in a single step out <- allocateRemote Z- liftIO $ executeOp ptx ks ptxModule gamma aenv stream (IE 0 n) out+ liftIO $ executeOp ptx ks gamma aenv stream (IE 0 n) out return out else do@@ -230,12 +229,12 @@ | otherwise = do let s = m `multipleOf` kernelThreadBlockSize km2 out <- allocateRemote (Z :. s)- liftIO $ executeOp ptx km2 ptxModule gamma aenv stream (IE 0 s) (tmp, out)+ liftIO $ executeOp ptx km2 gamma aenv stream (IE 0 s) (tmp, out) rec out -- let s = n `multipleOf` kernelThreadBlockSize km1 tmp <- allocateRemote (Z :. s)- liftIO $ executeOp ptx km1 ptxModule gamma aenv stream (IE 0 s) tmp+ liftIO $ executeOp ptx km1 gamma aenv stream (IE 0 s) tmp rec tmp @@ -247,16 +246,15 @@ -> Stream -> (sh :. Int) -> LLVM PTX (Array sh e)-foldDimOp exe@PTXR{..} gamma aenv stream (sh :. sz) = do+foldDimOp exe gamma aenv stream (sh :. sz) = withExecutable exe $ \ptxExecutable -> do let- kernel = fromMaybe ($internalError "foldDim" "kernel not found")- $ if sz > 0- then lookupKernel "fold" exe- else lookupKernel "generate" exe+ kernel = if sz > 0+ then ptxExecutable !# "fold"+ else ptxExecutable !# "generate" -- out <- allocateRemote sh ptx <- gets llvmTarget- liftIO $ executeOp ptx kernel ptxModule gamma aenv stream (IE 0 (size sh)) out+ liftIO $ executeOp ptx kernel gamma aenv stream (IE 0 (size sh)) out return out @@ -269,22 +267,21 @@ -> (sh :. Int) -> (Z :. Int) -> LLVM PTX (Array (sh :. Int) e)-foldSegOp exe@PTXR{..} gamma aenv stream (sh :. sz) (Z :. ss) = do- let n = ss - 1 -- segments array has been 'scanl (+) 0'`ed+foldSegOp exe gamma aenv stream (sh :. sz) (Z :. ss) = withExecutable exe $ \ptxExecutable -> do+ let+ n = ss - 1 -- segments array has been 'scanl (+) 0'`ed m = size sh * n foldseg = if (sz`quot`ss) < (2 * kernelThreadBlockSize foldseg_cta) then foldseg_warp else foldseg_cta --- err = $internalError "foldSeg" "kernel not found"- foldseg_cta = fromMaybe err $ lookupKernel "foldSeg_block" exe- foldseg_warp = fromMaybe err $ lookupKernel "foldSeg_warp" exe- -- qinit = fromMaybe err $ lookupKernel "qinit" exe+ foldseg_cta = ptxExecutable !# "foldSeg_block"+ foldseg_warp = ptxExecutable !# "foldSeg_warp"+ -- qinit = ptxExecutable !# "qinit" -- out <- allocateRemote (sh :. n) ptx <- gets llvmTarget- liftIO $ do- executeOp ptx foldseg ptxModule gamma aenv stream (IE 0 m) out+ liftIO $ executeOp ptx foldseg gamma aenv stream (IE 0 m) out return out @@ -340,12 +337,11 @@ -> Int -- input size -> Int -- output size -> LLVM PTX (Vector e)-scanAllOp exe@PTXR{..} gamma aenv stream n m = do+scanAllOp exe gamma aenv stream n m = withExecutable exe $ \ptxExecutable -> do let- err = $internalError "scanAllOp" "kernel not found"- k1 = fromMaybe err (lookupKernel "scanP1" exe)- k2 = fromMaybe err (lookupKernel "scanP2" exe)- k3 = fromMaybe err (lookupKernel "scanP3" exe)+ k1 = ptxExecutable !# "scanP1"+ k2 = ptxExecutable !# "scanP2"+ k3 = ptxExecutable !# "scanP3" -- c = kernelThreadBlockSize k1 s = n `multipleOf` c@@ -357,13 +353,13 @@ -- which can be computed by a single thread block will require no -- additional work. tmp <- allocateRemote (Z :. s) :: LLVM PTX (Vector e)- liftIO $ executeOp ptx k1 ptxModule gamma aenv stream (IE 0 s) (tmp, out)+ liftIO $ executeOp ptx k1 gamma aenv stream (IE 0 s) (tmp, out) -- Step 2: Multi-block reductions need to compute the per-block prefix, -- then apply those values to the partial results. when (s > 1) $ do- liftIO $ executeOp ptx k2 ptxModule gamma aenv stream (IE 0 s) tmp- liftIO $ executeOp ptx k3 ptxModule gamma aenv stream (IE 0 (s-1)) (tmp, out, i32 c)+ liftIO $ executeOp ptx k2 gamma aenv stream (IE 0 s) tmp+ liftIO $ executeOp ptx k3 gamma aenv stream (IE 0 (s-1)) (tmp, out, i32 c) return out @@ -377,14 +373,10 @@ -> sh -> Int -> LLVM PTX (Array (sh:.Int) e)-scanDimOp exe@PTXR{..} gamma aenv stream sz m = do- let- kernel = fromMaybe ($internalError "scanDimOp" "kernel not found")- $ lookupKernel "scan" exe- --+scanDimOp exe gamma aenv stream sz m = withExecutable exe $ \ptxExecutable -> do ptx <- gets llvmTarget out <- allocateRemote (sz :. m)- liftIO $ executeOp ptx kernel ptxModule gamma aenv stream (IE 0 (size sz)) out+ liftIO $ executeOp ptx (ptxExecutable !# "scan") gamma aenv stream (IE 0 (size sz)) out return out @@ -426,12 +418,11 @@ -> Stream -> DIM1 -> LLVM PTX (Vector e, Scalar e)-scan'AllOp exe@PTXR{..} gamma aenv stream (Z :. n) = do+scan'AllOp exe gamma aenv stream (Z :. n) = withExecutable exe $ \ptxExecutable -> do let- err = $internalError "scan'AllOp" "kernel not found"- k1 = fromMaybe err (lookupKernel "scanP1" exe)- k2 = fromMaybe err (lookupKernel "scanP2" exe)- k3 = fromMaybe err (lookupKernel "scanP3" exe)+ k1 = ptxExecutable !# "scanP1"+ k2 = ptxExecutable !# "scanP2"+ k3 = ptxExecutable !# "scanP3" -- c = kernelThreadBlockSize k1 s = n `multipleOf` c@@ -442,7 +433,7 @@ -- Step 1: independent thread-block-wide scans. Each block stores its partial -- sum to a temporary array.- liftIO $ executeOp ptx k1 ptxModule gamma aenv stream (IE 0 s) (tmp, out)+ liftIO $ executeOp ptx k1 gamma aenv stream (IE 0 s) (tmp, out) -- If this was a small array that was processed by a single thread block then -- we are done, otherwise compute the per-block prefix and apply those values@@ -452,8 +443,8 @@ Array _ ad -> return (out, Array () ad) else do sum <- allocateRemote Z- liftIO $ executeOp ptx k2 ptxModule gamma aenv stream (IE 0 s) (tmp, sum)- liftIO $ executeOp ptx k3 ptxModule gamma aenv stream (IE 0 (s-1)) (tmp, out, i32 c)+ liftIO $ executeOp ptx k2 gamma aenv stream (IE 0 s) (tmp, sum)+ liftIO $ executeOp ptx k3 gamma aenv stream (IE 0 (s-1)) (tmp, out, i32 c) return (out, sum) @@ -465,14 +456,11 @@ -> Stream -> sh :. Int -> LLVM PTX (Array (sh:.Int) e, Array sh e)-scan'DimOp exe@PTXR{..} gamma aenv stream sh@(sz :. _) = do- let kernel = fromMaybe ($internalError "scan'DimOp" "kernel not found")- $ lookupKernel "scan" exe- --+scan'DimOp exe gamma aenv stream sh@(sz :. _) = withExecutable exe $ \ptxExecutable -> do ptx <- gets llvmTarget out <- allocateRemote sh sum <- allocateRemote sz- liftIO $ executeOp ptx kernel ptxModule gamma aenv stream (IE 0 (size sz)) (out,sum)+ liftIO $ executeOp ptx (ptxExecutable !# "scan") gamma aenv stream (IE 0 (size sz)) (out,sum) return (out,sum) @@ -486,10 +474,11 @@ -> sh -> Array sh' e -> LLVM PTX (Array sh' e)-permuteOp PTXR{..} gamma aenv stream inplace shIn dfs = do- let n = size shIn+permuteOp exe gamma aenv stream inplace shIn dfs = withExecutable exe $ \ptxExecutable -> do+ let+ n = size shIn m = size (shape dfs)- kernel = case ptxKernel of+ kernel = case functionTable ptxExecutable of k:_ -> k _ -> $internalError "permute" "no kernels found" --@@ -499,11 +488,11 @@ else cloneArrayAsync stream dfs -- case kernelName kernel of- "permute_rmw" -> liftIO $ executeOp ptx kernel ptxModule gamma aenv stream (IE 0 n) out+ "permute_rmw" -> liftIO $ executeOp ptx kernel gamma aenv stream (IE 0 n) out "permute_mutex" -> do barrier@(Array _ ad) <- allocateRemote (Z :. m) :: LLVM PTX (Vector Word32) memsetArrayAsync stream m 0 ad- liftIO $ executeOp ptx kernel ptxModule gamma aenv stream (IE 0 n) (out, barrier)+ liftIO $ executeOp ptx kernel gamma aenv stream (IE 0 n) (out, barrier) _ -> $internalError "permute" "unexpected kernel image" -- return out@@ -550,54 +539,57 @@ -- | Retrieve the named kernel ---lookupKernel :: String -> ExecutableR PTX -> Maybe Kernel-lookupKernel name PTXR{..} =- find (\k -> kernelName k == name) ptxKernel+(!#) :: FunctionTable -> ShortByteString -> Kernel+(!#) exe name+ = fromMaybe ($internalError "lookupFunction" ("function not found: " ++ unpack name))+ $ lookupKernel name exe +lookupKernel :: ShortByteString -> FunctionTable -> Maybe Kernel+lookupKernel name ptxExecutable =+ find (\k -> kernelName k == name) (functionTable ptxExecutable) + -- Execute the function implementing this kernel. -- executeOp :: Marshalable args => PTX -> Kernel- -> ObjectCode -> Gamma aenv -> Aval aenv -> Stream -> Range -> args -> IO ()-executeOp ptx@PTX{..} kernel@Kernel{..} oc gamma aenv stream r args =+executeOp ptx@PTX{..} kernel@Kernel{..} gamma aenv stream r args = runExecutable fillP kernelName defaultPPT r $ \start end _ -> do argv <- marshal ptx stream (i32 start, i32 end, args, (gamma,aenv))- launch kernel oc stream (end-start) argv+ launch kernel stream (end-start) argv -- Execute a device function with the given thread configuration and function -- parameters. ---launch :: Kernel -> ObjectCode -> Stream -> Int -> [CUDA.FunParam] -> IO ()-launch Kernel{..} oc stream n args =+launch :: Kernel -> Stream -> Int -> [CUDA.FunParam] -> IO ()+launch Kernel{..} stream n args = when (n > 0) $- withLifetime oc $ \_ -> withLifetime stream $ \st -> Debug.monitorProcTime query msg (Just st) $ CUDA.launchKernel kernelFun grid cta smem (Just st) args where- cta = (kernelThreadBlockSize, 1, 1)- grid = (kernelThreadBlocks n, 1, 1)- smem = kernelSharedMemBytes+ cta = (kernelThreadBlockSize, 1, 1)+ grid = (kernelThreadBlocks n, 1, 1)+ smem = kernelSharedMemBytes -- Debugging/monitoring support- query = if Debug.monitoringIsEnabled- then return True- else Debug.queryFlag Debug.dump_exec+ query = if Debug.monitoringIsEnabled+ then return True+ else Debug.queryFlag Debug.dump_exec fst3 (x,_,_) = x msg wall cpu gpu = do Debug.addProcessorTime Debug.PTX gpu Debug.traceIO Debug.dump_exec $ printf "exec: %s <<< %d, %d, %d >>> %s"- kernelName (fst3 grid) (fst3 cta) smem (Debug.elapsed wall cpu gpu)+ (unpack kernelName) (fst3 grid) (fst3 cta) smem (Debug.elapsed wall cpu gpu)
Data/Array/Accelerate/LLVM/PTX/Execute/Async.hs view
@@ -1,5 +1,4 @@-{-# LANGUAGE RecordWildCards #-}-{-# LANGUAGE TypeFamilies #-}+{-# LANGUAGE TypeFamilies #-} {-# OPTIONS_GHC -fno-warn-orphans #-} -- | -- Module : Data.Array.Accelerate.LLVM.PTX.Execute.Async
Data/Array/Accelerate/LLVM/PTX/Execute/Event.hs view
@@ -1,4 +1,4 @@-{-# LANGUAGE RecordWildCards #-}+{-# LANGUAGE NamedFieldPuns #-} -- | -- Module : Data.Array.Accelerate.LLVM.PTX.Execute.Event -- Copyright : [2014..2017] Trevor L. McDonell@@ -56,7 +56,7 @@ create' :: LLVM PTX Event.Event create' = do- PTX{..} <- gets llvmTarget+ PTX{ptxMemoryTable} <- gets llvmTarget me <- attempt "create/new" (liftIO . catchOOM $ Event.create [Event.DisableTiming]) `orElse` do Remote.reclaim ptxMemoryTable
Data/Array/Accelerate/LLVM/PTX/Execute/Marshal.hs view
@@ -5,7 +5,6 @@ {-# LANGUAGE FlexibleInstances #-} {-# LANGUAGE GADTs #-} {-# LANGUAGE MultiParamTypeClasses #-}-{-# LANGUAGE RecordWildCards #-} {-# LANGUAGE ScopedTypeVariables #-} {-# LANGUAGE TypeFamilies #-} {-# OPTIONS_GHC -fno-warn-orphans #-}
Data/Array/Accelerate/LLVM/PTX/Execute/Stream/Reservoir.hs view
@@ -42,7 +42,7 @@ -- {-# INLINEABLE new #-} new :: Context -> IO Reservoir-new _ctx = newMVar ( Seq.empty )+new _ctx = newMVar Seq.empty -- | Retrieve an execution stream from the reservoir, if one is available.
Data/Array/Accelerate/LLVM/PTX/Foreign.hs view
@@ -21,7 +21,8 @@ -- useful re-exports LLVM,- PTX,+ PTX(..),+ Context(..), liftIO, withDevicePtr, module Data.Array.Accelerate.LLVM.PTX.Array.Data,@@ -37,8 +38,9 @@ import Data.Array.Accelerate.LLVM.Foreign import Data.Array.Accelerate.LLVM.PTX.Array.Data import Data.Array.Accelerate.LLVM.PTX.Array.Prim+import Data.Array.Accelerate.LLVM.PTX.Context import Data.Array.Accelerate.LLVM.PTX.Execute.Async-import Data.Array.Accelerate.LLVM.PTX.Target ( PTX )+import Data.Array.Accelerate.LLVM.PTX.Target import Control.Monad.State import Data.Typeable
+ Data/Array/Accelerate/LLVM/PTX/Link.hs view
@@ -0,0 +1,153 @@+{-# LANGUAGE RecordWildCards #-}+{-# LANGUAGE TypeFamilies #-}+{-# OPTIONS_GHC -fno-warn-orphans #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Link+-- Copyright : [2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Link (++ module Data.Array.Accelerate.LLVM.Link,+ ExecutableR(..), FunctionTable(..), Kernel(..), ObjectCode,+ withExecutable,+ linkFunctionQ,++) where++import Data.Array.Accelerate.Lifetime++import Data.Array.Accelerate.LLVM.Link+import Data.Array.Accelerate.LLVM.State++import Data.Array.Accelerate.LLVM.PTX.Analysis.Launch+import Data.Array.Accelerate.LLVM.PTX.Compile+import Data.Array.Accelerate.LLVM.PTX.Context+import Data.Array.Accelerate.LLVM.PTX.Link.Cache+import Data.Array.Accelerate.LLVM.PTX.Link.Object+import Data.Array.Accelerate.LLVM.PTX.Target+import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug++-- cuda+import qualified Foreign.CUDA.Analysis as CUDA+import qualified Foreign.CUDA.Driver as CUDA++-- standard library+import Control.Monad.State+import Data.ByteString.Short.Char8 ( ShortByteString, unpack )+import Foreign.Ptr+import Language.Haskell.TH+import Text.Printf ( printf )+import qualified Data.ByteString.Unsafe as B+import Prelude as P hiding ( lookup )+++instance Link PTX where+ data ExecutableR PTX = PTXR { ptxExecutable :: {-# UNPACK #-} !(Lifetime FunctionTable)+ }+ linkForTarget = link+++-- | Load the generated object code into the current CUDA context.+--+link :: ObjectR PTX -> LLVM PTX (ExecutableR PTX)+link (ObjectR uid cfg obj) = do+ target <- gets llvmTarget+ cache <- gets ptxKernelTable+ funs <- liftIO $ dlsym uid cache $ do+ -- Load the SASS object code into the current CUDA context+ jit <- B.unsafeUseAsCString obj $ \p -> CUDA.loadDataFromPtrEx (castPtr p) []+ let mdl = CUDA.jitModule jit++ -- Extract the kernel functions+ nm <- FunctionTable `fmap` mapM (uncurry (linkFunction mdl)) cfg+ oc <- newLifetime mdl++ -- Finalise the module by unloading it from the CUDA context+ addFinalizer oc $ do+ Debug.traceIO Debug.dump_gc ("gc: unload module: " ++ show nm)+ withContext (ptxContext target) (CUDA.unload mdl)++ return (nm, oc)+ --+ return $! PTXR funs+++-- | Extract the named function from the module and package into a Kernel+-- object, which includes meta-information on resource usage.+--+-- If we are in debug mode, print statistics on kernel resource usage, etc.+--+linkFunction+ :: CUDA.Module -- the compiled module+ -> ShortByteString -- __global__ entry function name+ -> LaunchConfig -- launch configuration for this global function+ -> IO Kernel+linkFunction mdl name configure =+ fst `fmap` linkFunctionQ mdl name configure++linkFunctionQ+ :: CUDA.Module+ -> ShortByteString+ -> LaunchConfig+ -> IO (Kernel, Q (TExp (Int -> Int)))+linkFunctionQ mdl name configure = do+ f <- CUDA.getFun mdl (unpack name)+ regs <- CUDA.requires f CUDA.NumRegs+ ssmem <- CUDA.requires f CUDA.SharedSizeBytes+ cmem <- CUDA.requires f CUDA.ConstSizeBytes+ lmem <- CUDA.requires f CUDA.LocalSizeBytes+ maxt <- CUDA.requires f CUDA.MaxKernelThreadsPerBlock++ let+ (occ, cta, grid, dsmem, gridQ) = configure maxt regs ssmem++ msg1, msg2 :: String+ msg1 = printf "kernel function '%s' used %d registers, %d bytes smem, %d bytes lmem, %d bytes cmem"+ (unpack name) regs (ssmem + dsmem) lmem cmem++ msg2 = printf "multiprocessor occupancy %.1f %% : %d threads over %d warps in %d blocks"+ (CUDA.occupancy100 occ)+ (CUDA.activeThreads occ)+ (CUDA.activeWarps occ)+ (CUDA.activeThreadBlocks occ)++ Debug.traceIO Debug.dump_cc (printf "cc: %s\n ... %s" msg1 msg2)+ return (Kernel name f dsmem cta grid, gridQ)+++-- | Execute some operation with the supplied executable functions+--+withExecutable :: ExecutableR PTX -> (FunctionTable -> LLVM PTX b) -> LLVM PTX b+withExecutable PTXR{..} f = do+ r <- f (unsafeGetValue ptxExecutable)+ liftIO $ touchLifetime ptxExecutable+ return r+++{--+-- | Extract the names of the function definitions from the module.+--+-- Note: [Extracting global function names]+--+-- It is important to run this on the module given to us by code generation.+-- After combining modules with 'libdevice', extra function definitions,+-- corresponding to basic maths operations, will be added to the module. These+-- functions will not be callable as __global__ functions.+--+-- The list of names will be exported in the order that they appear in the+-- module.+--+globalFunctions :: [Definition] -> [String]+globalFunctions defs =+ [ n | GlobalDefinition Function{..} <- defs+ , not (null basicBlocks)+ , let Name n = name+ ]+--}+
+ Data/Array/Accelerate/LLVM/PTX/Link/Cache.hs view
@@ -0,0 +1,22 @@+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Link.Cache+-- Copyright : [2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Link.Cache (++ KernelTable,+ LC.new, LC.dlsym,++) where++import Data.Array.Accelerate.LLVM.PTX.Link.Object+import qualified Data.Array.Accelerate.LLVM.Link.Cache as LC++type KernelTable = LC.LinkCache FunctionTable ObjectCode+
+ Data/Array/Accelerate/LLVM/PTX/Link/Object.hs view
@@ -0,0 +1,41 @@+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Link.Object+-- Copyright : [2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Link.Object+ where++import Data.Array.Accelerate.Lifetime+import Data.ByteString.Short.Char8 ( ShortByteString, unpack )+import Data.List+import qualified Foreign.CUDA.Driver as CUDA+++-- | The kernel function table is a list of the kernels implemented by a given+-- CUDA device module+--+data FunctionTable = FunctionTable { functionTable :: [Kernel] }+data Kernel = Kernel+ { kernelName :: {-# UNPACK #-} !ShortByteString+ , kernelFun :: {-# UNPACK #-} !CUDA.Fun+ , kernelSharedMemBytes :: {-# UNPACK #-} !Int+ , kernelThreadBlockSize :: {-# UNPACK #-} !Int+ , kernelThreadBlocks :: (Int -> Int)+ }++instance Show FunctionTable where+ showsPrec _ f+ = showString "<<"+ . showString (intercalate "," [ unpack (kernelName k) | k <- functionTable f ])+ . showString ">>"++-- | Object code consists of executable code in the device address space+--+type ObjectCode = Lifetime CUDA.Module+
Data/Array/Accelerate/LLVM/PTX/State.hs view
@@ -21,13 +21,14 @@ -- accelerate import Data.Array.Accelerate.Error -import Data.Array.Accelerate.LLVM.State import Data.Array.Accelerate.LLVM.PTX.Analysis.Device import Data.Array.Accelerate.LLVM.PTX.Target-import qualified Data.Array.Accelerate.LLVM.PTX.Context as CT+import Data.Array.Accelerate.LLVM.State import qualified Data.Array.Accelerate.LLVM.PTX.Array.Table as MT-import qualified Data.Array.Accelerate.LLVM.PTX.Execute.Stream as ST+import qualified Data.Array.Accelerate.LLVM.PTX.Context as CT import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug+import qualified Data.Array.Accelerate.LLVM.PTX.Execute.Stream as ST+import qualified Data.Array.Accelerate.LLVM.PTX.Link.Cache as LC import Data.Range.Range ( Range(..) ) import Control.Parallel.Meta ( Executable(..) )@@ -60,8 +61,9 @@ createTargetForDevice dev prp flags = do ctx <- CT.new dev prp flags mt <- MT.new ctx+ lc <- LC.new st <- ST.new ctx- return $! PTX ctx mt st simpleIO+ return $! PTX ctx mt lc st simpleIO -- | Create a PTX execute target for the given device context@@ -74,8 +76,9 @@ prp <- CUDA.props dev ctx <- CT.raw dev prp ctx' mt <- MT.new ctx+ lc <- LC.new st <- ST.new ctx- return $! PTX ctx mt st simpleIO+ return $! PTX ctx mt lc st simpleIO {-# INLINE simpleIO #-}
Data/Array/Accelerate/LLVM/PTX/Target.hs view
@@ -1,7 +1,8 @@-{-# LANGUAGE CPP #-}-{-# LANGUAGE EmptyDataDecls #-}-{-# LANGUAGE TemplateHaskell #-}-{-# LANGUAGE TypeFamilies #-}+{-# LANGUAGE CPP #-}+{-# LANGUAGE EmptyDataDecls #-}+{-# LANGUAGE OverloadedStrings #-}+{-# LANGUAGE TemplateHaskell #-}+{-# LANGUAGE TypeFamilies #-} -- | -- Module : Data.Array.Accelerate.LLVM.PTX.Target -- Copyright : [2014..2017] Trevor L. McDonell@@ -39,12 +40,15 @@ import Data.Array.Accelerate.LLVM.PTX.Array.Table ( MemoryTable ) import Data.Array.Accelerate.LLVM.PTX.Context ( Context, deviceProperties ) import Data.Array.Accelerate.LLVM.PTX.Execute.Stream.Reservoir ( Reservoir )+import Data.Array.Accelerate.LLVM.PTX.Link.Cache ( KernelTable ) -- CUDA import qualified Foreign.CUDA.Driver as CUDA -- standard library-import Control.Monad.Except+import Data.ByteString ( ByteString )+import Data.ByteString.Short ( ShortByteString )+import Data.String import System.IO.Unsafe import Text.Printf import qualified Data.Map as Map@@ -63,6 +67,7 @@ data PTX = PTX { ptxContext :: {-# UNPACK #-} !Context , ptxMemoryTable :: {-# UNPACK #-} !MemoryTable+ , ptxKernelTable :: {-# UNPACK #-} !KernelTable , ptxStreamReservoir :: {-# UNPACK #-} !Reservoir , fillP :: {-# UNPACK #-} !Executable }@@ -99,15 +104,15 @@ ptxDataLayout = DataLayout { endianness = LittleEndian , mangling = Nothing- , aggregateLayout = AlignmentInfo 0 (Just 64)+ , aggregateLayout = AlignmentInfo 0 64 , stackAlignment = Nothing , pointerLayouts = Map.fromList- [ (AddrSpace 0, (wordSize, AlignmentInfo wordSize (Just wordSize))) ]+ [ (AddrSpace 0, (wordSize, AlignmentInfo wordSize wordSize)) ] , typeLayouts = Map.fromList $- [ ((IntegerAlign, 1), AlignmentInfo 8 (Just 8)) ] ++- [ ((IntegerAlign, i), AlignmentInfo i (Just i)) | i <- [8,16,32,64]] ++- [ ((VectorAlign, v), AlignmentInfo v (Just v)) | v <- [16,32,64,128]] ++- [ ((FloatAlign, f), AlignmentInfo f (Just f)) | f <- [32,64] ]+ [ ((IntegerAlign, 1), AlignmentInfo 8 8) ] +++ [ ((IntegerAlign, i), AlignmentInfo i i) | i <- [8,16,32,64]] +++ [ ((VectorAlign, v), AlignmentInfo v v) | v <- [16,32,64,128]] +++ [ ((FloatAlign, f), AlignmentInfo f f) | f <- [32,64] ] , nativeSizes = Just $ Set.fromList [ 16,32,64 ] } where@@ -116,7 +121,7 @@ -- | String that describes the target host. ---ptxTargetTriple :: String+ptxTargetTriple :: ShortByteString ptxTargetTriple = case bitSize (undefined::Int) of 32 -> "nvptx-nvidia-cuda"@@ -132,8 +137,8 @@ -> IO a withPTXTargetMachine dev go = let CUDA.Compute m n = CUDA.computeCapability dev- isa = ptxISAVersion m n- sm = printf "sm_%d%d" m n+ isa = CPUFeature (ptxISAVersion m n)+ sm = fromString (printf "sm_%d%d" m n) in withTargetOptions $ \options -> do withTargetMachine@@ -152,15 +157,15 @@ -- -- https://github.com/llvm-mirror/llvm/blob/master/lib/Target/NVPTX/NVPTX.td#L72 ---ptxISAVersion :: Int -> Int -> CPUFeature-ptxISAVersion 2 _ = CPUFeature "ptx40"-ptxISAVersion 3 7 = CPUFeature "ptx41"-ptxISAVersion 3 _ = CPUFeature "ptx40"-ptxISAVersion 5 0 = CPUFeature "ptx40"-ptxISAVersion 5 2 = CPUFeature "ptx41"-ptxISAVersion 5 3 = CPUFeature "ptx42"-ptxISAVersion 6 _ = CPUFeature "ptx50"-ptxISAVersion _ _ = CPUFeature "ptx40"+ptxISAVersion :: Int -> Int -> ByteString+ptxISAVersion 2 _ = "ptx40"+ptxISAVersion 3 7 = "ptx41"+ptxISAVersion 3 _ = "ptx40"+ptxISAVersion 5 0 = "ptx40"+ptxISAVersion 5 2 = "ptx41"+ptxISAVersion 5 3 = "ptx42"+ptxISAVersion 6 _ = "ptx50"+ptxISAVersion _ _ = "ptx40" -- | The NVPTX target for this host.@@ -172,5 +177,5 @@ ptxTarget :: LLVM.Target ptxTarget = unsafePerformIO $ do initializeAllTargets- either error fst `fmap` runExceptT (lookupTarget Nothing ptxTargetTriple)+ fst `fmap` lookupTarget Nothing ptxTargetTriple
+ README.md view
@@ -0,0 +1,188 @@+An LLVM backend for the Accelerate Array Language+=================================================++[](https://travis-ci.org/AccelerateHS/accelerate-llvm)+[](https://hackage.haskell.org/package/accelerate-llvm)+[](https://hub.docker.com/r/tmcdonell/accelerate-llvm/)+[](https://microbadger.com/images/tmcdonell/accelerate-llvm)++This package compiles Accelerate code to LLVM IR, and executes that code on+multicore CPUs as well as NVIDIA GPUs. This avoids the need to go through `nvcc`+or `clang`. For details on Accelerate, refer to the [main repository][GitHub].++We love all kinds of contributions, so feel free to open issues for missing+features as well as report (or fix!) bugs on the [issue tracker][Issues].++ [GitHub]: https://github.com/AccelerateHS/accelerate+ [Issues]: https://github.com/AccelerateHS/accelerate/issues+++ * [Dependencies](#dependencies)+ * [Docker](#docker)+ * [Installing LLVM](#installing-llvm)+ * [Homebrew](#homebrew)+ * [Debian/Ubuntu](#debianubuntu)+ * [Building from source](#building-from-source)+ * [Installing Accelerate-LLVM](#installing-accelerate-llvm)+ * [libNVVM](#libNVVM)+++Dependencies+------------++Haskell dependencies are available from Hackage, but there are several external+library dependencies that you will need to install as well:++ * [`LLVM`](http://llvm.org)+ * [`libFFI`](http://sourceware.org/libffi/) (if using the `accelerate-llvm-native` backend for multicore CPUs)+ * [`CUDA`](https://developer.nvidia.com/cuda-downloads) (if using the `accelerate-llvm-ptx` backend for NVIDIA GPUs)+++Docker+------++A [docker](https://www.docker.com) container is provided with this package+preinstalled (via stack) at `/opt/accelerate-llvm`. Note that if you wish to use+the `accelerate-llvm-ptx` GPU backend, you will need to install the [NVIDIA+docker](https://github.com/NVIDIA/nvidia-docker) plugin; see that page for more+information.++```sh+$ docker run -it tmcdonell/accelerate-llvm+```+++Installing LLVM+---------------++When installing LLVM, make sure that it includes the `libLLVM` shared library.+If you want to use the GPU targeting `accelerate-llvm-ptx` backend, make sure+you install (or build) LLVM with the 'nvptx' target.++## Homebrew++Example using [Homebrew](http://brew.sh) on macOS:++```sh+$ brew install llvm-hs/homebrew-llvm/llvm-4.0+```++## Debian/Ubuntu++For Debian/Ubuntu based Linux distributions, the LLVM.org website provides+binary distribution packages. Check [apt.llvm.org](http://apt.llvm.org) for+instructions for adding the correct package database for your OS version, and+then:++```sh+$ apt-get install llvm-4.0-dev+```++## Building from source++If your OS does not have an appropriate LLVM distribution available, you can also build from source. Detailed build instructions are available on the [LLVM.org website](http://releases.llvm.org/4.0.0/docs/CMake.html). Note that you will require at least [CMake 3.4.3](http://www.cmake.org/cmake/resources/software.html) and a recent C++ compiler; at least Clang 3.1, GCC 4.8, or Visual Studio 2015 (update 3).++ 1. Download and unpack the [LLVM-4.0 source code](http://releases.llvm.org/4.0.0/llvm-4.0.0.src.tar.xz). We'll refer to+ the path that the source tree was unpacked to as `LLVM_SRC`. Only the main+ LLVM source tree is required, but you can optionally add other components+ such as the Clang compiler or Polly loop optimiser. See the [LLVM releases](http://releases.llvm.org/download.html#4.0.0)+ page for the complete list.++ 2. Create a temporary build directory and `cd` into it, for example:+ ```sh+ $ mkdir /tmp/build+ $ cd /tmp/build+ ```++ 3. Execute the following to configure the build. Here `INSTALL_PREFIX` is+ where LLVM is to be installed, for example `/usr/local` or+ `$HOME/opt/llvm`:+ ```sh+ $ cmake $LLVM_SRC -DCMAKE_INSTALL_PREFIX=$INSTALL_PREFIX -DCMAKE_BUILD_TYPE=Release -DLLVM_ENABLE_ASSERTIONS=ON -DLLVM_BUILD_LLVM_DYLIB=ON -DLLVM_LINK_LLVM_DYLIB=ON+ ```+ See [options and variables](http://llvm.org/docs/CMake.html#options-and-variables)+ for a list of additional build parameters you can specify.++ 4. Build and install:+ ```sh+ $ cmake --build .+ $ cmake --build . --target install+ ```++ 5. For macOS only, some additional steps are useful to work around issues related+ to [System Integrity Protection](https://en.wikipedia.org/wiki/System_Integrity_Protection):+ ```sh+ cd $INSTALL_PREFIX/lib+ ln -s libLLVM.dylib libLLVM-4.0.dylib+ install_name_tool -id $PWD/libLTO.dylib libLTO.dylib+ install_name_tool -id $PWD/libLLVM.dylib libLLVM.dylib+ install_name_tool -change '@rpath/libLLVM.dylib' $PWD/libLLVM.dylib libLTO.dylib+ ```+++Installing Accelerate-LLVM+--------------------------++Once the dependencies are installed, we are ready to install `accelerate-llvm`.++For example, installation using [`stack`](http://docs.haskellstack.org/en/stable/README.html)+just requires you to point it to the appropriate configuration file:+```sh+$ ln -s stack-8.0.yaml stack.yaml+$ stack setup+$ stack install+```++Note that the version of [`llvm-hs`](https://hackage.haskell.org/package/llvm-hs)+used must match the installed version of LLVM, which is currently 4.0.+++## libNVVM++The `accelerate-llvm-ptx` backend can optionally be compiled to generate GPU+code using the `libNVVM` library, rather than LLVM's inbuilt NVPTX code+generator. `libNVVM` is a closed-source library distributed as part of the+NVIDIA CUDA toolkit, and is what the `nvcc` compiler itself uses internally when+compiling CUDA C code.++Using `libNVVM` _may_ improve GPU performance compared to the code generator+built in to LLVM. One difficulty with using it however is that since `libNVVM`+is also based on LLVM, and typically lags LLVM by several releases, you must+install `accelerate-llvm` with a "compatible" version of LLVM, which will depend+on the version of the CUDA toolkit you have installed. The following table shows+some combinations:++| | LLVM-3.3 | LLVM-3.4 | LLVM-3.5 | LLVM-3.8 | LLVM-3.9 | LLVM-4.0 |+|:------------:|:--------:|:--------:|:--------:|:--------:|:--------:|:--------:|+| **CUDA-7.0** | ⭕ | ❌ | | | | |+| **CUDA-7.5** | | ⭕ | ⭕ | ❌ | | |+| **CUDA-8.0** | | | ⭕ | ⭕ | ❌ | ❌ |++Where ⭕ = Works, and ❌ = Does not work.++Note that the above restrictions on CUDA and LLVM version exist _only_ if you+want to use the NVVM component. Otherwise, you should be free to use any+combination of CUDA and LLVM.++Also note that `accelerate-llvm-ptx` itself currently requires at least LLVM-3.5.++Using `stack`, either edit the `stack.yaml` and add the following section:++```yaml+flags:+ accelerate-llvm-ptx:+ nvvm: true+```++Or install using the following option on the command line:++```sh+$ stack install accelerate-llvm-ptx --flag accelerate-llvm-ptx:nvvm+```++If installing via `cabal`:++```sh+$ cabal install accelerate-llvm-ptx -fnvvm+```+
accelerate-llvm-ptx.cabal view
@@ -1,15 +1,68 @@ name: accelerate-llvm-ptx-version: 1.0.0.1+version: 1.1.0.0 cabal-version: >= 1.10-tested-with: GHC == 7.8.*+tested-with: GHC >= 7.10 build-type: Simple -synopsis: Accelerate backend generating LLVM+synopsis: Accelerate backend for NVIDIA GPUs description: This library implements a backend for the /Accelerate/ language which generates LLVM-IR targeting CUDA capable GPUs. For further information,- refer to the main /Accelerate/ package:- <http://hackage.haskell.org/package/accelerate>+ refer to the main <http://hackage.haskell.org/package/accelerate accelerate>+ package.+ .+ [/Dependencies/]+ .+ Haskell dependencies are available from Hackage. The following external+ libraries are alse required:+ .+ * <http://llvm.org LLVM>+ .+ * <https://developer.nvidia.com/cuda-downloads CUDA>+ .+ [/Installing LLVM/]+ .+ /Homebrew/+ .+ Example using Homebrew on macOS:+ .+ > brew install llvm-hs/homebrew-llvm/llvm-5.0+ .+ /Debian & Ubuntu/+ .+ For Debian/Ubuntu based Linux distributions, the LLVM.org website provides+ binary distribution packages. Check <http://apt.llvm.org apt.llvm.org> for+ instructions for adding the correct package database for your OS version,+ and then:+ .+ > apt-get install llvm-5.0-dev+ .+ /Building from source/+ .+ If your OS does not have an appropriate LLVM distribution available, you can+ also build from source. Detailed build instructions are available on+ <http://releases.llvm.org/5.0.0/docs/CMake.html LLVM.org>. Make sure to+ include the cmake build options+ @-DLLVM_BUILD_LLVM_DYLIB=ON -DLLVM_LINK_LLVM_DYLIB=ON@ so that the @libLLVM@+ shared library will be built. Also ensure that the @LLVM_TARGETS_TO_BUILD@+ option includes the @NVPTX@ target (if not specified all targets are built).+ .+ [/Installing accelerate-llvm/]+ .+ To use @accelerate-llvm@ it is important that the @llvm-hs@ package is+ installed against the @libLLVM@ shared library, rather than statically+ linked, so that we can use LLVM from GHCi and Template Haskell. This is the+ default configuration, but you can also enforce this explicitly by adding+ the following to your @stack.yaml@ file:+ .+ > flags:+ > llvm-hs:+ > shared-llvm: true+ .+ Or by specifying the @shared-llvm@ flag to cabal:+ .+ > cabal install llvm-hs -fshared-llvm+ . license: BSD3 license-file: LICENSE@@ -18,7 +71,11 @@ bug-reports: https://github.com/AccelerateHS/accelerate/issues category: Compilers/Interpreters, Concurrency, Data, Parallelism +extra-source-files:+ CHANGELOG.md+ README.md + -- Configuration flags -- ------------------- @@ -76,6 +133,7 @@ Data.Array.Accelerate.LLVM.PTX.CodeGen.Fold Data.Array.Accelerate.LLVM.PTX.CodeGen.FoldSeg Data.Array.Accelerate.LLVM.PTX.CodeGen.Generate+ Data.Array.Accelerate.LLVM.PTX.CodeGen.Intrinsic Data.Array.Accelerate.LLVM.PTX.CodeGen.Loop Data.Array.Accelerate.LLVM.PTX.CodeGen.Map Data.Array.Accelerate.LLVM.PTX.CodeGen.Permute@@ -83,29 +141,44 @@ Data.Array.Accelerate.LLVM.PTX.CodeGen.Scan Data.Array.Accelerate.LLVM.PTX.Compile- Data.Array.Accelerate.LLVM.PTX.Compile.Link+ Data.Array.Accelerate.LLVM.PTX.Compile.Cache Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice+ Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice.Load+ Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice.TH + Data.Array.Accelerate.LLVM.PTX.Link+ Data.Array.Accelerate.LLVM.PTX.Link.Cache+ Data.Array.Accelerate.LLVM.PTX.Link.Object++ Data.Array.Accelerate.LLVM.PTX.Embed+ Data.Array.Accelerate.LLVM.PTX.Execute.Environment Data.Array.Accelerate.LLVM.PTX.Execute.Event Data.Array.Accelerate.LLVM.PTX.Execute.Marshal + Paths_accelerate_llvm_ptx+ build-depends:- base >= 4.7 && < 4.10- , accelerate == 1.0.*- , accelerate-llvm == 1.0.*- , bytestring >= 0.9+ base >= 4.7 && < 4.11+ , accelerate == 1.1.*+ , accelerate-llvm == 1.1.*+ , bytestring >= 0.10.4 , containers >= 0.5 && <0.6- , cuda >= 0.7+ , cuda >= 0.8+ , deepseq >= 1.3 , directory >= 1.0 , dlist >= 0.6 , fclabels >= 2.0+ , file-embed >= 0.0.8 , filepath >= 1.0 , hashable >= 1.2- , llvm-hs >= 3.9- , llvm-hs-pure >= 3.9+ , llvm-hs >= 4.1 && < 5.1+ , llvm-hs-pure >= 4.1 && < 5.1 , mtl >= 2.2.1+ , nvvm >= 0.7.5 , pretty >= 1.1+ , process >= 1.4.3+ , template-haskell , time >= 1.4 , unordered-containers >= 0.2 @@ -119,8 +192,6 @@ if flag(nvvm) cpp-options: -DACCELERATE_USE_NVVM- build-depends:- nvvm >= 0.7.5 if flag(debug) cpp-options: -DACCELERATE_DEBUG@@ -141,7 +212,7 @@ source-repository this type: git- tag: 1.0.0.0+ tag: 1.1.0.0 location: https://github.com/AccelerateHS/accelerate-llvm.git -- vim: nospell