diff --git a/CHANGELOG.md b/CHANGELOG.md new file mode 100644 --- /dev/null +++ b/CHANGELOG.md @@ -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 + diff --git a/Data/Array/Accelerate/LLVM/PTX.hs b/Data/Array/Accelerate/LLVM/PTX.hs --- a/Data/Array/Accelerate/LLVM/PTX.hs +++ b/Data/Array/Accelerate/LLVM/PTX.hs @@ -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 +-- project for an example. +-- +-- [/Note:/] +-- +-- Due to , 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. diff --git a/Data/Array/Accelerate/LLVM/PTX/Analysis/Device.hs b/Data/Array/Accelerate/LLVM/PTX/Analysis/Device.hs --- a/Data/Array/Accelerate/LLVM/PTX/Analysis/Device.hs +++ b/Data/Array/Accelerate/LLVM/PTX/Analysis/Device.hs @@ -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 diff --git a/Data/Array/Accelerate/LLVM/PTX/Analysis/Launch.hs b/Data/Array/Accelerate/LLVM/PTX/Analysis/Launch.hs --- a/Data/Array/Accelerate/LLVM/PTX/Analysis/Launch.hs +++ b/Data/Array/Accelerate/LLVM/PTX/Analysis/Launch.hs @@ -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 ||] diff --git a/Data/Array/Accelerate/LLVM/PTX/Array/Prim.hs b/Data/Array/Accelerate/LLVM/PTX/Array/Prim.hs --- a/Data/Array/Accelerate/LLVM/PTX/Array/Prim.hs +++ b/Data/Array/Accelerate/LLVM/PTX/Array/Prim.hs @@ -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 diff --git a/Data/Array/Accelerate/LLVM/PTX/Array/Remote.hs b/Data/Array/Accelerate/LLVM/PTX/Array/Remote.hs --- a/Data/Array/Accelerate/LLVM/PTX/Array/Remote.hs +++ b/Data/Array/Accelerate/LLVM/PTX/Array/Remote.hs @@ -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 diff --git a/Data/Array/Accelerate/LLVM/PTX/CodeGen.hs b/Data/Array/Accelerate/LLVM/PTX/CodeGen.hs --- a/Data/Array/Accelerate/LLVM/PTX/CodeGen.hs +++ b/Data/Array/Accelerate/LLVM/PTX/CodeGen.hs @@ -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 diff --git a/Data/Array/Accelerate/LLVM/PTX/CodeGen/Base.hs b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Base.hs --- a/Data/Array/Accelerate/LLVM/PTX/CodeGen/Base.hs +++ b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Base.hs @@ -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 diff --git a/Data/Array/Accelerate/LLVM/PTX/CodeGen/Fold.hs b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Fold.hs --- a/Data/Array/Accelerate/LLVM/PTX/CodeGen/Fold.hs +++ b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Fold.hs @@ -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 -- diff --git a/Data/Array/Accelerate/LLVM/PTX/CodeGen/FoldSeg.hs b/Data/Array/Accelerate/LLVM/PTX/CodeGen/FoldSeg.hs --- a/Data/Array/Accelerate/LLVM/PTX/CodeGen/FoldSeg.hs +++ b/Data/Array/Accelerate/LLVM/PTX/CodeGen/FoldSeg.hs @@ -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 diff --git a/Data/Array/Accelerate/LLVM/PTX/CodeGen/Intrinsic.hs b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Intrinsic.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Intrinsic.hs @@ -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 +-- 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" + ] + diff --git a/Data/Array/Accelerate/LLVM/PTX/CodeGen/Queue.hs b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Queue.hs --- a/Data/Array/Accelerate/LLVM/PTX/CodeGen/Queue.hs +++ b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Queue.hs @@ -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 diff --git a/Data/Array/Accelerate/LLVM/PTX/CodeGen/Scan.hs b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Scan.hs --- a/Data/Array/Accelerate/LLVM/PTX/CodeGen/Scan.hs +++ b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Scan.hs @@ -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) diff --git a/Data/Array/Accelerate/LLVM/PTX/Compile.hs b/Data/Array/Accelerate/LLVM/PTX/Compile.hs --- a/Data/Array/Accelerate/LLVM/PTX/Compile.hs +++ b/Data/Array/Accelerate/LLVM/PTX/Compile.hs @@ -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) diff --git a/Data/Array/Accelerate/LLVM/PTX/Compile/Cache.hs b/Data/Array/Accelerate/LLVM/PTX/Compile/Cache.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Compile/Cache.hs @@ -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 +-- 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" + diff --git a/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice.hs b/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice.hs --- a/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice.hs +++ b/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice.hs @@ -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, .. } diff --git a/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice/Load.hs b/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice/Load.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice/Load.hs @@ -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 +-- 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 + diff --git a/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice/TH.hs b/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice/TH.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice/TH.hs @@ -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 +-- 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 ) + diff --git a/Data/Array/Accelerate/LLVM/PTX/Compile/Link.hs b/Data/Array/Accelerate/LLVM/PTX/Compile/Link.hs deleted file mode 100644 --- a/Data/Array/Accelerate/LLVM/PTX/Compile/Link.hs +++ /dev/null @@ -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 --- 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, .. } - diff --git a/Data/Array/Accelerate/LLVM/PTX/Embed.hs b/Data/Array/Accelerate/LLVM/PTX/Embed.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Embed.hs @@ -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 +-- 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 ||] + diff --git a/Data/Array/Accelerate/LLVM/PTX/Execute.hs b/Data/Array/Accelerate/LLVM/PTX/Execute.hs --- a/Data/Array/Accelerate/LLVM/PTX/Execute.hs +++ b/Data/Array/Accelerate/LLVM/PTX/Execute.hs @@ -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) diff --git a/Data/Array/Accelerate/LLVM/PTX/Execute/Async.hs b/Data/Array/Accelerate/LLVM/PTX/Execute/Async.hs --- a/Data/Array/Accelerate/LLVM/PTX/Execute/Async.hs +++ b/Data/Array/Accelerate/LLVM/PTX/Execute/Async.hs @@ -1,5 +1,4 @@ -{-# LANGUAGE RecordWildCards #-} -{-# LANGUAGE TypeFamilies #-} +{-# LANGUAGE TypeFamilies #-} {-# OPTIONS_GHC -fno-warn-orphans #-} -- | -- Module : Data.Array.Accelerate.LLVM.PTX.Execute.Async diff --git a/Data/Array/Accelerate/LLVM/PTX/Execute/Event.hs b/Data/Array/Accelerate/LLVM/PTX/Execute/Event.hs --- a/Data/Array/Accelerate/LLVM/PTX/Execute/Event.hs +++ b/Data/Array/Accelerate/LLVM/PTX/Execute/Event.hs @@ -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 diff --git a/Data/Array/Accelerate/LLVM/PTX/Execute/Marshal.hs b/Data/Array/Accelerate/LLVM/PTX/Execute/Marshal.hs --- a/Data/Array/Accelerate/LLVM/PTX/Execute/Marshal.hs +++ b/Data/Array/Accelerate/LLVM/PTX/Execute/Marshal.hs @@ -5,7 +5,6 @@ {-# LANGUAGE FlexibleInstances #-} {-# LANGUAGE GADTs #-} {-# LANGUAGE MultiParamTypeClasses #-} -{-# LANGUAGE RecordWildCards #-} {-# LANGUAGE ScopedTypeVariables #-} {-# LANGUAGE TypeFamilies #-} {-# OPTIONS_GHC -fno-warn-orphans #-} diff --git a/Data/Array/Accelerate/LLVM/PTX/Execute/Stream/Reservoir.hs b/Data/Array/Accelerate/LLVM/PTX/Execute/Stream/Reservoir.hs --- a/Data/Array/Accelerate/LLVM/PTX/Execute/Stream/Reservoir.hs +++ b/Data/Array/Accelerate/LLVM/PTX/Execute/Stream/Reservoir.hs @@ -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. diff --git a/Data/Array/Accelerate/LLVM/PTX/Foreign.hs b/Data/Array/Accelerate/LLVM/PTX/Foreign.hs --- a/Data/Array/Accelerate/LLVM/PTX/Foreign.hs +++ b/Data/Array/Accelerate/LLVM/PTX/Foreign.hs @@ -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 diff --git a/Data/Array/Accelerate/LLVM/PTX/Link.hs b/Data/Array/Accelerate/LLVM/PTX/Link.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Link.hs @@ -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 +-- 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 + ] +--} + diff --git a/Data/Array/Accelerate/LLVM/PTX/Link/Cache.hs b/Data/Array/Accelerate/LLVM/PTX/Link/Cache.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Link/Cache.hs @@ -0,0 +1,22 @@ +-- | +-- Module : Data.Array.Accelerate.LLVM.PTX.Link.Cache +-- Copyright : [2017] Trevor L. McDonell +-- License : BSD3 +-- +-- Maintainer : Trevor L. McDonell +-- 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 + diff --git a/Data/Array/Accelerate/LLVM/PTX/Link/Object.hs b/Data/Array/Accelerate/LLVM/PTX/Link/Object.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Link/Object.hs @@ -0,0 +1,41 @@ +-- | +-- Module : Data.Array.Accelerate.LLVM.PTX.Link.Object +-- Copyright : [2017] Trevor L. McDonell +-- License : BSD3 +-- +-- Maintainer : Trevor L. McDonell +-- 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 + diff --git a/Data/Array/Accelerate/LLVM/PTX/State.hs b/Data/Array/Accelerate/LLVM/PTX/State.hs --- a/Data/Array/Accelerate/LLVM/PTX/State.hs +++ b/Data/Array/Accelerate/LLVM/PTX/State.hs @@ -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 #-} diff --git a/Data/Array/Accelerate/LLVM/PTX/Target.hs b/Data/Array/Accelerate/LLVM/PTX/Target.hs --- a/Data/Array/Accelerate/LLVM/PTX/Target.hs +++ b/Data/Array/Accelerate/LLVM/PTX/Target.hs @@ -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 diff --git a/README.md b/README.md new file mode 100644 --- /dev/null +++ b/README.md @@ -0,0 +1,188 @@ +An LLVM backend for the Accelerate Array Language +================================================= + +[![Build Status](https://travis-ci.org/AccelerateHS/accelerate-llvm.svg)](https://travis-ci.org/AccelerateHS/accelerate-llvm) +[![Hackage](https://img.shields.io/hackage/v/accelerate-llvm.svg)](https://hackage.haskell.org/package/accelerate-llvm) +[![Docker Automated build](https://img.shields.io/docker/automated/tmcdonell/accelerate-llvm.svg)](https://hub.docker.com/r/tmcdonell/accelerate-llvm/) +[![Docker status](https://images.microbadger.com/badges/image/tmcdonell/accelerate-llvm.svg)](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 +``` + diff --git a/accelerate-llvm-ptx.cabal b/accelerate-llvm-ptx.cabal --- a/accelerate-llvm-ptx.cabal +++ b/accelerate-llvm-ptx.cabal @@ -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: - + refer to the main + package. + . + [/Dependencies/] + . + Haskell dependencies are available from Hackage. The following external + libraries are alse required: + . + * + . + * + . + [/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 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 + . 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