accelerate-llvm-ptx (empty) → 1.0.0.0
raw patch · 37 files changed
+7993/−0 lines, 37 filesdep +acceleratedep +accelerate-llvmdep +basesetup-changed
Dependencies added: accelerate, accelerate-llvm, base, bytestring, containers, cuda, directory, dlist, fclabels, filepath, hashable, llvm-hs, llvm-hs-pure, mtl, nvvm, pretty, time, unordered-containers
Files
- Data/Array/Accelerate/LLVM/PTX.hs +246/−0
- Data/Array/Accelerate/LLVM/PTX/Analysis/Device.hs +50/−0
- Data/Array/Accelerate/LLVM/PTX/Analysis/Launch.hs +72/−0
- Data/Array/Accelerate/LLVM/PTX/Array/Data.hs +213/−0
- Data/Array/Accelerate/LLVM/PTX/Array/Prim.hs +503/−0
- Data/Array/Accelerate/LLVM/PTX/Array/Remote.hs +160/−0
- Data/Array/Accelerate/LLVM/PTX/Array/Table.hs +60/−0
- Data/Array/Accelerate/LLVM/PTX/CodeGen.hs +46/−0
- Data/Array/Accelerate/LLVM/PTX/CodeGen/Base.hs +406/−0
- Data/Array/Accelerate/LLVM/PTX/CodeGen/Fold.hs +625/−0
- Data/Array/Accelerate/LLVM/PTX/CodeGen/FoldSeg.hs +461/−0
- Data/Array/Accelerate/LLVM/PTX/CodeGen/Generate.hs +60/−0
- Data/Array/Accelerate/LLVM/PTX/CodeGen/Loop.hs +47/−0
- Data/Array/Accelerate/LLVM/PTX/CodeGen/Map.hs +60/−0
- Data/Array/Accelerate/LLVM/PTX/CodeGen/Permute.hs +366/−0
- Data/Array/Accelerate/LLVM/PTX/CodeGen/Queue.hs +118/−0
- Data/Array/Accelerate/LLVM/PTX/CodeGen/Scan.hs +1332/−0
- Data/Array/Accelerate/LLVM/PTX/Compile.hs +289/−0
- Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice.hs +575/−0
- Data/Array/Accelerate/LLVM/PTX/Compile/Link.hs +183/−0
- Data/Array/Accelerate/LLVM/PTX/Context.hs +147/−0
- Data/Array/Accelerate/LLVM/PTX/Debug.hs +101/−0
- Data/Array/Accelerate/LLVM/PTX/Execute.hs +601/−0
- Data/Array/Accelerate/LLVM/PTX/Execute/Async.hs +64/−0
- Data/Array/Accelerate/LLVM/PTX/Execute/Environment.hs +22/−0
- Data/Array/Accelerate/LLVM/PTX/Execute/Event.hs +162/−0
- Data/Array/Accelerate/LLVM/PTX/Execute/Event.hs-boot +26/−0
- Data/Array/Accelerate/LLVM/PTX/Execute/Marshal.hs +148/−0
- Data/Array/Accelerate/LLVM/PTX/Execute/Stream.hs +181/−0
- Data/Array/Accelerate/LLVM/PTX/Execute/Stream.hs-boot +32/−0
- Data/Array/Accelerate/LLVM/PTX/Execute/Stream/Reservoir.hs +102/−0
- Data/Array/Accelerate/LLVM/PTX/Foreign.hs +79/−0
- Data/Array/Accelerate/LLVM/PTX/State.hs +108/−0
- Data/Array/Accelerate/LLVM/PTX/Target.hs +176/−0
- LICENSE +23/−0
- Setup.hs +2/−0
- accelerate-llvm-ptx.cabal +147/−0
+ Data/Array/Accelerate/LLVM/PTX.hs view
@@ -0,0 +1,246 @@+{-# LANGUAGE BangPatterns #-}+{-# LANGUAGE CPP #-}+{-# LANGUAGE FlexibleInstances #-}+{-# LANGUAGE GADTs #-}+{-# LANGUAGE TemplateHaskell #-}+{-# LANGUAGE TypeSynonymInstances #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX+-- Copyright : [2014..2017] Trevor L. McDonell+-- [2014..2014] Vinod Grover (NVIDIA Corporation)+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--+-- This module implements a backend for the /Accelerate/ language targeting+-- NVPTX for execution on NVIDIA GPUs. Expressions are on-line translated into+-- LLVM code, which is just-in-time executed in parallel on the GPU.+--++module Data.Array.Accelerate.LLVM.PTX (++ Acc, Arrays,++ -- * Synchronous execution+ run, runWith,+ run1, run1With,+ stream, streamWith,++ -- * Asynchronous execution+ Async,+ wait, poll, cancel,++ runAsync, runAsyncWith,+ run1Async, run1AsyncWith,++ -- * Execution targets+ PTX, createTargetForDevice, createTargetFromContext,++ -- * Controlling host-side allocation+ registerPinnedAllocator, registerPinnedAllocatorWith,++) where++-- accelerate+import Data.Array.Accelerate.Array.Sugar ( Arrays )+import Data.Array.Accelerate.Async+import Data.Array.Accelerate.Debug as Debug+import Data.Array.Accelerate.Error+import Data.Array.Accelerate.Smart ( Acc )+import Data.Array.Accelerate.Trafo++import Data.Array.Accelerate.LLVM.PTX.Compile+import Data.Array.Accelerate.LLVM.PTX.Execute+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 Foreign.CUDA.Driver as CUDA ( CUDAException, mallocHostForeignPtr )++-- standard library+import Control.Exception+import Control.Monad.Trans+import System.IO.Unsafe+import Text.Printf+++-- Accelerate: LLVM backend for NVIDIA GPUs+-- ----------------------------------------++-- | Compile and run a complete embedded array program.+--+-- Note that it is recommended that you use 'run1' 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.+--+-- Contexts passed to this function may all target to the same device, or to+-- separate devices of differing compute capabilities.+--+runWith :: Arrays a => PTX -> Acc a -> a+runWith target a+ = unsafePerformIO+ $ wait =<< runAsyncWith target a+++-- | As 'run', but run the computation asynchronously and return immediately+-- without waiting for the result. The status of the computation can be queried+-- using 'wait', 'poll', and 'cancel'.+--+-- Note that a CUDA context can be active on only one host thread at a time. If+-- you want to execute multiple computations in parallel, on the same or+-- different devices, use 'runAsyncWith'.+--+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.+--+runAsyncWith :: Arrays a => PTX -> Acc a -> IO (Async a)+runAsyncWith target a = asyncBound execute+ where+ !acc = convertAccWith config a+ execute = do+ dumpGraph acc+ evalPTX target $ do+ acc `seq` dumpSimplStats+ exec <- phase "compile" (compileAcc acc)+ res <- phase "execute" (executeAcc exec >>= AD.copyToHostLazy)+ return res+++-- | Prepare and execute an embedded array program of one argument.+--+-- 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.+-- 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.+--+-- At an example, once your program is expressed as a function of one argument,+-- instead of the usual:+--+-- > step :: Acc (Vector a) -> Acc (Vector b)+-- > step = ...+-- >+-- > simulate :: Vector a -> Vector b+-- > simulate xs = run $ step (use xs)+--+-- Instead write:+--+-- > simulate xs = run1 step xs+--+-- 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.+--+-- See the programs in the 'accelerate-examples' package for examples.+--+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 = run1' unsafePerformIO+++-- | 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++run1' :: (Arrays a, Arrays b) => (IO b -> c) -> PTX -> (Acc a -> Acc b) -> a -> c+run1' using target f = \a -> using (execute a)+ 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))+++-- | 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+++-- How the Accelerate program should be evaluated.+--+-- TODO: make sharing/fusion runtime configurable via debug flags or otherwise.+--+config :: Phase+config = phases+ { convertOffsetOfSegment = True+ }+++-- Controlling host-side allocation+-- --------------------------------++-- | Configure the default execution target to allocate all future host-side+-- arrays using (CUDA) pinned memory. Any newly allocated arrays will be+-- page-locked and directly accessible from the device, enabling high-speed+-- (asynchronous) DMA.+--+-- Note that since the amount of available pageable memory will be reduced,+-- overall system performance can suffer.+--+registerPinnedAllocator :: IO ()+registerPinnedAllocator = registerPinnedAllocatorWith defaultTarget+++-- | As with 'registerPinnedAllocator', but configure the given execution+-- context.+--+registerPinnedAllocatorWith :: PTX -> IO ()+registerPinnedAllocatorWith target =+ AD.registerForeignPtrAllocator $ \bytes ->+ CT.withContext (ptxContext target) (CUDA.mallocHostForeignPtr [] bytes)+ `catch`+ \e -> $internalError "registerPinnedAlocator" (show (e :: CUDAException))+++-- Debugging+-- =========++dumpStats :: MonadIO m => a -> m a+dumpStats x = dumpSimplStats >> return x++phase :: MonadIO m => String -> m a -> m a+phase n go = timed dump_phases (\wall cpu -> printf "phase %s: %s" n (elapsed wall cpu)) go+
+ Data/Array/Accelerate/LLVM/PTX/Analysis/Device.hs view
@@ -0,0 +1,50 @@+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Analysis.Device+-- Copyright : [2008..2017] Manuel M T Chakravarty, Gabriele Keller+-- [2009..2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Analysis.Device+ where++import Data.Ord+import Data.List+import Data.Function+import Foreign.CUDA.Driver.Device+import Foreign.CUDA.Analysis.Device+import qualified Foreign.CUDA.Driver as CUDA+++-- Select the best of the available CUDA capable devices. This prefers devices+-- with higher compute capability, followed by maximum throughput. This does not+-- take into account any other factors, such as whether the device is currently+-- in use by another process.+--+-- Ignore the possibility of emulation-mode devices, as this has been deprecated+-- as of CUDA v3.0 (compute-capability == 9999.9999)+--+selectBestDevice :: IO (Device, DeviceProperties)+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+ where+ compute = computeCapability+ flops d = multiProcessorCount d * coresPerMultiProcessor d * clockRate d+ cmp x y+ | compute x == compute y = comparing flops x y+ | otherwise = comparing compute x y+++-- Number of CUDA cores per streaming multiprocessor for a given architecture+-- revision. This is the number of SIMD arithmetic units per multiprocessor,+-- executing in lockstep in half-warp groupings (16 ALUs).+--+coresPerMultiProcessor :: DeviceProperties -> Int+coresPerMultiProcessor = coresPerMP . deviceResources+
+ Data/Array/Accelerate/LLVM/PTX/Analysis/Launch.hs view
@@ -0,0 +1,72 @@+{-# LANGUAGE RecordWildCards #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Analysis.Launch+-- Copyright : [2008..2017] Manuel M T Chakravarty, Gabriele Keller+-- [2009..2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Analysis.Launch (++ DeviceProperties, Occupancy, LaunchConfig,+ simpleLaunchConfig, launchConfig,+ multipleOf,++) where++-- library+import Foreign.CUDA.Analysis as CUDA+++-- 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+ -> ( Occupancy+ , Int -- thread block size+ , Int -> Int -- grid size required to process the given input size+ , Int -- #bytes dynamic shared memory+ )++-- | Analytics for a simple kernel which requires no additional shared memory or+-- have other constraints on launch configuration. The smallest thread block+-- 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+++-- | 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)+ -> LaunchConfig+launchConfig dev candidates dynamic_smem grid_size 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+ in+ ( occ, cta, grid, dynamic_smem cta )+++-- | The next highest multiple of 'y' from 'x'.+--+multipleOf :: Int -> Int -> Int+multipleOf x y = ((x + y - 1) `quot` y)+
+ Data/Array/Accelerate/LLVM/PTX/Array/Data.hs view
@@ -0,0 +1,213 @@+{-# LANGUAGE BangPatterns #-}+{-# LANGUAGE GADTs #-}+{-# LANGUAGE RecordWildCards #-}+{-# OPTIONS_GHC -fno-warn-orphans #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Array.Data+-- Copyright : [2014..2017] Trevor L. McDonell+-- [2014..2014] Vinod Grover (NVIDIA Corporation)+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Array.Data (++ module Data.Array.Accelerate.LLVM.Array.Data,+ module Data.Array.Accelerate.LLVM.PTX.Array.Data,++) where++-- accelerate+import Data.Array.Accelerate.Array.Sugar+import Data.Array.Accelerate.Array.Unique ( UniqueArray(..) )+import Data.Array.Accelerate.Lifetime ( Lifetime(..) )+import qualified Data.Array.Accelerate.Array.Representation as R++import Data.Array.Accelerate.LLVM.Array.Data+import Data.Array.Accelerate.LLVM.State++import Data.Array.Accelerate.LLVM.PTX.State+import Data.Array.Accelerate.LLVM.PTX.Target+import Data.Array.Accelerate.LLVM.PTX.Execute.Async+import qualified Data.Array.Accelerate.LLVM.PTX.Array.Prim as Prim++-- standard library+import Control.Applicative+import Control.Monad.State ( liftIO, gets )+import Data.Typeable+import Foreign.Ptr+import Foreign.Storable+import System.IO.Unsafe+import Prelude+++-- Instance of remote array memory management for the PTX target+--+instance Remote PTX where++ {-# INLINEABLE allocateRemote #-}+ allocateRemote !sh = do+ arr <- liftIO $ allocateArray sh+ runArray arr (\ad -> Prim.mallocArray (size sh) ad >> return ad)++ {-# INLINEABLE useRemoteR #-}+ useRemoteR !n !mst !ad = do+ case mst of+ Nothing -> Prim.useArray n ad+ Just st -> Prim.useArrayAsync st n ad++ {-# INLINEABLE copyToRemoteR #-}+ copyToRemoteR !from !to !mst !ad = do+ case mst of+ Nothing -> Prim.pokeArrayR from to ad+ Just st -> Prim.pokeArrayAsyncR st from to ad++ {-# INLINEABLE copyToHostR #-}+ copyToHostR !from !to !mst !ad = do+ case mst of+ Nothing -> Prim.peekArrayR from to ad+ Just st -> Prim.peekArrayAsyncR st from to ad++ {-# INLINEABLE copyToPeerR #-}+ copyToPeerR !from !to !dst !mst !ad = do+ case mst of+ Nothing -> Prim.copyArrayPeerR (ptxContext dst) (ptxMemoryTable dst) from to ad+ Just st -> Prim.copyArrayPeerAsyncR (ptxContext dst) (ptxMemoryTable dst) st from to ad++ {-# INLINEABLE indexRemote #-}+ indexRemote arr i =+ runIndexArray Prim.indexArray arr i+++-- | Copy an array from the remote device to the host. Although the Accelerate+-- program is hyper-strict and will evaluate the computation as soon as any part+-- of it is demanded, the individual array payloads are copied back to the host+-- _only_ as they are demanded by the Haskell program. This has several+-- consequences:+--+-- 1. If the device has multiple memcpy engines, only one will be used. The+-- transfers are however associated with a non-default stream.+--+-- 2. Using 'seq' to force an Array to head-normal form will initiate the+-- computation, but not transfer the results back to the host. Requesting+-- an array element or using 'deepseq' to force to normal form is required+-- to actually transfer the data.+--+copyToHostLazy+ :: Arrays arrs+ => arrs+ -> LLVM PTX arrs+copyToHostLazy arrs = do+ ptx <- gets llvmTarget+ liftIO $ runArrays arrs $ \(Array sh adata) ->+ let+ n :: Int+ n = R.size sh++ peekR :: (ArrayElt e, ArrayPtrs e ~ Ptr a, Storable a, Typeable a, Typeable e)+ => ArrayData e+ -> UniqueArray a+ -> IO (UniqueArray a)+ peekR ad (UniqueArray uid (Lifetime ref weak fp)) = do+ fp' <- unsafeInterleaveIO $+ evalPTX ptx $ do+ s <- fork+ copyToHostR 0 n (Just s) ad+ e <- checkpoint s+ block e+ join s+ return fp+ return $ UniqueArray uid (Lifetime ref weak fp')++ runR :: ArrayEltR e -> ArrayData e -> IO (ArrayData e)+ runR ArrayEltRunit AD_Unit = return AD_Unit+ runR (ArrayEltRpair aeR2 aeR1) (AD_Pair ad2 ad1) = AD_Pair <$> runR aeR2 ad2 <*> runR aeR1 ad1+ runR ArrayEltRint ad@(AD_Int ua) = AD_Int <$> peekR ad ua+ runR ArrayEltRint8 ad@(AD_Int8 ua) = AD_Int8 <$> peekR ad ua+ runR ArrayEltRint16 ad@(AD_Int16 ua) = AD_Int16 <$> peekR ad ua+ runR ArrayEltRint32 ad@(AD_Int32 ua) = AD_Int32 <$> peekR ad ua+ runR ArrayEltRint64 ad@(AD_Int64 ua) = AD_Int64 <$> peekR ad ua+ runR ArrayEltRword ad@(AD_Word ua) = AD_Word <$> peekR ad ua+ runR ArrayEltRword8 ad@(AD_Word8 ua) = AD_Word8 <$> peekR ad ua+ runR ArrayEltRword16 ad@(AD_Word16 ua) = AD_Word16 <$> peekR ad ua+ runR ArrayEltRword32 ad@(AD_Word32 ua) = AD_Word32 <$> peekR ad ua+ runR ArrayEltRword64 ad@(AD_Word64 ua) = AD_Word64 <$> peekR ad ua+ runR ArrayEltRcshort ad@(AD_CShort ua) = AD_CShort <$> peekR ad ua+ runR ArrayEltRcushort ad@(AD_CUShort ua) = AD_CUShort <$> peekR ad ua+ runR ArrayEltRcint ad@(AD_CInt ua) = AD_CInt <$> peekR ad ua+ runR ArrayEltRcuint ad@(AD_CUInt ua) = AD_CUInt <$> peekR ad ua+ runR ArrayEltRclong ad@(AD_CLong ua) = AD_CLong <$> peekR ad ua+ runR ArrayEltRculong ad@(AD_CULong ua) = AD_CULong <$> peekR ad ua+ runR ArrayEltRcllong ad@(AD_CLLong ua) = AD_CLLong <$> peekR ad ua+ runR ArrayEltRcullong ad@(AD_CULLong ua) = AD_CULLong <$> peekR ad ua+ runR ArrayEltRfloat ad@(AD_Float ua) = AD_Float <$> peekR ad ua+ runR ArrayEltRdouble ad@(AD_Double ua) = AD_Double <$> peekR ad ua+ runR ArrayEltRcfloat ad@(AD_CFloat ua) = AD_CFloat <$> peekR ad ua+ runR ArrayEltRcdouble ad@(AD_CDouble ua) = AD_CDouble <$> peekR ad ua+ runR ArrayEltRbool ad@(AD_Bool ua) = AD_Bool <$> peekR ad ua+ runR ArrayEltRchar ad@(AD_Char ua) = AD_Char <$> peekR ad ua+ runR ArrayEltRcchar ad@(AD_CChar ua) = AD_CChar <$> peekR ad ua+ runR ArrayEltRcschar ad@(AD_CSChar ua) = AD_CSChar <$> peekR ad ua+ runR ArrayEltRcuchar ad@(AD_CUChar ua) = AD_CUChar <$> peekR ad ua+ in+ Array sh <$> runR arrayElt adata+++-- | Clone an array into a newly allocated array on the device.+--+cloneArrayAsync+ :: (Shape sh, Elt e)+ => Stream+ -> Array sh e+ -> LLVM PTX (Array sh e)+cloneArrayAsync stream arr@(Array _ src) = do+ out@(Array _ dst) <- allocateRemote sh+ copyR arrayElt src dst+ return out+ where+ sh = shape arr+ n = size sh++ copyR :: ArrayEltR e -> ArrayData e -> ArrayData e -> LLVM PTX ()+ copyR ArrayEltRunit _ _ = return ()+ copyR (ArrayEltRpair aeR1 aeR2) ad1 ad2 = copyR aeR1 (fstArrayData ad1) (fstArrayData ad2) >>+ copyR aeR2 (sndArrayData ad1) (sndArrayData ad2)+ --+ copyR ArrayEltRint ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRint8 ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRint16 ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRint32 ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRint64 ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRword ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRword8 ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRword16 ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRword32 ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRword64 ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRfloat ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRdouble ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRbool ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRchar ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRcshort ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRcushort ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRcint ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRcuint ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRclong ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRculong ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRcllong ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRcullong ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRcfloat ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRcdouble ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRcchar ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRcschar ad1 ad2 = copyPrim ad1 ad2+ copyR ArrayEltRcuchar ad1 ad2 = copyPrim ad1 ad2++ copyPrim+ :: (ArrayElt e, ArrayPtrs e ~ Ptr a, Typeable e, Storable a, Typeable a)+ => ArrayData e+ -> ArrayData e+ -> LLVM PTX ()+ copyPrim a1 a2 = Prim.copyArrayAsync stream n a1 a2+
+ Data/Array/Accelerate/LLVM/PTX/Array/Prim.hs view
@@ -0,0 +1,503 @@+{-# LANGUAGE BangPatterns #-}+{-# LANGUAGE CPP #-}+{-# LANGUAGE DataKinds #-}+{-# LANGUAGE GADTs #-}+{-# LANGUAGE RecordWildCards #-}+{-# LANGUAGE ScopedTypeVariables #-}+{-# LANGUAGE TemplateHaskell #-}+{-# LANGUAGE TypeOperators #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Array.Prim+-- Copyright : [2014..2017] Trevor L. McDonell+-- [2014..2014] Vinod Grover (NVIDIA Corporation)+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Array.Prim (++ mallocArray,+ memsetArray, memsetArrayAsync,+ useArray, useArrayAsync,+ indexArray,+ peekArray, peekArrayR, peekArrayAsync, peekArrayAsyncR,+ pokeArray, pokeArrayR, pokeArrayAsync, pokeArrayAsyncR,+ copyArray, copyArrayR, copyArrayAsync, copyArrayAsyncR,+ copyArrayPeer, copyArrayPeerR, copyArrayPeerAsync, copyArrayPeerAsyncR,+ withDevicePtr,++) where++-- accelerate+import Data.Array.Accelerate.Array.Data+import Data.Array.Accelerate.Error+import Data.Array.Accelerate.Lifetime+import Data.Array.Accelerate.Type++import Data.Array.Accelerate.LLVM.State++import Data.Array.Accelerate.LLVM.PTX.Context+import Data.Array.Accelerate.LLVM.PTX.Target+import Data.Array.Accelerate.LLVM.PTX.Execute.Event+import Data.Array.Accelerate.LLVM.PTX.Execute.Stream+import Data.Array.Accelerate.LLVM.PTX.Array.Table+import Data.Array.Accelerate.LLVM.PTX.Array.Remote as Remote+import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug++-- CUDA+import qualified Foreign.CUDA.Driver as CUDA+import qualified Foreign.CUDA.Driver.Stream as CUDA++-- standard library+import Control.Exception+import Control.Monad+import Control.Monad.State+import Data.Typeable+import Foreign.Ptr+import Foreign.Storable+import GHC.TypeLits+import Text.Printf+import Prelude hiding ( lookup )+++-- | Allocate a device-side array associated with the given host array. If the+-- allocation fails due to a memory error, we attempt some last-ditch memory+-- cleanup before trying again.+--+{-# INLINEABLE mallocArray #-}+mallocArray+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, Storable a, Typeable a, Typeable e)+ => Int+ -> ArrayData e+ -> LLVM PTX ()+mallocArray !n !ad = do+ message ("mallocArray: " ++ showBytes (n * sizeOf (undefined::a)))+ void $ malloc ad n False+++-- | A combination of 'mallocArray' and 'pokeArray', that allocates remotes+-- memory and uploads an existing array. This is specialised because we tell the+-- allocator that the host-side array is frozen, and thus it is safe to evict+-- the remote memory and re-upload the data at any time.+--+{-# INLINEABLE useArray #-}+useArray+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, Storable a, Typeable a, Typeable e)+ => Int+ -> ArrayData e+ -> LLVM PTX ()+useArray !n !ad =+ blocking $ \st -> useArrayAsync st n ad++{-# INLINEABLE useArrayAsync #-}+useArrayAsync+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, Storable a, Typeable a, Typeable e)+ => Stream+ -> Int+ -> ArrayData e+ -> LLVM PTX ()+useArrayAsync !st !n !ad = do+ alloc <- malloc ad n True+ when alloc $ pokeArrayAsync st n ad+++-- | Copy data from the host to an existing array on the device+--+{-# INLINEABLE pokeArray #-}+pokeArray+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, Typeable e, Storable a, Typeable a)+ => Int+ -> ArrayData e+ -> LLVM PTX ()+pokeArray !n !ad =+ blocking $ \st -> pokeArrayAsync st n ad++{-# INLINEABLE pokeArrayAsync #-}+pokeArrayAsync+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, Typeable e, Storable a, Typeable a)+ => Stream+ -> Int+ -> ArrayData e+ -> LLVM PTX ()+pokeArrayAsync !stream !n !ad = do+ let !src = CUDA.HostPtr (ptrsOfArrayData ad)+ !bytes = n * sizeOf (undefined :: a)+ !st = unsafeGetValue stream+ --+ withDevicePtr ad $ \dst ->+ nonblocking stream $+ transfer "pokeArray" bytes (Just st) $ CUDA.pokeArrayAsync n src dst (Just st)+ liftIO (touchLifetime stream)+++{-# INLINEABLE pokeArrayR #-}+pokeArrayR+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, Typeable e, Typeable a, Storable a)+ => Int+ -> Int+ -> ArrayData e+ -> LLVM PTX ()+pokeArrayR !from !to !ad =+ blocking $ \st -> pokeArrayAsyncR st from to ad++{-# INLINEABLE pokeArrayAsyncR #-}+pokeArrayAsyncR+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, Typeable e, Typeable a, Storable a)+ => Stream+ -> Int+ -> Int+ -> ArrayData e+ -> LLVM PTX ()+pokeArrayAsyncR !stream !from !to !ad = do+ let !n = to - from+ !bytes = n * sizeOf (undefined :: a)+ !offset = from * sizeOf (undefined :: a)+ !src = CUDA.HostPtr (ptrsOfArrayData ad)+ !st = unsafeGetValue stream+ --+ withDevicePtr ad $ \dst ->+ nonblocking stream $+ transfer "pokeArray" bytes (Just st) $+ CUDA.pokeArrayAsync n (src `CUDA.plusHostPtr` offset) (dst `CUDA.plusDevPtr` offset) (Just st)+ liftIO (touchLifetime stream)+++-- | Read a single element from an array at a given row-major index+--+{-# INLINEABLE indexArray #-}+indexArray+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, Typeable e, Typeable a, Storable a)+ => ArrayData e+ -> Int+ -> LLVM PTX a+indexArray !ad !i =+ blocking $ \stream ->+ withDevicePtr ad $ \src -> liftIO $+ bracket (CUDA.mallocHostArray [] 1) CUDA.freeHost $ \dst -> do+ let !st = unsafeGetValue stream+ message $ "indexArray: " ++ showBytes (sizeOf (undefined::a))+ CUDA.peekArrayAsync 1 (src `CUDA.advanceDevPtr` i) dst (Just st)+ CUDA.block st+ touchLifetime stream+ r <- peek (CUDA.useHostPtr dst)+ return (Nothing, r)+++-- | Copy data from the device into the associated host-side Accelerate array+--+{-# INLINEABLE peekArray #-}+peekArray+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, Typeable e, Typeable a, Storable a)+ => Int+ -> ArrayData e+ -> LLVM PTX ()+peekArray !n !ad =+ blocking $ \st -> peekArrayAsync st n ad++{-# INLINEABLE peekArrayAsync #-}+peekArrayAsync+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, Typeable e, Typeable a, Storable a)+ => Stream+ -> Int+ -> ArrayData e+ -> LLVM PTX ()+peekArrayAsync !stream !n !ad = do+ let !bytes = n * sizeOf (undefined :: a)+ !dst = CUDA.HostPtr (ptrsOfArrayData ad)+ !st = unsafeGetValue stream+ --+ withDevicePtr ad $ \src ->+ nonblocking stream $+ transfer "peekArray" bytes (Just st) $ CUDA.peekArrayAsync n src dst (Just st)+ liftIO (touchLifetime stream)++{-# INLINEABLE peekArrayR #-}+peekArrayR+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, Typeable a, Typeable e, Storable a)+ => Int+ -> Int+ -> ArrayData e+ -> LLVM PTX ()+peekArrayR !from !to !ad =+ blocking $ \st -> peekArrayAsyncR st from to ad++{-# INLINEABLE peekArrayAsyncR #-}+peekArrayAsyncR+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, Typeable e, Typeable a, Storable a)+ => Stream+ -> Int+ -> Int+ -> ArrayData e+ -> LLVM PTX ()+peekArrayAsyncR !stream !from !to !ad = do+ let !n = to - from+ !bytes = n * sizeOf (undefined :: a)+ !offset = from * sizeOf (undefined :: a)+ !dst = CUDA.HostPtr (ptrsOfArrayData ad)+ !st = unsafeGetValue stream+ --+ withDevicePtr ad $ \src ->+ nonblocking stream $+ transfer "peekArray" bytes (Just st) $+ CUDA.peekArrayAsync n (src `CUDA.plusDevPtr` offset) (dst `CUDA.plusHostPtr` offset) (Just st)+ liftIO (touchLifetime stream)+++-- | Copy data between arrays in the same context+--+{-# INLINEABLE copyArray #-}+copyArray+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, Typeable e, Storable a, Typeable a)+ => Int+ -> ArrayData e+ -> ArrayData e+ -> LLVM PTX ()+copyArray !n !src !dst =+ blocking $ \st -> copyArrayAsync st n src dst++{-# INLINEABLE copyArrayAsync #-}+copyArrayAsync+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, Typeable e, Storable a, Typeable a)+ => Stream+ -> Int+ -> ArrayData e+ -> ArrayData e+ -> LLVM PTX ()+copyArrayAsync !stream !n !ad_src !ad_dst = do+ let !bytes = n * sizeOf (undefined :: a)+ !st = unsafeGetValue stream+ --+ withDevicePtr ad_src $ \src -> do+ e <- withDevicePtr ad_dst $ \dst -> do+ (e,()) <- nonblocking stream+ $ transfer "copyArray" bytes (Just st) $ CUDA.copyArrayAsync n src dst (Just st)+ return (e,e)+ return (e,())+ liftIO (touchLifetime stream)++{-# INLINEABLE copyArrayR #-}+copyArrayR+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, Typeable e, Storable a, Typeable a)+ => Int+ -> Int+ -> ArrayData e+ -> ArrayData e+ -> LLVM PTX ()+copyArrayR !from !to !src !dst =+ blocking $ \st -> copyArrayAsyncR st from to src dst++{-# INLINEABLE copyArrayAsyncR #-}+copyArrayAsyncR+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, Typeable e, Storable a, Typeable a)+ => Stream+ -> Int+ -> Int+ -> ArrayData e+ -> ArrayData e+ -> LLVM PTX ()+copyArrayAsyncR !stream !from !to !ad_src !ad_dst = do+ let !n = to - from+ !bytes = n * sizeOf (undefined :: a)+ !offset = from * sizeOf (undefined :: a)+ !st = unsafeGetValue stream+ --+ withDevicePtr ad_src $ \src -> do+ e <- withDevicePtr ad_dst $ \dst -> do+ (e,()) <- nonblocking stream+ $ transfer "copyArray" bytes (Just st)+ $ CUDA.copyArrayAsync n (src `CUDA.plusDevPtr` offset) (dst `CUDA.plusDevPtr` offset) (Just st)+ return (e,e)+ return (e,())+ liftIO (touchLifetime stream)+++-- | Copy data from one device context into a _new_ array on the second context.+-- It is an error if the destination array already exists.+--+{-# INLINEABLE copyArrayPeer #-}+copyArrayPeer+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, Storable a, Typeable a)+ => Context -- destination context+ -> MemoryTable -- destination memory table+ -> Int+ -> ArrayData e+ -> LLVM PTX ()+copyArrayPeer !ctx2 !mt2 !n !ad =+ blocking $ \st -> copyArrayPeerAsync ctx2 mt2 st n ad++{-# INLINEABLE copyArrayPeerAsync #-}+copyArrayPeerAsync+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, Storable a, Typeable a)+ => Context -- destination context+ -> MemoryTable -- destination memory table+ -> Stream+ -> Int+ -> ArrayData e+ -> LLVM PTX ()+copyArrayPeerAsync = error "copyArrayPeerAsync"+{--+copyArrayPeerAsync !ctx2 !mt2 !st !n !ad = do+ let !bytes = n * sizeOf (undefined :: a)+ src <- devicePtr mt1 ad+ dst <- mallocArray ctx2 mt2 n ad+ transfer "copyArrayPeer" bytes (Just st) $+ CUDA.copyArrayPeerAsync n src (deviceContext ctx1) dst (deviceContext ctx2) (Just st)+--}++-- | Copy part of an array from one device context to another. Both source and+-- destination arrays must exist.+--+{-# INLINEABLE copyArrayPeerR #-}+copyArrayPeerR+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, Storable a, Typeable a)+ => Context -- destination context+ -> MemoryTable -- destination memory table+ -> Int+ -> Int+ -> ArrayData e+ -> LLVM PTX ()+copyArrayPeerR !ctx2 !mt2 !from !to !ad =+ blocking $ \st -> copyArrayPeerAsyncR ctx2 mt2 st from to ad++{-# INLINEABLE copyArrayPeerAsyncR #-}+copyArrayPeerAsyncR+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, Storable a, Typeable a)+ => Context -- destination context+ -> MemoryTable -- destination memory table+ -> Stream+ -> Int+ -> Int+ -> ArrayData e+ -> LLVM PTX ()+copyArrayPeerAsyncR = error "copyArrayPeerAsyncR"+{--+copyArrayPeerAsyncR !ctx2 !mt2 !st !from !to !ad = do+ let !n = to - from+ !bytes = n * sizeOf (undefined :: a)+ !offset = from * sizeOf (undefined :: a)+ src <- devicePtr mt1 ad :: IO (CUDA.DevicePtr a)+ dst <- devicePtr mt2 ad :: IO (CUDA.DevicePtr a)+ transfer "copyArrayPeer" bytes (Just st) $+ CUDA.copyArrayPeerAsync n (src `CUDA.plusDevPtr` offset) (deviceContext ctx1)+ (dst `CUDA.plusDevPtr` offset) (deviceContext ctx2) (Just st)+--}+++-- | Set elements of the array to the specified value. Only 8-, 16-, and 32-bit+-- values are supported.+--+{-# INLINEABLE memsetArray #-}+memsetArray+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, Typeable e, Typeable a, Storable a, BitSize a <= 32)+ => Int+ -> a+ -> ArrayData e+ -> LLVM PTX ()+memsetArray !n !v !ad =+ blocking $ \st -> memsetArrayAsync st n v ad++{-# INLINEABLE memsetArrayAsync #-}+memsetArrayAsync+ :: forall e a. (ArrayElt e, ArrayPtrs e ~ Ptr a, Typeable e, Typeable a, Storable a, BitSize a <= 32)+ => Stream+ -> Int+ -> a+ -> ArrayData e+ -> LLVM PTX ()+memsetArrayAsync !stream !n !v !ad = do+ let !bytes = n * sizeOf (undefined :: a)+ !st = unsafeGetValue stream+ --+ withDevicePtr ad $ \ptr ->+ nonblocking stream $+ transfer "memset" bytes (Just st) $ CUDA.memsetAsync ptr n v (Just st)+ liftIO (touchLifetime stream)+++{--+-- | Lookup the device memory associated with a given host array+--+{-# INLINEABLE devicePtr #-}+devicePtr+ :: (ArrayElt e, ArrayPtrs e ~ Ptr a, Typeable a, Typeable b)+ => ArrayData e+ -> LLVM PTX (CUDA.DevicePtr b)+devicePtr !ad = do+ undefined+ {--+ mv <- Table.lookup mt ad+ case mv of+ Just v -> return v+ Nothing -> $internalError "devicePtr" "lost device memory"+ --}+--}++-- Auxiliary+-- ---------++-- | Lookup the device memory associated with a given host array and do+-- something with it.+--+{-# INLINEABLE withDevicePtr #-}+withDevicePtr+ :: (ArrayElt e, ArrayPtrs e ~ Ptr a, Typeable e, Typeable a, Storable a)+ => ArrayData e+ -> (CUDA.DevicePtr a -> LLVM PTX (Maybe Event, r))+ -> LLVM PTX r+withDevicePtr !ad !f = do+ mr <- withRemote ad f+ case mr of+ Nothing -> $internalError "withDevicePtr" "array does not exist on the device"+ Just r -> return r++-- | Execute the given operation in a new stream, and wait for the operation to+-- complete before returning.+--+{-# INLINE blocking #-}+blocking :: (Stream -> LLVM PTX a) -> LLVM PTX a+blocking !f =+ streaming f $ \e r -> do+ liftIO $ block e+ return r++-- | Execute a (presumable asynchronous) operation and return the result+-- together with an event recorded immediately afterwards in the given stream.+--+{-# INLINE nonblocking #-}+nonblocking :: Stream -> LLVM PTX a -> LLVM PTX (Maybe Event, a)+nonblocking !stream !f = do+ r <- f+ e <- waypoint stream+ return (Just e, r)+++-- Debug+-- -----++{-# INLINE showBytes #-}+showBytes :: Int -> String+showBytes x = Debug.showFFloatSIBase (Just 0) 1024 (fromIntegral x :: Double) "B"++{-# INLINE trace #-}+trace :: MonadIO m => String -> m a -> m a+trace msg next = liftIO (Debug.traceIO Debug.dump_gc ("gc: " ++ msg)) >> next++{-# INLINE message #-}+message :: MonadIO m => String -> m ()+message s = s `trace` return ()++{-# INLINE transfer #-}+transfer :: MonadIO m => String -> Int -> Maybe CUDA.Stream -> IO () -> m ()+transfer name bytes stream action+ = let showRate x t = Debug.showFFloatSIBase (Just 3) 1024 (fromIntegral x / t) "B/s"+ msg wall cpu gpu = printf "gc: %s: %s bytes @ %s, %s"+ name+ (showBytes bytes)+ (showRate bytes wall)+ (Debug.elapsed wall cpu gpu)+ in+ liftIO (Debug.timed Debug.dump_gc msg stream action)+
+ Data/Array/Accelerate/LLVM/PTX/Array/Remote.hs view
@@ -0,0 +1,160 @@+{-# LANGUAGE BangPatterns #-}+{-# LANGUAGE FlexibleInstances #-}+{-# LANGUAGE RecordWildCards #-}+{-# LANGUAGE ScopedTypeVariables #-}+{-# LANGUAGE TypeFamilies #-}+{-# OPTIONS_GHC -fno-warn-orphans #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Array.Remote+-- Copyright : [2014..2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Array.Remote (++ withRemote, malloc,++) where++import Data.Array.Accelerate.LLVM.State+import Data.Array.Accelerate.LLVM.PTX.Target+import {-# SOURCE #-} Data.Array.Accelerate.LLVM.PTX.Execute.Event+import {-# SOURCE #-} Data.Array.Accelerate.LLVM.PTX.Execute.Stream++import Data.Array.Accelerate.Lifetime+import Data.Array.Accelerate.Array.Data+import qualified Data.Array.Accelerate.Array.Remote as Remote+import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug++import Foreign.CUDA.Driver.Error+import qualified Foreign.CUDA.Ptr as CUDA+import qualified Foreign.CUDA.Driver as CUDA+import qualified Foreign.CUDA.Driver.Stream as CUDA++import Control.Exception+import Control.Monad.State+import Data.Typeable+import Foreign.Ptr+import Foreign.Storable+import Text.Printf+++-- Events signal once a computation has completed+--+instance Remote.Task (Maybe Event) where+ completed Nothing = return True+ completed (Just e) = query e++instance Remote.RemoteMemory (LLVM PTX) where+ type RemotePtr (LLVM PTX) = CUDA.DevicePtr+ --+ mallocRemote n+ | n <= 0 = return (Just CUDA.nullDevPtr)+ | otherwise = liftIO $ do+ ep <- try (CUDA.mallocArray n)+ case ep of+ Right p -> return (Just p)+ Left (ExitCode OutOfMemory) -> return Nothing+ Left e -> do message ("malloc failed with error: " ++ show e)+ throwIO e++ peekRemote n src ad =+ let bytes = n * sizeOfPtr src+ dst = CUDA.HostPtr (ptrsOfArrayData ad)+ in+ blocking $ \stream ->+ withLifetime stream $ \st ->+ transfer "peekRemote" bytes (Just st) $ CUDA.peekArrayAsync n src dst (Just st)++ pokeRemote n dst ad =+ let bytes = n * sizeOfPtr dst+ src = CUDA.HostPtr (ptrsOfArrayData ad)+ in+ blocking $ \stream ->+ withLifetime stream $ \st ->+ transfer "pokeRemote" bytes (Just st) $ CUDA.pokeArrayAsync n src dst (Just st)++ castRemotePtr _ = CUDA.castDevPtr+ availableRemoteMem = liftIO $ fst `fmap` CUDA.getMemInfo+ totalRemoteMem = liftIO $ snd `fmap` CUDA.getMemInfo+ remoteAllocationSize = return 4096++++-- | Allocate an array in the remote memory space sufficient to hold the given+-- number of elements, and associated with the given host side array. Space will+-- be freed from the remote device if necessary.+--+{-# INLINEABLE malloc #-}+malloc+ :: (ArrayElt e, ArrayPtrs e ~ Ptr a, Typeable e, Typeable a, Storable a)+ => ArrayData e+ -> Int+ -> Bool+ -> LLVM PTX Bool+malloc !ad !n !frozen = do+ PTX{..} <- gets llvmTarget+ Remote.malloc ptxMemoryTable ad frozen n+++-- | Lookup up the remote array pointer for the given host-side array+--+{-# INLINEABLE withRemote #-}+withRemote+ :: (ArrayElt e, ArrayPtrs e ~ Ptr a, Typeable e, Typeable a, Storable a)+ => ArrayData e+ -> (CUDA.DevicePtr a -> LLVM PTX (Maybe Event, r))+ -> LLVM PTX (Maybe r)+withRemote !ad !f = do+ PTX{..} <- gets llvmTarget+ Remote.withRemote ptxMemoryTable ad f+++-- Auxiliary+-- ---------++-- | Execute the given operation in a new stream, and wait for the operation to+-- complete before returning.+--+{-# INLINE blocking #-}+blocking :: (Stream -> IO a) -> LLVM PTX a+blocking !fun =+ streaming (liftIO . fun) $ \e r -> do+ liftIO $ block e+ return r++{-# INLINE sizeOfPtr #-}+sizeOfPtr :: forall a. Storable a => CUDA.DevicePtr a -> Int+sizeOfPtr _ = sizeOf (undefined :: a)++-- Debugging+-- ---------++{-# INLINE showBytes #-}+showBytes :: Int -> String+showBytes x = Debug.showFFloatSIBase (Just 0) 1024 (fromIntegral x :: Double) "B"++{-# INLINE trace #-}+trace :: String -> IO a -> IO a+trace msg next = Debug.traceIO Debug.dump_gc ("gc: " ++ msg) >> next++{-# INLINE message #-}+message :: String -> IO ()+message s = s `trace` return ()++{-# INLINE transfer #-}+transfer :: String -> Int -> Maybe CUDA.Stream -> IO () -> IO ()+transfer name bytes stream action+ = let showRate x t = Debug.showFFloatSIBase (Just 3) 1024 (fromIntegral x / t) "B/s"+ msg wall cpu gpu = printf "gc: %s: %s bytes @ %s, %s"+ name+ (showBytes bytes)+ (showRate bytes wall)+ (Debug.elapsed wall cpu gpu)+ in+ Debug.timed Debug.dump_gc msg stream action+
+ Data/Array/Accelerate/LLVM/PTX/Array/Table.hs view
@@ -0,0 +1,60 @@+{-# LANGUAGE BangPatterns #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Array.Table+-- Copyright : [2014..2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Array.Table (++ MemoryTable,+ new,++) where++import Data.Array.Accelerate.LLVM.PTX.Context ( Context, withContext )+import qualified Data.Array.Accelerate.Array.Remote as Remote+import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug+import {-# SOURCE #-} Data.Array.Accelerate.LLVM.PTX.Execute.Event++import qualified Foreign.CUDA.Ptr as CUDA+import qualified Foreign.CUDA.Driver as CUDA++import Text.Printf+++-- Remote memory tables. This builds upon the LRU-cached memory tables provided+-- by the base Accelerate package.+--+type MemoryTable = Remote.MemoryTable CUDA.DevicePtr (Maybe Event)+++-- | Create a new PTX memory table. This is specific to a given PTX target, as+-- devices arrays are unique to a CUDA context.+--+{-# INLINEABLE new #-}+new :: Context -> IO MemoryTable+new !ctx = Remote.new freeRemote+ where+ freeRemote :: CUDA.DevicePtr a -> IO ()+ freeRemote !ptr = do+ message (printf "freeRemote %s" (show ptr))+ withContext ctx (CUDA.free ptr)+++-- Debugging+-- ---------++{-# INLINE trace #-}+trace :: String -> IO a -> IO a+trace msg next = Debug.traceIO Debug.dump_gc ("gc: " ++ msg) >> next++{-# INLINE message #-}+message :: String -> IO ()+message s = s `trace` return ()++
+ Data/Array/Accelerate/LLVM/PTX/CodeGen.hs view
@@ -0,0 +1,46 @@+{-# OPTIONS_GHC -fno-warn-orphans #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.CodeGen+-- Copyright : [2014..2017] Trevor L. McDonell+-- [2014..2014] Vinod Grover (NVIDIA Corporation)+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.CodeGen (++ KernelMetadata(..),++) where++-- accelerate+import Data.Array.Accelerate.LLVM.CodeGen++import Data.Array.Accelerate.LLVM.PTX.CodeGen.Base+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.Map+import Data.Array.Accelerate.LLVM.PTX.CodeGen.Permute+import Data.Array.Accelerate.LLVM.PTX.CodeGen.Scan+import Data.Array.Accelerate.LLVM.PTX.Target+++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+
+ Data/Array/Accelerate/LLVM/PTX/CodeGen/Base.hs view
@@ -0,0 +1,406 @@+{-# LANGUAGE GADTs #-}+{-# LANGUAGE OverloadedStrings #-}+{-# LANGUAGE ScopedTypeVariables #-}+{-# LANGUAGE TemplateHaskell #-}+{-# LANGUAGE TypeFamilies #-}+{-# LANGUAGE ViewPatterns #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.CodeGen.Base+-- Copyright : [2014..2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.CodeGen.Base (++ -- Types+ DeviceProperties, KernelMetadata(..),++ -- Thread identifiers+ blockDim, gridDim, threadIdx, blockIdx, warpSize,+ gridSize, globalThreadIdx,+ gangParam,++ -- Other intrinsics+ laneId, warpId,+ laneMask_eq, laneMask_lt, laneMask_le, laneMask_gt, laneMask_ge,+ atomicAdd_f,++ -- Barriers and synchronisation+ __syncthreads,+ __threadfence_block, __threadfence_grid,++ -- Shared memory+ staticSharedMem,+ dynamicSharedMem,+ sharedMemAddrSpace,++ -- Kernel definitions+ (+++),+ makeOpenAcc, makeOpenAccWith,++) where++-- llvm+import LLVM.AST.Type.AddrSpace+import LLVM.AST.Type.Constant+import LLVM.AST.Type.Global+import LLVM.AST.Type.Instruction+import LLVM.AST.Type.Instruction.Volatile+import LLVM.AST.Type.Metadata+import LLVM.AST.Type.Name+import LLVM.AST.Type.Operand+import LLVM.AST.Type.Representation+import qualified LLVM.AST.Global as LLVM+import qualified LLVM.AST.Linkage as LLVM+import qualified LLVM.AST.Name as LLVM+import qualified LLVM.AST.Type as LLVM++-- accelerate+import Data.Array.Accelerate.Analysis.Type+import Data.Array.Accelerate.Array.Sugar ( Elt, Vector, eltType )+import Data.Array.Accelerate.Error++import Data.Array.Accelerate.LLVM.CodeGen.Arithmetic as A+import Data.Array.Accelerate.LLVM.CodeGen.Base+import Data.Array.Accelerate.LLVM.CodeGen.Constant+import Data.Array.Accelerate.LLVM.CodeGen.Downcast+import Data.Array.Accelerate.LLVM.CodeGen.IR+import Data.Array.Accelerate.LLVM.CodeGen.Module+import Data.Array.Accelerate.LLVM.CodeGen.Monad+import Data.Array.Accelerate.LLVM.CodeGen.Ptr+import Data.Array.Accelerate.LLVM.CodeGen.Sugar+import Data.Array.Accelerate.LLVM.CodeGen.Type++import Data.Array.Accelerate.LLVM.PTX.Analysis.Launch+import Data.Array.Accelerate.LLVM.PTX.Context+import Data.Array.Accelerate.LLVM.PTX.Target++-- standard library+import Control.Applicative+import Control.Monad ( void )+import Text.Printf+import Prelude as P+++-- Thread identifiers+-- ------------------++-- | Read the builtin registers that store CUDA thread and grid identifiers+--+-- <https://github.com/llvm-mirror/llvm/blob/master/include/llvm/IR/IntrinsicsNVVM.td>+--+specialPTXReg :: Label -> CodeGen (IR Int32)+specialPTXReg f =+ call (Body type' f) [NoUnwind, ReadNone]++blockDim, gridDim, threadIdx, blockIdx, warpSize :: CodeGen (IR Int32)+blockDim = specialPTXReg "llvm.nvvm.read.ptx.sreg.ntid.x"+gridDim = specialPTXReg "llvm.nvvm.read.ptx.sreg.nctaid.x"+threadIdx = specialPTXReg "llvm.nvvm.read.ptx.sreg.tid.x"+blockIdx = specialPTXReg "llvm.nvvm.read.ptx.sreg.ctaid.x"+warpSize = specialPTXReg "llvm.nvvm.read.ptx.sreg.warpsize"++laneId :: CodeGen (IR Int32)+laneId = specialPTXReg "llvm.nvvm.read.ptx.sreg.laneid"++laneMask_eq, laneMask_lt, laneMask_le, laneMask_gt, laneMask_ge :: CodeGen (IR Int32)+laneMask_eq = specialPTXReg "llvm.nvvm.read.ptx.sreg.lanemask.eq"+laneMask_lt = specialPTXReg "llvm.nvvm.read.ptx.sreg.lanemask.lt"+laneMask_le = specialPTXReg "llvm.nvvm.read.ptx.sreg.lanemask.le"+laneMask_gt = specialPTXReg "llvm.nvvm.read.ptx.sreg.lanemask.gt"+laneMask_ge = specialPTXReg "llvm.nvvm.read.ptx.sreg.lanemask.ge"+++-- | NOTE: The special register %warpid as volatile value and is not guaranteed+-- to be constant over the lifetime of a thread or thread block.+--+-- http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#sm-id-and-warp-id+--+-- http://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-warpid+--+-- We might consider passing in the (constant) warp size from device properties,+-- so that the division can be optimised to a shift.+--+warpId :: CodeGen (IR Int32)+warpId = do+ tid <- threadIdx+ ws <- warpSize+ A.quot integralType tid ws++_warpId :: CodeGen (IR Int32)+_warpId = specialPTXReg "llvm.ptx.read.warpid"+++-- | The size of the thread grid+--+-- > gridDim.x * blockDim.x+--+gridSize :: CodeGen (IR Int32)+gridSize = do+ ncta <- gridDim+ nt <- blockDim+ mul numType ncta nt+++-- | The global thread index+--+-- > blockDim.x * blockIdx.x + threadIdx.x+--+globalThreadIdx :: CodeGen (IR Int32)+globalThreadIdx = do+ ntid <- blockDim+ ctaid <- blockIdx+ tid <- threadIdx+ --+ u <- mul numType ntid ctaid+ v <- add numType tid u+ return v+++-- | Generate function parameters that will specify the first and last (linear)+-- index of the array this kernel should evaluate.+--+gangParam :: (IR Int32, IR Int32, [LLVM.Parameter])+gangParam =+ let t = scalarType+ start = "ix.start"+ end = "ix.end"+ in+ (local t start, local t end, [ scalarParameter t start, scalarParameter t end ] )+++-- Barriers and synchronisation+-- ----------------------------++-- | Call a builtin CUDA synchronisation intrinsic+--+barrier :: Label -> CodeGen ()+barrier f = void $ call (Body VoidType f) [NoUnwind, NoDuplicate, Convergent]+++-- | Wait until all threads in the thread block have reached this point and all+-- global and shared memory accesses made by these threads prior to+-- __syncthreads() are visible to all threads in the block.+--+-- <http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#synchronization-functions>+--+__syncthreads :: CodeGen ()+__syncthreads = barrier "llvm.nvvm.barrier0"+++-- | Ensure that all writes to shared and global memory before the call to+-- __threadfence_block() are observed by all threads in the *block* of the+-- calling thread as occurring before all writes to shared and global memory+-- made by the calling thread after the call.+--+-- <http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#memory-fence-functions>+--+__threadfence_block :: CodeGen ()+__threadfence_block = barrier "llvm.nvvm.membar.cta"+++-- | As __threadfence_block(), but the synchronisation is for *all* thread blocks.+-- In CUDA this is known simply as __threadfence().+--+__threadfence_grid :: CodeGen ()+__threadfence_grid = barrier "llvm.nvvm.membar.gl"+++-- Atomic functions+-- ----------------++-- LLVM provides atomic instructions for integer arguments only. CUDA provides+-- additional support for atomic add on floating point types, which can be+-- accessed through the following intrinsics.+--+-- Double precision is only supported on Compute 6.0 devices and later. LLVM-4.0+-- currently lacks support for this intrinsic, however it may be possible to use+-- inline assembly.+--+-- <https://github.com/AccelerateHS/accelerate/issues/363>+--+atomicAdd_f :: FloatingType a -> Operand (Ptr a) -> Operand a -> CodeGen ()+atomicAdd_f t addr val =+ let+ width :: Int+ width =+ case t of+ TypeFloat{} -> 32+ TypeDouble{} -> 64+ TypeCFloat{} -> 32+ TypeCDouble{} -> 64++ addrspace :: Word32+ (t_addr, t_val, addrspace) =+ case typeOf addr of+ PrimType ta@(PtrPrimType (ScalarPrimType tv) (AddrSpace as))+ -> (ta, tv, as)+ _ -> $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+ in+ void $ call (Lam t_addr addr (Lam (ScalarPrimType t_val) val (Body t_ret fun))) [NoUnwind]+++-- Shared memory+-- -------------++sharedMemAddrSpace :: AddrSpace+sharedMemAddrSpace = AddrSpace 3++sharedMemVolatility :: Volatility+sharedMemVolatility = Volatile+++-- Declare a new statically allocated array in the __shared__ memory address+-- space, with enough storage to contain the given number of elements.+--+staticSharedMem+ :: forall e. Elt e+ => Word64+ -> CodeGen (IRArray (Vector e))+staticSharedMem n = do+ ad <- go (eltType (undefined::e))+ return $ IRArray { irArrayShape = IR (OP_Pair OP_Unit (OP_Int (integral integralType (P.fromIntegral n))))+ , irArrayData = IR ad+ , irArrayAddrSpace = sharedMemAddrSpace+ , irArrayVolatility = sharedMemVolatility+ }+ where+ go :: TupleType s -> CodeGen (Operands s)+ go UnitTuple = return OP_Unit+ go (PairTuple t1 t2) = OP_Pair <$> go t1 <*> go t2+ go tt@(SingleTuple t) = do+ -- Declare a new global reference for the statically allocated array+ -- located in the __shared__ memory space.+ nm <- freshName+ sm <- return $ ConstantOperand $ GlobalReference (PrimType (PtrPrimType (ArrayType n t) sharedMemAddrSpace)) nm+ declare $ LLVM.globalVariableDefaults+ { LLVM.addrSpace = sharedMemAddrSpace+ , LLVM.type' = LLVM.ArrayType n (downcast t)+ , LLVM.linkage = LLVM.Internal+ , LLVM.name = downcast nm+ , LLVM.alignment = 4 `P.max` P.fromIntegral (sizeOf tt)+ }++ -- Return a pointer to the first element of the __shared__ memory array.+ -- We do this rather than just returning the global reference directly due+ -- to how __shared__ memory needs to be indexed with the GEP instruction.+ p <- instr' $ GetElementPtr sm [num numType 0, num numType 0 :: Operand Int32]+ q <- instr' $ PtrCast (PtrPrimType (ScalarPrimType t) sharedMemAddrSpace) p++ return $ ir' t (unPtr q)+++-- External declaration in shared memory address space. This must be declared in+-- order to access memory allocated dynamically by the CUDA driver. This results+-- in the following global declaration:+--+-- > @__shared__ = external addrspace(3) global [0 x i8]+--+initialiseDynamicSharedMemory :: CodeGen (Operand (Ptr Word8))+initialiseDynamicSharedMemory = do+ declare $ LLVM.globalVariableDefaults+ { LLVM.addrSpace = sharedMemAddrSpace+ , LLVM.type' = LLVM.ArrayType 0 (LLVM.IntegerType 8)+ , LLVM.linkage = LLVM.External+ , LLVM.name = LLVM.Name "__shared__"+ , LLVM.alignment = 4+ }+ return $ ConstantOperand $ GlobalReference type' "__shared__"+++-- Declared a new dynamically allocated array in the __shared__ memory space+-- with enough space to contain the given number of elements.+--+dynamicSharedMem+ :: forall e int. (Elt e, IsIntegral int)+ => IR int -- number of array elements+ -> IR int -- #bytes of shared memory the have already been allocated+ -> CodeGen (IRArray (Vector e))+dynamicSharedMem n@(op integralType -> m) (op integralType -> offset) = do+ smem <- initialiseDynamicSharedMemory+ let+ go :: TupleType s -> Operand int -> CodeGen (Operand int, Operands s)+ go UnitTuple i = return (i, OP_Unit)+ go (PairTuple t2 t1) i0 = do+ (i1, p1) <- go t1 i0+ (i2, p2) <- go t2 i1+ return $ (i2, OP_Pair p2 p1)+ go (SingleTuple t) i = do+ p <- instr' $ GetElementPtr smem [num numType 0, i] -- TLM: note initial zero index!!+ q <- instr' $ PtrCast (PtrPrimType (ScalarPrimType t) sharedMemAddrSpace) p+ a <- instr' $ Mul numType m (integral integralType (P.fromIntegral (sizeOf (SingleTuple t))))+ b <- instr' $ Add numType i a+ return (b, ir' t (unPtr q))+ --+ (_, ad) <- go (eltType (undefined::e)) offset+ IR sz <- A.fromIntegral integralType (numType :: NumType Int) n+ return $ IRArray { irArrayShape = IR $ OP_Pair OP_Unit sz+ , irArrayData = IR ad+ , irArrayAddrSpace = sharedMemAddrSpace+ , irArrayVolatility = sharedMemVolatility+ }+++-- Global kernel definitions+-- -------------------------++data instance KernelMetadata PTX = KM_PTX LaunchConfig++-- | Combine kernels into a single program+--+(+++) :: IROpenAcc PTX aenv a -> IROpenAcc PTX aenv a -> IROpenAcc PTX aenv a+IROpenAcc k1 +++ IROpenAcc k2 = IROpenAcc (k1 ++ k2)+++-- | Create a single kernel program with the default launch configuration.+--+makeOpenAcc+ :: PTX+ -> Label+ -> [LLVM.Parameter]+ -> CodeGen ()+ -> CodeGen (IROpenAcc PTX aenv a)+makeOpenAcc (deviceProperties . ptxContext -> dev) =+ makeOpenAccWith (simpleLaunchConfig dev)++-- | Create a single kernel program with the given launch analysis information.+--+makeOpenAccWith+ :: LaunchConfig+ -> Label+ -> [LLVM.Parameter]+ -> CodeGen ()+ -> CodeGen (IROpenAcc PTX aenv a)+makeOpenAccWith config name param kernel = do+ body <- makeKernel config name param kernel+ return $ IROpenAcc [body]++-- | Create a complete kernel function by running the code generation process+-- specified in the final parameter.+--+makeKernel :: LaunchConfig -> Label -> [LLVM.Parameter] -> CodeGen () -> CodeGen (Kernel PTX aenv a)+makeKernel config name@(Label l) param kernel = do+ _ <- kernel+ code <- createBlocks+ addMetadata "nvvm.annotations"+ [ Just . MetadataOperand $ ConstantOperand (GlobalReference VoidType (Name l))+ , Just . MetadataStringOperand $ "kernel"+ , Just . MetadataOperand $ scalar scalarType (1::Int)+ ]+ return $ Kernel+ { kernelMetadata = KM_PTX config+ , unKernel = LLVM.functionDefaults+ { LLVM.returnType = LLVM.VoidType+ , LLVM.name = downcast name+ , LLVM.parameters = (param, False)+ , LLVM.basicBlocks = code+ }+ }+
+ Data/Array/Accelerate/LLVM/PTX/CodeGen/Fold.hs view
@@ -0,0 +1,625 @@+{-# LANGUAGE GADTs #-}+{-# LANGUAGE OverloadedStrings #-}+{-# LANGUAGE RebindableSyntax #-}+{-# LANGUAGE RecordWildCards #-}+{-# LANGUAGE ScopedTypeVariables #-}+{-# LANGUAGE TypeOperators #-}+{-# LANGUAGE ViewPatterns #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.CodeGen.Fold+-- Copyright : [2016..2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.CodeGen.Fold+ where++-- accelerate+import Data.Array.Accelerate.Analysis.Match+import Data.Array.Accelerate.Analysis.Type+import Data.Array.Accelerate.Array.Sugar ( Array, Scalar, Vector, Shape, Z, (:.), Elt(..) )++-- accelerate-llvm-*+import Data.Array.Accelerate.LLVM.Analysis.Match+import Data.Array.Accelerate.LLVM.CodeGen.Arithmetic as A+import Data.Array.Accelerate.LLVM.CodeGen.Array+import Data.Array.Accelerate.LLVM.CodeGen.Base+import Data.Array.Accelerate.LLVM.CodeGen.Environment+import Data.Array.Accelerate.LLVM.CodeGen.Exp+import Data.Array.Accelerate.LLVM.CodeGen.IR+import Data.Array.Accelerate.LLVM.CodeGen.Loop as Loop+import Data.Array.Accelerate.LLVM.CodeGen.Monad+import Data.Array.Accelerate.LLVM.CodeGen.Sugar++import Data.Array.Accelerate.LLVM.PTX.Analysis.Launch+import Data.Array.Accelerate.LLVM.PTX.CodeGen.Base+import Data.Array.Accelerate.LLVM.PTX.CodeGen.Generate+import Data.Array.Accelerate.LLVM.PTX.Context+import Data.Array.Accelerate.LLVM.PTX.Target++import LLVM.AST.Type.Representation++-- cuda+import qualified Foreign.CUDA.Analysis as CUDA++import Control.Applicative ( (<$>), (<*>) )+import Control.Monad ( (>=>), (<=<) )+import Data.String ( fromString )+import Data.Bits as P+import Prelude as P+++-- Reduce an array along the innermost dimension. The reduction function must be+-- associative to allow for an efficient parallel implementation, but the+-- initial element does /not/ need to be a neutral element of operator.+--+-- TODO: Specialise for commutative operations (such as (+)) and those with+-- a neutral element {(+), 0}+--+mkFold+ :: forall aenv sh e. (Shape sh, Elt e)+ => PTX+ -> Gamma aenv+ -> IRFun2 PTX aenv (e -> e -> e)+ -> IRExp PTX aenv e+ -> IRDelayed PTX aenv (Array (sh :. Int) e)+ -> CodeGen (IROpenAcc PTX aenv (Array sh e))+mkFold ptx@(deviceProperties . ptxContext -> dev) aenv f z acc+ | Just Refl <- matchShapeType (undefined::sh) (undefined::Z)+ = (+++) <$> mkFoldAll dev aenv f (Just z) acc+ <*> mkFoldFill ptx aenv z++ | otherwise+ = (+++) <$> mkFoldDim dev aenv f (Just z) acc+ <*> mkFoldFill ptx aenv z+++-- Reduce a non-empty array along the innermost dimension. The reduction+-- function must be associative to allow for an efficient parallel+-- implementation.+--+-- TODO: Specialise for commutative operations (such as (+)) and those with+-- a neutral element {(+), 0}+--+mkFold1+ :: forall aenv sh e. (Shape sh, Elt e)+ => PTX+ -> Gamma aenv+ -> IRFun2 PTX aenv (e -> e -> e)+ -> IRDelayed PTX aenv (Array (sh :. Int) e)+ -> CodeGen (IROpenAcc PTX aenv (Array sh e))+mkFold1 (deviceProperties . ptxContext -> dev) aenv f acc+ | Just Refl <- matchShapeType (undefined::sh) (undefined::Z)+ = mkFoldAll dev aenv f Nothing acc++ | otherwise+ = mkFoldDim dev aenv f Nothing acc+++-- Reduce an array to a single element.+--+-- Since reductions consume arrays that have been fused into them, parallel+-- reduction requires two separate kernels. At an example, take vector dot+-- product:+--+-- > dotp xs ys = fold (+) 0 (zipWith (*) xs ys)+--+-- 1. The first pass reads in the fused array data, in this case corresponding+-- to the function (\i -> (xs!i) * (ys!i)).+--+-- 2. The second pass reads in the manifest array data from the first step and+-- directly reduces the array. This can be done recursively in-place until+-- only a single element remains.+--+-- In both phases, thread blocks cooperatively reduce a stripe of the input (one+-- element per thread) to a single element, which is stored to the output array.+--+mkFoldAll+ :: forall aenv e. Elt e+ => DeviceProperties -- ^ properties of the target GPU+ -> Gamma aenv -- ^ array environment+ -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function+ -> Maybe (IRExp PTX aenv e) -- ^ seed element, if this is an exclusive reduction+ -> IRDelayed PTX aenv (Vector e) -- ^ input data+ -> CodeGen (IROpenAcc PTX aenv (Scalar e))+mkFoldAll dev aenv combine mseed acc =+ foldr1 (+++) <$> sequence [ mkFoldAllS dev aenv combine mseed acc+ , mkFoldAllM1 dev aenv combine acc+ , mkFoldAllM2 dev aenv combine mseed+ ]+++-- Reduction to an array to a single element, for small arrays which can be+-- processed by a single thread block.+--+mkFoldAllS+ :: forall aenv e. Elt e+ => DeviceProperties -- ^ properties of the target GPU+ -> Gamma aenv -- ^ array environment+ -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function+ -> Maybe (IRExp PTX aenv e)+ -> IRDelayed PTX aenv (Vector e) -- ^ input data+ -> CodeGen (IROpenAcc PTX aenv (Scalar e))+mkFoldAllS dev aenv combine mseed IRDelayed{..} =+ let+ (start, end, paramGang) = gangParam+ (arrOut, paramOut) = mutableArray ("out" :: Name (Scalar e))+ paramEnv = envParam aenv+ --+ config = launchConfig dev (CUDA.incWarp dev) smem multipleOf+ smem n = warps * (1 + per_warp) * bytes+ where+ ws = CUDA.warpSize dev+ warps = n `div` ws+ per_warp = ws + ws `div` 2+ bytes = sizeOf (eltType (undefined :: e))+ in+ makeOpenAccWith config "foldAllS" (paramGang ++ paramOut ++ paramEnv) $ do++ tid <- threadIdx+ bd <- blockDim++ -- We can assume that there is only a single thread block+ i0 <- A.add numType start tid+ sz <- A.sub numType end start+ when (A.lt scalarType i0 sz) $ do++ -- Thread reads initial element and then participates in block-wide+ -- reduction.+ x0 <- app1 delayedLinearIndex =<< A.fromIntegral integralType numType i0+ r0 <- if A.eq scalarType sz bd+ then reduceBlockSMem dev combine Nothing x0+ else reduceBlockSMem dev combine (Just sz) x0++ when (A.eq scalarType tid (lift 0)) $+ writeArray arrOut tid =<<+ case mseed of+ Nothing -> return r0+ Just z -> flip (app2 combine) r0 =<< z -- Note: initial element on the left++ return_+++-- Reduction of an entire array to a single element. This kernel implements step+-- one for reducing large arrays which must be processed by multiple thread+-- blocks.+--+mkFoldAllM1+ :: forall aenv e. Elt e+ => DeviceProperties -- ^ properties of the target GPU+ -> Gamma aenv -- ^ array environment+ -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function+ -> IRDelayed PTX aenv (Vector e) -- ^ input data+ -> CodeGen (IROpenAcc PTX aenv (Scalar e))+mkFoldAllM1 dev aenv combine IRDelayed{..} =+ let+ (start, end, paramGang) = gangParam+ (arrTmp, paramTmp) = mutableArray ("tmp" :: Name (Vector e))+ paramEnv = envParam aenv+ --+ config = launchConfig dev (CUDA.incWarp dev) smem const+ smem n = warps * (1 + per_warp) * bytes+ where+ ws = CUDA.warpSize dev+ warps = n `div` ws+ per_warp = ws + ws `div` 2+ bytes = sizeOf (eltType (undefined :: e))+ in+ makeOpenAccWith config "foldAllM1" (paramGang ++ paramTmp ++ paramEnv) $ do++ -- Each thread block cooperatively reduces a stripe of the input and stores+ -- that value into a temporary array at a corresponding index. Since the+ -- order of operations remains fixed, this method supports non-commutative+ -- reductions.+ --+ tid <- threadIdx+ bd <- blockDim+ sz <- i32 . indexHead =<< delayedExtent++ imapFromTo start end $ \seg -> do++ -- Wait for all threads to catch up before beginning the stripe+ __syncthreads++ -- Bounds of the input array we will reduce between+ from <- A.mul numType seg bd+ step <- A.add numType from bd+ to <- A.min scalarType sz step++ -- Threads cooperatively reduce this stripe+ reduceFromTo dev from to combine+ (app1 delayedLinearIndex <=< A.fromIntegral integralType numType)+ (when (A.eq scalarType tid (lift 0)) . writeArray arrTmp seg)++ return_+++-- Reduction of an array to a single element, (recursive) step 2 of multi-block+-- reduction algorithm.+--+mkFoldAllM2+ :: forall aenv e. Elt e+ => DeviceProperties+ -> Gamma aenv+ -> IRFun2 PTX aenv (e -> e -> e)+ -> Maybe (IRExp PTX aenv e)+ -> CodeGen (IROpenAcc PTX aenv (Scalar e))+mkFoldAllM2 dev aenv combine mseed =+ let+ (start, end, paramGang) = gangParam+ (arrTmp, paramTmp) = mutableArray ("tmp" :: Name (Vector e))+ (arrOut, paramOut) = mutableArray ("out" :: Name (Vector e))+ paramEnv = envParam aenv+ --+ config = launchConfig dev (CUDA.incWarp dev) smem const+ smem n = warps * (1 + per_warp) * bytes+ where+ ws = CUDA.warpSize dev+ warps = n `div` ws+ per_warp = ws + ws `div` 2+ bytes = sizeOf (eltType (undefined :: e))+ in+ makeOpenAccWith config "foldAllM2" (paramGang ++ paramTmp ++ paramOut ++ paramEnv) $ do++ -- Threads cooperatively reduce a stripe of the input (temporary) array+ -- output from the first phase, storing the results into another temporary.+ -- When only a single thread block remains, we have reached the final+ -- reduction step and add the initial element (for exclusive reductions).+ --+ tid <- threadIdx+ bd <- blockDim+ gd <- gridDim+ sz <- i32 . indexHead $ irArrayShape arrTmp++ imapFromTo start end $ \seg -> do++ -- Wait for all threads to catch up before beginning the stripe+ __syncthreads++ -- Bounds of the input we will reduce between+ from <- A.mul numType seg bd+ step <- A.add numType from bd+ to <- A.min scalarType sz step++ -- Threads cooperatively reduce this stripe+ reduceFromTo dev from to combine (readArray arrTmp) $ \r ->+ when (A.eq scalarType tid (lift 0)) $+ writeArray arrOut seg =<<+ case mseed of+ Nothing -> return r+ Just z -> if A.eq scalarType gd (lift 1)+ then flip (app2 combine) r =<< z -- Note: initial element on the left+ else return r++ return_+++-- Reduce an array of arbitrary rank along the innermost dimension only.+--+-- For simplicity, each element of the output (reduction along an+-- innermost-dimension index) is computed by a single thread block, meaning we+-- don't have to worry about inter-block synchronisation. A more balanced method+-- would be a segmented reduction (specialised, since the length of each segment+-- is known a priori).+--+mkFoldDim+ :: forall aenv sh e. (Shape sh, Elt e)+ => DeviceProperties -- ^ properties of the target GPU+ -> Gamma aenv -- ^ array environment+ -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function+ -> Maybe (IRExp PTX aenv e) -- ^ seed element, if this is an exclusive reduction+ -> IRDelayed PTX aenv (Array (sh :. Int) e) -- ^ input data+ -> CodeGen (IROpenAcc PTX aenv (Array sh e))+mkFoldDim dev aenv combine mseed IRDelayed{..} =+ let+ (start, end, paramGang) = gangParam+ (arrOut, paramOut) = mutableArray ("out" :: Name (Array sh e))+ paramEnv = envParam aenv+ --+ config = launchConfig dev (CUDA.incWarp dev) smem const+ smem n = warps * (1 + per_warp) * bytes+ where+ ws = CUDA.warpSize dev+ warps = n `div` ws+ per_warp = ws + ws `div` 2+ bytes = sizeOf (eltType (undefined :: e))+ in+ makeOpenAccWith config "fold" (paramGang ++ paramOut ++ paramEnv) $ do++ -- If the innermost dimension is smaller than the number of threads in the+ -- block, those threads will never contribute to the output.+ tid <- threadIdx+ sz <- i32 . indexHead =<< delayedExtent+ when (A.lt scalarType tid sz) $ do++ -- Thread blocks iterate over the outer dimensions, each thread block+ -- cooperatively reducing along each outermost index to a single value.+ --+ imapFromTo start end $ \seg -> do++ -- Wait for threads to catch up before starting this segment. We could+ -- also place this at the bottom of the loop, but here allows threads to+ -- exit quickly on the last iteration.+ __syncthreads++ -- Step 1: initialise local sums+ from <- A.mul numType seg sz -- first linear index this block will reduce+ to <- A.add numType from sz -- last linear index this block will reduce (exclusive)++ i0 <- A.add numType from tid+ x0 <- app1 delayedLinearIndex =<< A.fromIntegral integralType numType i0+ bd <- blockDim+ r0 <- if A.gte scalarType sz bd+ then reduceBlockSMem dev combine Nothing x0+ else reduceBlockSMem dev combine (Just sz) x0++ -- Step 2: keep walking over the input+ next <- A.add numType from bd+ r <- iterFromStepTo next bd to r0 $ \offset r -> do++ -- Wait for all threads to catch up before starting the next stripe+ __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++ -- 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++ if A.eq scalarType tid (lift 0)+ then app2 combine r r'+ else return r'++ -- Step 3: Thread 0 writes the aggregate reduction of this dimension to+ -- memory. If this is an exclusive fold, combine with the initial element.+ --+ when (A.eq scalarType tid (lift 0)) $+ writeArray arrOut seg =<<+ case mseed of+ Nothing -> return r+ Just z -> flip (app2 combine) r =<< z -- Note: initial element on the left++ return_+++-- Exclusive reductions over empty arrays (of any dimension) fill the lower+-- dimensions with the initial element.+--+mkFoldFill+ :: (Shape sh, Elt e)+ => PTX+ -> Gamma aenv+ -> IRExp PTX aenv e+ -> CodeGen (IROpenAcc PTX aenv (Array sh e))+mkFoldFill ptx aenv seed =+ mkGenerate ptx aenv (IRFun1 (const seed))+++-- Efficient threadblock-wide reduction using the specified operator. The+-- aggregate reduction value is stored in thread zero. Supports non-commutative+-- operators.+--+-- Requires dynamically allocated memory: (#warps * (1 + 1.5 * warp size)).+--+-- Example: https://github.com/NVlabs/cub/blob/1.5.2/cub/block/specializations/block_reduce_warp_reductions.cuh+--+reduceBlockSMem+ :: forall aenv e. Elt e+ => DeviceProperties -- ^ properties of the target device+ -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function+ -> Maybe (IR Int32) -- ^ number of valid elements (may be less than block size)+ -> IR e -- ^ calling thread's input element+ -> CodeGen (IR e) -- ^ thread-block-wide reduction using the specified operator (lane 0 only)+reduceBlockSMem dev combine size = warpReduce >=> warpAggregate+ where+ int32 :: Integral a => a -> IR Int32+ int32 = lift . P.fromIntegral++ -- Temporary storage required for each warp+ bytes = sizeOf (eltType (undefined::e))+ warp_smem_elems = CUDA.warpSize dev + (CUDA.warpSize dev `div` 2)++ -- Step 1: Reduction in every warp+ --+ warpReduce :: IR e -> CodeGen (IR e)+ warpReduce input = do+ -- Allocate (1.5 * warpSize) elements of shared memory for each warp+ wid <- warpId+ skip <- A.mul numType wid (int32 (warp_smem_elems * bytes))+ smem <- dynamicSharedMem (int32 warp_smem_elems) skip++ -- Are we doing bounds checking for this warp?+ --+ case size of+ -- The entire thread block is valid, so skip bounds checks.+ Nothing ->+ reduceWarpSMem dev combine smem Nothing input++ -- Otherwise check how many elements are valid for this warp. If it is+ -- full then we can still skip bounds checks for it.+ Just n -> do+ offset <- A.mul numType wid (int32 (CUDA.warpSize dev))+ valid <- A.sub numType n offset+ if A.gte scalarType valid (int32 (CUDA.warpSize dev))+ then reduceWarpSMem dev combine smem Nothing input+ else reduceWarpSMem dev combine smem (Just valid) input++ -- Step 2: Aggregate per-warp reductions+ --+ warpAggregate :: IR e -> CodeGen (IR e)+ warpAggregate input = do+ -- Allocate #warps elements of shared memory+ bd <- blockDim+ warps <- A.quot integralType bd (int32 (CUDA.warpSize dev))+ skip <- A.mul numType warps (int32 (warp_smem_elems * bytes))+ smem <- dynamicSharedMem warps skip++ -- Share the per-lane aggregates+ wid <- warpId+ lane <- laneId+ when (A.eq scalarType lane (lift 0)) $ do+ writeArray smem wid input++ -- Wait for each warp to finish its local reduction+ __syncthreads++ -- Update the total aggregate. Thread 0 just does this sequentially (as is+ -- done in CUB), but we could also do this cooperatively (better for+ -- larger thread blocks?)+ tid <- threadIdx+ if A.eq scalarType tid (lift 0)+ then do+ steps <- case size of+ Nothing -> return warps+ Just n -> do+ a <- A.add numType n (int32 (CUDA.warpSize dev - 1))+ b <- A.quot integralType a (int32 (CUDA.warpSize dev))+ return b+ iterFromStepTo (lift 1) (lift 1) steps input $ \step x ->+ app2 combine x =<< readArray smem step+ else+ return input+++-- Efficient warp-wide reduction using shared memory. The aggregate reduction+-- value for the warp is stored in thread lane zero.+--+-- Each warp requires 48 (1.5 x warp size) elements of shared memory. The+-- routine assumes that is is allocated individually per-warp (i.e. can be+-- indexed in the range [0,warp size)).+--+-- Example: https://github.com/NVlabs/cub/blob/1.5.2/cub/warp/specializations/warp_reduce_smem.cuh#L128+--+reduceWarpSMem+ :: forall aenv e. Elt e+ => DeviceProperties -- ^ properties of the target device+ -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function+ -> IRArray (Vector e) -- ^ temporary storage array in shared memory (1.5 warp size elements)+ -> Maybe (IR Int32) -- ^ number of items that will be reduced by this warp, otherwise all lanes are valid+ -> IR e -- ^ calling thread's input element+ -> CodeGen (IR e) -- ^ warp-wide reduction using the specified operator (lane 0 only)+reduceWarpSMem dev combine smem size = reduce 0+ where+ log2 :: Double -> Double+ log2 = P.logBase 2++ -- Number steps required to reduce warp+ steps = P.floor . log2 . P.fromIntegral . CUDA.warpSize $ dev++ -- Return whether the index is valid. Assume that constant branches are+ -- optimised away.+ valid i =+ case size of+ Nothing -> return (lift True)+ Just n -> A.lt scalarType i n++ -- Unfold the reduction as a recursive code generation function.+ reduce :: Int -> IR e -> CodeGen (IR e)+ reduce step x+ | step >= steps = return x+ | offset <- 1 `P.shiftL` step = do+ -- share input through buffer+ lane <- laneId+ writeArray smem lane x++ -- update input if in range+ i <- A.add numType lane (lift offset)+ x' <- if valid i+ then app2 combine x =<< readArray smem i+ else return x++ reduce (step+1) x'+++-- Efficient warp reduction using __shfl_up instruction (compute >= 3.0)+--+-- Example: https://github.com/NVlabs/cub/blob/1.5.2/cub/warp/specializations/warp_reduce_shfl.cuh#L310+--+-- reduceWarpShfl+-- :: IRFun2 PTX aenv (e -> e -> e) -- ^ combination function+-- -> IR e -- ^ this thread's input value+-- -> CodeGen (IR e) -- ^ final result+-- reduceWarpShfl combine input =+-- error "TODO: PTX.reduceWarpShfl"+++-- Reduction loops+-- ---------------++reduceFromTo+ :: Elt a+ => DeviceProperties+ -> IR Int32 -- ^ starting index+ -> IR Int32 -- ^ final index (exclusive)+ -> (IRFun2 PTX aenv (a -> a -> a)) -- ^ combination function+ -> (IR Int32 -> CodeGen (IR a)) -- ^ function to retrieve element at index+ -> (IR a -> CodeGen ()) -- ^ what to do with the value+ -> CodeGen ()+reduceFromTo dev from to combine get set = do++ tid <- threadIdx+ bd <- blockDim++ valid <- A.sub numType to from+ i <- A.add numType from tid++ _ <- if A.gte scalarType valid bd+ then do+ -- All threads in the block will participate in the reduction, so+ -- we can avoid bounds checks+ x <- get i+ r <- reduceBlockSMem dev combine Nothing x+ set r++ return (IR OP_Unit :: IR ()) -- unsightly, but free+ else do+ -- Only in-bounds threads can read their input and participate in+ -- the reduction+ when (A.lt scalarType i to) $ do+ x <- get i+ r <- reduceBlockSMem dev combine (Just valid) x+ set r++ return (IR OP_Unit :: IR ())++ return ()++++-- Utilities+-- ---------++i32 :: IR Int -> CodeGen (IR Int32)+i32 = A.fromIntegral integralType numType+++imapFromTo+ :: IR Int32+ -> IR Int32+ -> (IR Int32 -> CodeGen ())+ -> CodeGen ()+imapFromTo start end body = do+ bid <- blockIdx+ gd <- gridDim+ i0 <- A.add numType start bid+ imapFromStepTo i0 gd end body+
+ Data/Array/Accelerate/LLVM/PTX/CodeGen/FoldSeg.hs view
@@ -0,0 +1,461 @@+{-# LANGUAGE GADTs #-}+{-# LANGUAGE OverloadedStrings #-}+{-# LANGUAGE RebindableSyntax #-}+{-# LANGUAGE RecordWildCards #-}+{-# LANGUAGE ScopedTypeVariables #-}+{-# LANGUAGE TypeOperators #-}+{-# LANGUAGE ViewPatterns #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.CodeGen.FoldSeg+-- Copyright : [2016..2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.CodeGen.FoldSeg+ where++-- accelerate+import Data.Array.Accelerate.Analysis.Type+import Data.Array.Accelerate.Array.Sugar ( Array, Segments, Shape(rank), (:.), Elt(..) )++-- accelerate-llvm-*+import LLVM.AST.Type.Representation++import Data.Array.Accelerate.LLVM.CodeGen.Arithmetic as A+import Data.Array.Accelerate.LLVM.CodeGen.Array+import Data.Array.Accelerate.LLVM.CodeGen.Base+import Data.Array.Accelerate.LLVM.CodeGen.Constant+import Data.Array.Accelerate.LLVM.CodeGen.Environment+import Data.Array.Accelerate.LLVM.CodeGen.Exp+import Data.Array.Accelerate.LLVM.CodeGen.IR+import Data.Array.Accelerate.LLVM.CodeGen.Loop as Loop+import Data.Array.Accelerate.LLVM.CodeGen.Monad+import Data.Array.Accelerate.LLVM.CodeGen.Sugar++import Data.Array.Accelerate.LLVM.PTX.Analysis.Launch+import Data.Array.Accelerate.LLVM.PTX.CodeGen.Base+import Data.Array.Accelerate.LLVM.PTX.CodeGen.Fold ( reduceBlockSMem, reduceWarpSMem, imapFromTo )+-- import Data.Array.Accelerate.LLVM.PTX.CodeGen.Queue+import Data.Array.Accelerate.LLVM.PTX.Context+import Data.Array.Accelerate.LLVM.PTX.Target++-- cuda+import qualified Foreign.CUDA.Analysis as CUDA++import Control.Applicative ( (<$>), (<*>) )+import Control.Monad ( void )+import Data.String ( fromString )+import Prelude as P+++-- Segmented reduction along the innermost dimension of an array. Performs one+-- reduction per segment of the source array.+--+mkFoldSeg+ :: forall aenv sh i e. (Shape sh, IsIntegral i, Elt i, Elt e)+ => PTX+ -> Gamma aenv+ -> IRFun2 PTX aenv (e -> e -> e)+ -> IRExp PTX aenv e+ -> IRDelayed PTX aenv (Array (sh :. Int) e)+ -> IRDelayed PTX aenv (Segments i)+ -> CodeGen (IROpenAcc PTX aenv (Array (sh :. Int) e))+mkFoldSeg (deviceProperties . ptxContext -> dev) aenv combine seed arr seg =+ (+++) <$> mkFoldSegP_block dev aenv combine (Just seed) arr seg+ <*> mkFoldSegP_warp dev aenv combine (Just seed) arr seg+++-- Segmented reduction along the innermost dimension of an array, where /all/+-- segments are non-empty.+--+mkFold1Seg+ :: forall aenv sh i e. (Shape sh, IsIntegral i, Elt i, Elt e)+ => PTX+ -> Gamma aenv+ -> IRFun2 PTX aenv (e -> e -> e)+ -> IRDelayed PTX aenv (Array (sh :. Int) e)+ -> IRDelayed PTX aenv (Segments i)+ -> CodeGen (IROpenAcc PTX aenv (Array (sh :. Int) e))+mkFold1Seg (deviceProperties . ptxContext -> dev) aenv combine arr seg =+ (+++) <$> mkFoldSegP_block dev aenv combine Nothing arr seg+ <*> mkFoldSegP_warp dev aenv combine Nothing arr seg+++-- This implementation assumes that the segments array represents the offset+-- indices to the source array, rather than the lengths of each segment. The+-- segment-offset approach is required for parallel implementations.+--+-- Each segment is computed by a single thread block, meaning we don't have to+-- worry about inter-block synchronisation.+--+mkFoldSegP_block+ :: forall aenv sh i e. (Shape sh, IsIntegral i, Elt i, Elt e)+ => DeviceProperties+ -> Gamma aenv+ -> IRFun2 PTX aenv (e -> e -> e)+ -> Maybe (IRExp PTX aenv e)+ -> IRDelayed PTX aenv (Array (sh :. Int) e)+ -> IRDelayed PTX aenv (Segments i)+ -> CodeGen (IROpenAcc PTX aenv (Array (sh :. Int) e))+mkFoldSegP_block dev aenv combine mseed arr seg =+ let+ (start, end, paramGang) = gangParam+ (arrOut, paramOut) = mutableArray ("out" :: Name (Array (sh :. Int) e))+ paramEnv = envParam aenv+ --+ config = launchConfig dev (CUDA.decWarp dev) dsmem const+ dsmem n = warps * (1 + per_warp) * bytes+ where+ ws = CUDA.warpSize dev+ warps = n `div` ws+ per_warp = ws + ws `div` 2+ bytes = sizeOf (eltType (undefined :: e))+ in+ makeOpenAccWith config "foldSeg_block" (paramGang ++ paramOut ++ paramEnv) $ do++ -- We use a dynamically scheduled work queue in order to evenly distribute+ -- the uneven workload, due to the variable length of each segment, over the+ -- available thread blocks.+ -- queue <- globalWorkQueue++ -- All threads in the block need to know what the start and end indices of+ -- this segment are in order to participate in the reduction. We use+ -- variables in __shared__ memory to communicate these values between+ -- threads in the block. Furthermore, by using a 2-element array, we can+ -- have the first two threads of the block read the start and end indices as+ -- a single coalesced read, since they will be sequential in the+ -- segment-offset array.+ --+ smem <- staticSharedMem 2++ -- Compute the number of segments and size of the innermost dimension. These+ -- are required if we are reducing a rank-2 or higher array, to properly+ -- compute the start and end indices of the portion of the array this thread+ -- block reduces. Note that this is a segment-offset array computed by+ -- 'scanl (+) 0' of the segment length array, so its size has increased by+ -- one.+ --+ sz <- i32 . indexHead =<< delayedExtent arr+ ss <- do n <- i32 . indexHead =<< delayedExtent seg+ A.sub numType n (lift 1)++ -- Each thread block cooperatively reduces a segment.+ -- s0 <- dequeue queue (lift 1)+ -- for s0 (\s -> A.lt scalarType s end) (\_ -> dequeue queue (lift 1)) $ \s -> do+ imapFromTo start end $ \s -> do++ -- The first two threads of the block determine the indices of the+ -- segments array that we will reduce between and distribute those values+ -- to the other threads in the block.+ tid <- threadIdx+ when (A.lt scalarType tid (lift 2)) $ do+ i <- case rank (undefined::sh) of+ 0 -> return s+ _ -> A.rem integralType s ss+ j <- A.add numType i tid+ v <- app1 (delayedLinearIndex seg) =<< A.fromIntegral integralType numType j+ writeArray smem tid =<< i32 v++ -- Once all threads have caught up, begin work on the new segment.+ __syncthreads++ u <- readArray smem (lift 0 :: IR Int32)+ v <- readArray smem (lift 1 :: IR Int32)++ -- Determine the index range of the input array we will reduce over.+ -- Necessary for multidimensional segmented reduction.+ (inf,sup) <- A.unpair <$> case rank (undefined::sh) of+ 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++ void $+ if A.eq scalarType inf sup+ -- This segment is empty. If this is an exclusive reduction the+ -- first thread writes out the initial element for this segment.+ then do+ case mseed of+ Nothing -> return (IR OP_Unit :: IR ())+ Just z -> do+ when (A.eq scalarType tid (lift 0)) $ writeArray arrOut s =<< z+ return (IR OP_Unit)++ -- This is a non-empty segment.+ else do+ -- Step 1: initialise local sums+ --+ -- NOTE: We require all threads to enter this branch and execute the+ -- first step, even if they do not have a valid element and must+ -- return 'undef'. If we attempt to skip this entire section for+ -- non-participating threads (i.e. 'when (i0 < sup)'), it seems that+ -- 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.+ --+ bd <- blockDim+ i0 <- A.add numType inf tid+ x0 <- if A.lt scalarType i0 sup+ then app1 (delayedLinearIndex arr) =<< A.fromIntegral integralType numType i0+ else let+ go :: TupleType a -> Operands a+ go UnitTuple = OP_Unit+ go (PairTuple a b) = OP_Pair (go a) (go b)+ go (SingleTuple t) = ir' t (undef t)+ in+ return . IR $ go (eltType (undefined::e))++ v0 <- A.sub numType sup inf+ r0 <- if A.gte scalarType v0 bd+ then reduceBlockSMem dev combine Nothing x0+ else reduceBlockSMem dev combine (Just v0) x0++ -- Step 2: keep walking over the input+ nxt <- A.add numType inf bd+ r <- iterFromStepTo nxt bd sup r0 $ \offset r -> do++ -- Wait for threads to catch up before starting the next stripe+ __syncthreads++ i' <- A.add numType offset tid+ v' <- A.sub numType sup offset+ r' <- if A.gte scalarType v' bd+ -- 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'++ -- 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++ -- first thread incorporates the result from the previous+ -- iteration+ if A.eq scalarType tid (lift 0)+ then app2 combine r r'+ else return r'++ -- Step 3: Thread zero writes the aggregate reduction for this+ -- segment to memory. If this is an exclusive fold combine with the+ -- initial element as well.+ when (A.eq scalarType tid (lift 0)) $+ writeArray arrOut s =<<+ case mseed of+ Nothing -> return r+ Just z -> flip (app2 combine) r =<< z -- Note: initial element on the left++ return (IR OP_Unit)++ return_+++-- This implementation assumes that the segments array represents the offset+-- indices to the source array, rather than the lengths of each segment. The+-- segment-offset approach is required for parallel implementations.+--+-- Each segment is computed by a single warp, meaning we don't have to worry+-- about inter- or intra-block synchronisation.+--+mkFoldSegP_warp+ :: forall aenv sh i e. (Shape sh, IsIntegral i, Elt i, Elt e)+ => DeviceProperties+ -> Gamma aenv+ -> IRFun2 PTX aenv (e -> e -> e)+ -> Maybe (IRExp PTX aenv e)+ -> IRDelayed PTX aenv (Array (sh :. Int) e)+ -> IRDelayed PTX aenv (Segments i)+ -> CodeGen (IROpenAcc PTX aenv (Array (sh :. Int) e))+mkFoldSegP_warp dev aenv combine mseed arr seg =+ let+ (start, end, paramGang) = gangParam+ (arrOut, paramOut) = mutableArray ("out" :: Name (Array (sh :. Int) e))+ paramEnv = envParam aenv+ --+ config = launchConfig dev (CUDA.decWarp dev) dsmem grid+ dsmem n = warps * (2 + per_warp_elems) * bytes+ where+ warps = n `div` ws+ --+ grid n m = multipleOf n (m `div` ws)+ --+ per_warp_bytes = per_warp_elems * bytes+ per_warp_elems = ws + (ws `div` 2)+ ws = CUDA.warpSize dev+ bytes = sizeOf (eltType (undefined :: e))++ int32 :: Integral a => a -> IR Int32+ int32 = lift . P.fromIntegral+ in+ makeOpenAccWith config "foldSeg_warp" (paramGang ++ paramOut ++ paramEnv) $ do++ -- Each warp works independently.+ -- Determine the ID of this warp within the thread block.+ tid <- threadIdx+ wid <- A.quot integralType tid (int32 ws)++ -- Number of warps per thread block+ bd <- blockDim+ wpb <- A.quot integralType bd (int32 ws)++ -- ID of this warp within the grid+ bid <- blockIdx+ gwid <- do a <- A.mul numType bid wpb+ b <- A.add numType wid a+ return b++ -- All threads in the warp need to know what the start and end indices of+ -- this segment are in order to participate in the reduction. We use+ -- variables in __shared__ memory to communicate these values between+ -- threads. Furthermore, by using a 2-element array, we can have the first+ -- two threads of the warp read the start and end indices as a single+ -- coalesced read, as these elements will be adjacent in the segment-offset+ -- array.+ --+ lim <- do+ a <- A.mul numType wid (int32 (2 * bytes))+ b <- dynamicSharedMem (lift 2) a+ return b++ -- Allocate (1.5 * warpSize) elements of share memory for each warp+ smem <- do+ a <- A.mul numType wpb (int32 (2 * bytes))+ b <- A.mul numType wid (int32 per_warp_bytes)+ c <- A.add numType a b+ d <- dynamicSharedMem (int32 per_warp_elems) c+ return d++ -- Compute the number of segments and size of the innermost dimension. These+ -- are required if we are reducing a rank-2 or higher array, to properly+ -- compute the start and end indices of the portion of the array this warp+ -- reduces. Note that this is a segment-offset array computed by 'scanl (+) 0'+ -- of the segment length array, so its size has increased by one.+ --+ sz <- i32 . indexHead =<< delayedExtent arr+ ss <- do a <- i32 . indexHead =<< delayedExtent seg+ b <- A.sub numType a (lift 1)+ return b++ -- Each thread reduces a segment independently+ s0 <- A.add numType start gwid+ gd <- gridDim+ step <- A.mul numType wpb gd+ imapFromStepTo s0 step end $ \s -> do++ -- The first two threads of the warp determine the indices of the segments+ -- array that we will reduce between and distribute those values to the+ -- other threads in the warp+ lane <- laneId+ when (A.lt scalarType lane (lift 2)) $ do+ a <- case rank (undefined::sh) of+ 0 -> return s+ _ -> A.rem integralType s ss+ b <- A.add numType a lane+ c <- app1 (delayedLinearIndex seg) =<< A.fromIntegral integralType numType b+ writeArray lim lane =<< i32 c++ -- Determine the index range of the input array we will reduce over.+ -- Necessary for multidimensional segmented reduction.+ (inf,sup) <- do+ u <- readArray lim (lift 0 :: IR Int32)+ v <- readArray lim (lift 1 :: IR Int32)+ A.unpair <$> case rank (undefined::sh) of+ 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++ -- TLM: I don't think this should be necessary...+ __syncthreads++ void $+ if A.eq scalarType inf sup+ -- This segment is empty. If this is an exclusive reduction the first+ -- lane writes out the initial element for this segment.+ then do+ case mseed of+ Nothing -> return (IR OP_Unit :: IR ())+ Just z -> do+ when (A.eq scalarType lane (lift 0)) $ writeArray arrOut s =<< z+ return (IR OP_Unit)++ -- This is a non-empty segment.+ else do+ -- Step 1: initialise local sums+ --+ -- See comment above why we initialise the loop in this way+ --+ i0 <- A.add numType inf lane+ x0 <- if A.lt scalarType i0 sup+ then app1 (delayedLinearIndex arr) =<< A.fromIntegral integralType numType i0+ else let+ go :: TupleType a -> Operands a+ go UnitTuple = OP_Unit+ go (PairTuple a b) = OP_Pair (go a) (go b)+ go (SingleTuple t) = ir' t (undef t)+ in+ return . IR $ go (eltType (undefined::e))++ v0 <- A.sub numType sup inf+ r0 <- if A.gte scalarType v0 (int32 ws)+ then reduceWarpSMem dev combine smem Nothing x0+ else reduceWarpSMem dev combine smem (Just v0) x0++ -- Step 2: Keep walking over the rest of the segment+ nx <- A.add numType inf (int32 ws)+ r <- iterFromStepTo nx (int32 ws) sup r0 $ \offset r -> do++ -- TLM: Similarly, I think this is unnecessary...+ __syncthreads++ i' <- A.add numType offset lane+ v' <- A.sub numType sup offset+ r' <- if A.gte scalarType v' (int32 ws)+ then do+ -- All lanes are in bounds, so avoid bounds checks+ x <- app1 (delayedLinearIndex arr) =<< A.fromIntegral integralType numType i'+ y <- reduceWarpSMem dev combine smem Nothing x+ 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++ -- The first lane incorporates the result from the previous+ -- iteration+ if A.eq scalarType lane (lift 0)+ then app2 combine r r'+ else return r'++ -- Step 3: Lane zero writes the aggregate reduction for this+ -- segment to memory. If this is an exclusive reduction, also+ -- combine with the initial element+ when (A.eq scalarType lane (lift 0)) $+ writeArray arrOut s =<<+ case mseed of+ Nothing -> return r+ Just z -> flip (app2 combine) r =<< z -- Note: initial element on the left++ return (IR OP_Unit)++ return_+++i32 :: IsIntegral a => IR a -> CodeGen (IR Int32)+i32 = A.fromIntegral integralType numType+
+ Data/Array/Accelerate/LLVM/PTX/CodeGen/Generate.hs view
@@ -0,0 +1,60 @@+{-# LANGUAGE OverloadedStrings #-}+{-# LANGUAGE ScopedTypeVariables #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.CodeGen.Generate+-- Copyright : [2014..2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.CodeGen.Generate+ where++import Prelude hiding ( fromIntegral )++-- accelerate+import Data.Array.Accelerate.Array.Sugar ( Array, Shape, Elt )+import Data.Array.Accelerate.Type++import Data.Array.Accelerate.LLVM.CodeGen.Arithmetic+import Data.Array.Accelerate.LLVM.CodeGen.Array+import Data.Array.Accelerate.LLVM.CodeGen.Base+import Data.Array.Accelerate.LLVM.CodeGen.Environment+import Data.Array.Accelerate.LLVM.CodeGen.Exp+import Data.Array.Accelerate.LLVM.CodeGen.Monad+import Data.Array.Accelerate.LLVM.CodeGen.Sugar++import Data.Array.Accelerate.LLVM.PTX.CodeGen.Base+import Data.Array.Accelerate.LLVM.PTX.CodeGen.Loop+import Data.Array.Accelerate.LLVM.PTX.Target ( PTX )+++-- Construct a new array by applying a function to each index. Each thread+-- processes multiple adjacent elements.+--+mkGenerate+ :: forall aenv sh e. (Shape sh, Elt e)+ => PTX+ -> Gamma aenv+ -> IRFun1 PTX aenv (sh -> e)+ -> CodeGen (IROpenAcc PTX aenv (Array sh e))+mkGenerate ptx aenv apply =+ let+ (start, end, paramGang) = gangParam+ (arrOut, paramOut) = mutableArray ("out" :: Name (Array sh e))+ paramEnv = envParam aenv+ in+ makeOpenAcc ptx "generate" (paramGang ++ paramOut ++ paramEnv) $ do++ imapFromTo start end $ \i -> do+ i' <- fromIntegral integralType numType i -- loop counter is Int32+ ix <- indexOfInt (irArrayShape arrOut) i' -- convert to multidimensional index+ r <- app1 apply ix -- apply generator function+ writeArray arrOut i' r -- store result++ return_++
+ Data/Array/Accelerate/LLVM/PTX/CodeGen/Loop.hs view
@@ -0,0 +1,47 @@+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.CodeGen.Loop+-- Copyright : [2015..2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.CodeGen.Loop+ where++-- accelerate+import Data.Array.Accelerate.Type++import Data.Array.Accelerate.LLVM.CodeGen.Arithmetic+import Data.Array.Accelerate.LLVM.CodeGen.IR+import Data.Array.Accelerate.LLVM.CodeGen.Monad+import qualified Data.Array.Accelerate.LLVM.CodeGen.Loop as Loop++import Data.Array.Accelerate.LLVM.PTX.CodeGen.Base+++-- | A standard loop where the CUDA threads cooperatively step over an index+-- space from the start to end indices. The threads stride the array in a way+-- that maintains memory coalescing.+--+-- The start and end array indices are given as natural array indexes, and the+-- thread specific indices are calculated by the loop.+--+-- > for ( int32 i = blockDim.x * blockIdx.x + threadIdx.x + start+-- > ; i < end+-- > ; i += blockDim.x * gridDim.x )+--+-- TODO: This assumes that the starting offset retains alignment to the warp+-- boundary. This might not always be the case, so provide a version that+-- explicitly aligns reads to the warp boundary.+--+imapFromTo :: IR Int32 -> IR Int32 -> (IR Int32 -> CodeGen ()) -> CodeGen ()+imapFromTo start end body = do+ step <- gridSize+ tid <- globalThreadIdx+ i0 <- add numType tid start+ --+ Loop.imapFromStepTo i0 step end body+
+ Data/Array/Accelerate/LLVM/PTX/CodeGen/Map.hs view
@@ -0,0 +1,60 @@+{-# LANGUAGE GADTs #-}+{-# LANGUAGE OverloadedStrings #-}+{-# LANGUAGE RecordWildCards #-}+{-# LANGUAGE ScopedTypeVariables #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.CodeGen.Map+-- Copyright : [2014..2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.CodeGen.Map+ where++import Prelude hiding ( fromIntegral )++-- accelerate+import Data.Array.Accelerate.Array.Sugar ( Array, Elt )+import Data.Array.Accelerate.Type++import Data.Array.Accelerate.LLVM.CodeGen.Arithmetic+import Data.Array.Accelerate.LLVM.CodeGen.Array+import Data.Array.Accelerate.LLVM.CodeGen.Base+import Data.Array.Accelerate.LLVM.CodeGen.Environment+import Data.Array.Accelerate.LLVM.CodeGen.Monad+import Data.Array.Accelerate.LLVM.CodeGen.Sugar++import Data.Array.Accelerate.LLVM.PTX.CodeGen.Base+import Data.Array.Accelerate.LLVM.PTX.CodeGen.Loop+import Data.Array.Accelerate.LLVM.PTX.Target ( PTX )+++-- Apply a unary function to each element of an array. Each thread processes+-- multiple elements, striding the array by the grid size.+--+mkMap :: forall aenv sh a b. Elt b+ => PTX+ -> Gamma aenv+ -> IRFun1 PTX aenv (a -> b)+ -> IRDelayed PTX aenv (Array sh a)+ -> CodeGen (IROpenAcc PTX aenv (Array sh b))+mkMap ptx aenv apply IRDelayed{..} =+ let+ (start, end, paramGang) = gangParam+ (arrOut, paramOut) = mutableArray ("out" :: Name (Array sh b))+ paramEnv = envParam aenv+ in+ makeOpenAcc ptx "map" (paramGang ++ paramOut ++ paramEnv) $ do++ imapFromTo start end $ \i -> do+ i' <- fromIntegral integralType numType i+ xs <- app1 delayedLinearIndex i'+ ys <- app1 apply xs+ writeArray arrOut i' ys++ return_+
+ Data/Array/Accelerate/LLVM/PTX/CodeGen/Permute.hs view
@@ -0,0 +1,366 @@+{-# LANGUAGE GADTs #-}+{-# LANGUAGE OverloadedStrings #-}+{-# LANGUAGE RecordWildCards #-}+{-# LANGUAGE ScopedTypeVariables #-}+{-# LANGUAGE TemplateHaskell #-}+{-# LANGUAGE ViewPatterns #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.CodeGen.Permute+-- Copyright : [2016..2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.CodeGen.Permute (++ mkPermute,++) where++-- accelerate+import Data.Array.Accelerate.Analysis.Type+import Data.Array.Accelerate.Array.Sugar ( Array, Vector, Shape, Elt, eltType )+import Data.Array.Accelerate.Error+import qualified Data.Array.Accelerate.Array.Sugar as S++import Data.Array.Accelerate.LLVM.CodeGen.Arithmetic as A+import Data.Array.Accelerate.LLVM.CodeGen.Array+import Data.Array.Accelerate.LLVM.CodeGen.Base+import Data.Array.Accelerate.LLVM.CodeGen.Constant+import Data.Array.Accelerate.LLVM.CodeGen.Environment+import Data.Array.Accelerate.LLVM.CodeGen.Exp+import Data.Array.Accelerate.LLVM.CodeGen.IR+import Data.Array.Accelerate.LLVM.CodeGen.Monad+import Data.Array.Accelerate.LLVM.CodeGen.Permute+import Data.Array.Accelerate.LLVM.CodeGen.Ptr+import Data.Array.Accelerate.LLVM.CodeGen.Sugar++import Data.Array.Accelerate.LLVM.PTX.CodeGen.Base+import Data.Array.Accelerate.LLVM.PTX.CodeGen.Loop+import Data.Array.Accelerate.LLVM.PTX.Context+import Data.Array.Accelerate.LLVM.PTX.Target++import LLVM.AST.Type.AddrSpace+import LLVM.AST.Type.Instruction+import LLVM.AST.Type.Instruction.Atomic+import LLVM.AST.Type.Instruction.RMW as RMW+import LLVM.AST.Type.Instruction.Volatile+import LLVM.AST.Type.Operand+import LLVM.AST.Type.Representation++import Foreign.CUDA.Analysis++import Data.Typeable+import Control.Monad ( void )+import Prelude+++-- Forward permutation specified by an indexing mapping. The resulting array is+-- initialised with the given defaults, and any further values that are permuted+-- into the result array are added to the current value using the combination+-- function.+--+-- The combination function must be /associative/ and /commutative/. Elements+-- that are mapped to the magic index 'ignore' are dropped.+--+-- Parallel forward permutation has to take special care because different+-- threads could concurrently try to update the same memory location. Where+-- available we make use of special atomic instructions and other optimisations,+-- but in the general case each element of the output array has a lock which+-- must be obtained by the thread before it can update that memory location.+--+-- TODO: After too many failures to acquire the lock on an element, the thread+-- should back off and try a different element, adding this failed element to+-- a queue or some such.+--+mkPermute+ :: forall aenv sh sh' e. (Shape sh, Shape sh', Elt e)+ => PTX+ -> Gamma aenv+ -> IRPermuteFun PTX aenv (e -> e -> e)+ -> IRFun1 PTX aenv (sh -> sh')+ -> IRDelayed PTX aenv (Array sh e)+ -> CodeGen (IROpenAcc PTX aenv (Array sh' e))+mkPermute ptx aenv IRPermuteFun{..} project arr =+ let+ bytes = sizeOf (eltType (undefined :: e))+ sizeOk = bytes == 4 || bytes == 8+ in+ case atomicRMW of+ Just (rmw, f) | sizeOk -> mkPermute_rmw ptx aenv rmw f project arr+ _ -> mkPermute_mutex ptx aenv combine project arr+++-- Parallel forward permutation function which uses atomic instructions to+-- implement lock-free array updates.+--+-- Atomic instruction support on CUDA devices is a bit patchy, so depending on+-- the element type and compute capability of the target hardware we may need to+-- emulate the operation using atomic compare-and-swap.+--+-- Int32 Int64 Float32 Float64+-- +----------------------------------------+-- (+) | 2.0 2.0 2.0 6.0+-- (-) | 2.0 2.0 x x+-- (.&.) | 2.0 3.2+-- (.|.) | 2.0 3.2+-- xor | 2.0 3.2+-- min | 2.0 3.2 x x+-- max | 2.0 3.2 x x+-- CAS | 2.0 2.0+--+-- Note that NVPTX requires at least compute 2.0, so we can always implement the+-- lockfree update operations in terms of compare-and-swap.+--+mkPermute_rmw+ :: forall aenv sh sh' e. (Shape sh, Shape sh', Elt e)+ => PTX+ -> Gamma aenv+ -> RMWOperation+ -> IRFun1 PTX aenv (e -> e)+ -> IRFun1 PTX aenv (sh -> sh')+ -> IRDelayed PTX aenv (Array sh e)+ -> CodeGen (IROpenAcc PTX aenv (Array sh' e))+mkPermute_rmw ptx@(deviceProperties . ptxContext -> dev) aenv rmw update project IRDelayed{..} =+ let+ (start, end, paramGang) = gangParam+ (arrOut, paramOut) = mutableArray ("out" :: Name (Array sh' e))+ paramEnv = envParam aenv+ --+ bytes = sizeOf (eltType (undefined :: e))+ compute = computeCapability dev+ compute32 = Compute 3 2+ -- compute60 = Compute 6 0+ in+ makeOpenAcc ptx "permute_rmw" (paramGang ++ paramOut ++ paramEnv) $ do++ sh <- delayedExtent++ imapFromTo start end $ \i -> do++ i' <- A.fromIntegral integralType numType i+ ix <- indexOfInt sh i'+ ix' <- app1 project ix++ unless (ignore ix') $ do+ j <- intOfIndex (irArrayShape arrOut) ix'+ x <- app1 delayedLinearIndex i'+ r <- app1 update x++ case rmw of+ Exchange+ -> writeArray arrOut j r+ --+ _ | SingleTuple s <- eltType (undefined::e)+ , Just adata <- gcast (irArrayData arrOut)+ , Just r' <- gcast r+ -> do+ addr <- instr' $ GetElementPtr (asPtr defaultAddrSpace (op s adata)) [op integralType j]+ --+ let+ rmw_integral :: IntegralType t -> Operand (Ptr t) -> Operand t -> CodeGen ()+ rmw_integral t ptr val+ | primOk = void . instr' $ AtomicRMW t NonVolatile rmw ptr val (CrossThread, AcquireRelease)+ | otherwise =+ case rmw of+ RMW.And -> atomicCAS_rmw s' (A.band t (ir t val)) ptr+ RMW.Or -> atomicCAS_rmw s' (A.bor t (ir t val)) ptr+ RMW.Xor -> atomicCAS_rmw s' (A.xor t (ir t val)) ptr+ RMW.Min -> atomicCAS_cmp s' A.lt ptr val+ RMW.Max -> atomicCAS_cmp s' A.gt ptr val+ _ -> $internalError "mkPermute_rmw.integral" "unexpected transition"+ where+ s' = NumScalarType (IntegralNumType t)+ primOk = compute >= compute32+ || bytes == 4+ || case rmw of+ RMW.Add -> True+ RMW.Sub -> True+ _ -> False++ rmw_floating :: FloatingType t -> Operand (Ptr t) -> Operand t -> CodeGen ()+ rmw_floating t ptr val =+ case rmw of+ RMW.Min -> atomicCAS_cmp s' A.lt ptr val+ RMW.Max -> atomicCAS_cmp s' A.gt ptr val+ RMW.Sub -> atomicCAS_rmw s' (A.sub n (ir t val)) ptr+ RMW.Add+ | primAdd -> atomicAdd_f t ptr val+ | otherwise -> atomicCAS_rmw s' (A.add n (ir t val)) ptr+ _ -> $internalError "mkPermute_rmw.floating" "unexpected transition"+ where+ n = FloatingNumType t+ s' = NumScalarType n+ primAdd = bytes == 4+ -- Disabling due to missing support from llvm-4.0.+ -- <https://github.com/AccelerateHS/accelerate/issues/363>+ -- compute >= compute60++ rmw_nonnum :: NonNumType t -> Operand (Ptr t) -> Operand t -> CodeGen ()+ rmw_nonnum TypeChar{} ptr val = do+ ptr32 <- instr' $ PtrCast (primType :: PrimType (Ptr Word32)) ptr+ val32 <- instr' $ BitCast (scalarType :: ScalarType Word32) val+ void $ instr' $ AtomicRMW (integralType :: IntegralType Word32) NonVolatile rmw ptr32 val32 (CrossThread, AcquireRelease)+ rmw_nonnum _ _ _ = -- C character types are 8-bit, and thus not supported+ $internalError "mkPermute_rmw.nonnum" "unexpected transition"+ case s of+ NumScalarType (IntegralNumType t) -> rmw_integral t addr (op t r')+ NumScalarType (FloatingNumType t) -> rmw_floating t addr (op t r')+ NonNumScalarType t -> rmw_nonnum t addr (op t r')+ --+ _ -> $internalError "mkPermute_rmw" "unexpected transition"++ return_+++-- Parallel forward permutation function which uses a spinlock to acquire+-- a mutex before updating the value at that location.+--+mkPermute_mutex+ :: forall aenv sh sh' e. (Shape sh, Shape sh', Elt e)+ => PTX+ -> Gamma aenv+ -> IRFun2 PTX aenv (e -> e -> e)+ -> IRFun1 PTX aenv (sh -> sh')+ -> IRDelayed PTX aenv (Array sh e)+ -> CodeGen (IROpenAcc PTX aenv (Array sh' e))+mkPermute_mutex ptx aenv combine project IRDelayed{..} =+ let+ (start, end, paramGang) = gangParam+ (arrOut, paramOut) = mutableArray ("out" :: Name (Array sh' e))+ (arrLock, paramLock) = mutableArray ("lock" :: Name (Vector Word32))+ paramEnv = envParam aenv+ in+ makeOpenAcc ptx "permute_mutex" (paramGang ++ paramOut ++ paramLock ++ paramEnv) $ do++ sh <- delayedExtent++ imapFromTo start end $ \i -> do++ i' <- A.fromIntegral integralType numType i+ ix <- indexOfInt sh i'+ ix' <- app1 project ix++ -- project element onto the destination array and (atomically) update+ unless (ignore ix') $ do+ j <- intOfIndex (irArrayShape arrOut) ix'+ x <- app1 delayedLinearIndex i'++ atomically arrLock j $ do+ y <- readArray arrOut j+ r <- app2 combine x y+ writeArray arrOut j r++ return_+++-- Atomically execute the critical section only when the lock at the given array+-- index is obtained. The thread spins waiting for the lock to be released and+-- there is no backoff strategy in case the lock is contended.+--+-- The canonical implementation of a spin-lock looks like this:+--+-- > do {+-- > old = atomic_exchange(&lock[i], 1);+-- > } while (old == 1);+-- >+-- > /* critical section */+-- >+-- > atomic_exchange(&lock[i], 0);+--+-- The initial loop repeatedly attempts to take the lock by writing a 1 (locked)+-- into the lock slot. Once the 'old' state of the lock returns 0 (unlocked),+-- then we just acquired the lock and the atomic section can be computed.+-- Finally, the lock is released by writing 0 back to the lock slot.+--+-- However, there is a complication with CUDA devices because all threads in+-- a warp must execute in lockstep (with predicated execution). In the above+-- setup, once a thread acquires a lock, then it will be disabled and stop+-- participating in the loop, waiting for all other threads (to acquire their+-- locks) before continuing program execution. If two threads in the same warp+-- attempt to acquire the same lock, then once the lock is acquired by one+-- thread then it will sit idle waiting while the second thread spins attempting+-- to grab a lock that will never be released because the first thread (which+-- holds the lock) can not make progress. DEADLOCK.+--+-- To prevent this situation we must invert the algorithm so that threads can+-- always make progress, until each warp in the thread has committed their+-- result.+--+-- > done = 0;+-- > do {+-- > if ( atomic_exchange(&lock[i], 1) == 0 ) {+-- >+-- > /* critical section */+-- >+-- > done = 1;+-- > atomic_exchange(&lock[i], 0);+-- > }+-- > } while ( done == 0 );+--+atomically+ :: IRArray (Vector Word32)+ -> IR Int+ -> CodeGen a+ -> CodeGen a+atomically barriers i action = do+ let+ lock = integral integralType 1+ unlock = integral integralType 0+ unlock' = lift 0+ --+ spin <- newBlock "spinlock.entry"+ crit <- newBlock "spinlock.critical-start"+ skip <- newBlock "spinlock.critical-end"+ exit <- newBlock "spinlock.exit"++ addr <- instr' $ GetElementPtr (asPtr defaultAddrSpace (op integralType (irArrayData barriers))) [op integralType i]+ _ <- br spin++ -- Loop until this thread has completed its critical section. If the slot was+ -- unlocked then we just acquired the lock and the thread can perform the+ -- critical section, otherwise skip to the bottom of the critical section.+ setBlock spin+ old <- instr $ AtomicRMW integralType NonVolatile Exchange addr lock (CrossThread, Acquire)+ ok <- A.eq scalarType old unlock'+ no <- cbr ok crit skip++ -- If we just acquired the lock, execute the critical section+ setBlock crit+ r <- action+ _ <- instr $ AtomicRMW integralType NonVolatile Exchange addr unlock (CrossThread, Release)+ yes <- br skip++ -- At the base of the critical section, threads participate in a memory fence+ -- to ensure the lock state is committed to memory. Depending on which+ -- incoming edge the thread arrived at this block from determines whether they+ -- have completed their critical section.+ setBlock skip+ done <- phi [(lift True, yes), (lift False, no)]++ __syncthreads+ _ <- cbr done exit spin++ setBlock exit+ return r+++-- Helper functions+-- ----------------++-- Test whether the given index is the magic value 'ignore'. This operates+-- strictly rather than performing short-circuit (&&).+--+ignore :: forall ix. Shape ix => IR ix -> CodeGen (IR Bool)+ignore (IR ix) = go (S.eltType (undefined::ix)) (S.fromElt (S.ignore::ix)) ix+ where+ go :: TupleType t -> t -> Operands t -> CodeGen (IR Bool)+ go UnitTuple () OP_Unit = return (lift True)+ go (PairTuple tsh tsz) (ish, isz) (OP_Pair sh sz) = do x <- go tsh ish sh+ y <- go tsz isz sz+ land' x y+ go (SingleTuple t) ig sz = A.eq t (ir t (scalar t ig)) (ir t (op' t sz))+
+ Data/Array/Accelerate/LLVM/PTX/CodeGen/Queue.hs view
@@ -0,0 +1,118 @@+{-# LANGUAGE OverloadedStrings #-}+{-# LANGUAGE ScopedTypeVariables #-}+{-# LANGUAGE ViewPatterns #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.CodeGen.Queue+-- Copyright : [2014..2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--+-- Abstractions for creating simply dynamically scheduled work queues. This+-- works by atomically incrementing a global counter (in global memory) and+-- distributing this result to each thread in the block (via shared memory).+-- Thus there is an additional ~1000 cycle overhead for a thread block to+-- determine their next work item. This also implies all thread blocks are+-- contending for the same global counter.+--+-- In practice this extra overhead is not always worth paying. We use it for+-- segmented reductions, because the length of each segment is unknown apriori+-- and the entire thread block participates in the reduction of a segment. On+-- the other hand, the arithmetically unbalanced mandelbrot fractal program was+-- (generally) slower with this addition, so for now at least keep (morally)+-- balanced operations (map, generate) with a static schedule. (Admittidely this+-- test was on my very old 650M, so newer/more powerful GPUs with faster atomic+-- instructions or more inflight thread blocks could benefit more.)+--++module Data.Array.Accelerate.LLVM.PTX.CodeGen.Queue+ where++import Data.Array.Accelerate.LLVM.CodeGen.Arithmetic as A+import Data.Array.Accelerate.LLVM.CodeGen.Downcast+import Data.Array.Accelerate.LLVM.CodeGen.IR+import Data.Array.Accelerate.LLVM.CodeGen.Monad+import Data.Array.Accelerate.LLVM.CodeGen.Sugar++import Data.Array.Accelerate.LLVM.PTX.Analysis.Launch+import Data.Array.Accelerate.LLVM.PTX.CodeGen.Base+import Data.Array.Accelerate.LLVM.PTX.Target++import LLVM.AST.Type.Constant+import LLVM.AST.Type.Instruction+import LLVM.AST.Type.Instruction.Atomic+import LLVM.AST.Type.Instruction.Volatile+import LLVM.AST.Type.Operand+import LLVM.AST.Type.Representation+import qualified LLVM.AST.Global as LLVM+import qualified LLVM.AST.Linkage as LLVM+import qualified LLVM.AST.Name as LLVM+import qualified LLVM.AST.Type as LLVM+import qualified LLVM.AST.Type.Instruction.RMW as RMW+++-- Interface+-- ---------++type WorkQueue = (Operand (Ptr Int32), Operand (Ptr Int32))++-- Declare a new dynamically scheduled global work queue. Don't forget to+-- initialise the queue with the kernel generated by 'mkQueueInit'.+--+globalWorkQueue :: CodeGen WorkQueue+globalWorkQueue = do+ sn <- freshName+ declare $ LLVM.globalVariableDefaults+ { LLVM.name = LLVM.Name "__queue__"+ , LLVM.type' = LLVM.IntegerType 32+ , LLVM.alignment = 4+ }+ declare $ LLVM.globalVariableDefaults+ { LLVM.name = downcast sn+ , LLVM.addrSpace = sharedMemAddrSpace+ , LLVM.type' = LLVM.IntegerType 32+ , LLVM.linkage = LLVM.Internal+ , LLVM.alignment = 4+ }+ return ( ConstantOperand (GlobalReference type' "__queue__")+ , ConstantOperand (GlobalReference type' sn) )+++-- Dequeue the next 'n' items from the work queue for evaluation by the calling+-- thread block. Each thread in the thread block receives the index of the start+-- of the newly acquired range.+--+dequeue :: WorkQueue -> IR Int32 -> CodeGen (IR Int32)+dequeue (queue, smem) n = do+ tid <- threadIdx+ when (A.eq scalarType tid (lift 0)) $ do+ v <- instr' $ AtomicRMW integralType NonVolatile RMW.Add queue (op integralType n) (CrossThread, AcquireRelease)+ _ <- instr' $ Store Volatile smem v+ return ()+ --+ __syncthreads+ v <- instr' $ Load scalarType Volatile smem+ return (ir integralType v)+++-- Initialisation kernel+-- ---------------------++-- This kernel is used to initialise the dynamically scheduled work queue. It+-- must be called before the main kernel, which uses the work queue, is invoked.+--+mkQueueInit+ :: DeviceProperties+ -> CodeGen (IROpenAcc PTX aenv a)+mkQueueInit dev =+ let+ (start, _end, paramGang) = gangParam+ config = launchConfig dev [1] (\_ -> 0) (\_ _ -> 1)+ in+ makeOpenAccWith config "qinit" paramGang $ do+ (queue,_) <- globalWorkQueue+ _ <- instr' $ Store Volatile queue (op integralType start)+ return_+
+ Data/Array/Accelerate/LLVM/PTX/CodeGen/Scan.hs view
@@ -0,0 +1,1332 @@+{-# LANGUAGE GADTs #-}+{-# LANGUAGE OverloadedStrings #-}+{-# LANGUAGE PatternGuards #-}+{-# LANGUAGE RebindableSyntax #-}+{-# LANGUAGE RecordWildCards #-}+{-# LANGUAGE ScopedTypeVariables #-}+{-# LANGUAGE TypeOperators #-}+{-# LANGUAGE ViewPatterns #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.CodeGen.Scan+-- Copyright : [2016..2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.CodeGen.Scan (++ mkScanl, mkScanl1, mkScanl',+ mkScanr, mkScanr1, mkScanr',++) where++-- accelerate+import Data.Array.Accelerate.Analysis.Type+import Data.Array.Accelerate.Array.Sugar++import Data.Array.Accelerate.LLVM.Analysis.Match+import Data.Array.Accelerate.LLVM.CodeGen.Arithmetic as A+import Data.Array.Accelerate.LLVM.CodeGen.Array+import Data.Array.Accelerate.LLVM.CodeGen.Base+import Data.Array.Accelerate.LLVM.CodeGen.Constant+import Data.Array.Accelerate.LLVM.CodeGen.Environment+import Data.Array.Accelerate.LLVM.CodeGen.Exp+import Data.Array.Accelerate.LLVM.CodeGen.IR+import Data.Array.Accelerate.LLVM.CodeGen.Loop+import Data.Array.Accelerate.LLVM.CodeGen.Monad+import Data.Array.Accelerate.LLVM.CodeGen.Sugar+import Data.Array.Accelerate.LLVM.PTX.Analysis.Launch+import Data.Array.Accelerate.LLVM.PTX.CodeGen.Base+import Data.Array.Accelerate.LLVM.PTX.CodeGen.Generate+import Data.Array.Accelerate.LLVM.PTX.Context+import Data.Array.Accelerate.LLVM.PTX.Target++import LLVM.AST.Type.Representation++import qualified Foreign.CUDA.Analysis as CUDA++import Control.Applicative+import Control.Monad ( (>=>), void )+import Data.String ( fromString )+import Data.Coerce as Safe+import Data.Bits as P+import Prelude as P hiding ( last )+++data Direction = L | R++-- 'Data.List.scanl' style left-to-right exclusive scan, but with the+-- restriction that the combination function must be associative to enable+-- efficient parallel implementation.+--+-- > scanl (+) 10 (use $ fromList (Z :. 10) [0..])+-- >+-- > ==> Array (Z :. 11) [10,10,11,13,16,20,25,31,38,46,55]+--+mkScanl+ :: forall aenv sh e. (Shape sh, Elt e)+ => PTX+ -> Gamma aenv+ -> IRFun2 PTX aenv (e -> e -> e)+ -> IRExp PTX aenv e+ -> IRDelayed PTX aenv (Array (sh:.Int) e)+ -> CodeGen (IROpenAcc PTX aenv (Array (sh:.Int) e))+mkScanl ptx@(deviceProperties . ptxContext -> dev) aenv combine seed arr+ | Just Refl <- matchShapeType (undefined::sh) (undefined::Z)+ = foldr1 (+++) <$> sequence [ mkScanAllP1 L dev aenv combine (Just seed) arr+ , mkScanAllP2 L dev aenv combine+ , mkScanAllP3 L dev aenv combine (Just seed)+ , mkScanFill ptx aenv seed+ ]+ --+ | otherwise+ = (+++) <$> mkScanDim L dev aenv combine (Just seed) arr+ <*> mkScanFill ptx aenv seed+++-- 'Data.List.scanl1' style left-to-right inclusive scan, but with the+-- restriction that the combination function must be associative to enable+-- efficient parallel implementation. The array must not be empty.+--+-- > scanl1 (+) (use $ fromList (Z :. 10) [0..])+-- >+-- > ==> Array (Z :. 10) [0,1,3,6,10,15,21,28,36,45]+--+mkScanl1+ :: forall aenv sh e. (Shape sh, Elt e)+ => PTX+ -> Gamma aenv+ -> IRFun2 PTX aenv (e -> e -> e)+ -> IRDelayed PTX aenv (Array (sh:.Int) e)+ -> CodeGen (IROpenAcc PTX aenv (Array (sh:.Int) e))+mkScanl1 (deviceProperties . ptxContext -> dev) aenv combine arr+ | Just Refl <- matchShapeType (undefined::sh) (undefined::Z)+ = foldr1 (+++) <$> sequence [ mkScanAllP1 L dev aenv combine Nothing arr+ , mkScanAllP2 L dev aenv combine+ , mkScanAllP3 L dev aenv combine Nothing+ ]+ --+ | otherwise+ = mkScanDim L dev aenv combine Nothing arr+++-- Variant of 'scanl' where the final result is returned in a separate array.+--+-- > scanr' (+) 10 (use $ fromList (Z :. 10) [0..])+-- >+-- > ==> ( Array (Z :. 10) [10,10,11,13,16,20,25,31,38,46]+-- , Array Z [55]+-- )+--+mkScanl'+ :: forall aenv sh e. (Shape sh, Elt e)+ => PTX+ -> Gamma aenv+ -> IRFun2 PTX aenv (e -> e -> e)+ -> IRExp PTX aenv e+ -> IRDelayed PTX aenv (Array (sh:.Int) e)+ -> CodeGen (IROpenAcc PTX aenv (Array (sh:.Int) e, Array sh e))+mkScanl' ptx@(deviceProperties . ptxContext -> dev) aenv combine seed arr+ | Just Refl <- matchShapeType (undefined::sh) (undefined::Z)+ = foldr1 (+++) <$> sequence [ mkScan'AllP1 L dev aenv combine seed arr+ , mkScan'AllP2 L dev aenv combine+ , mkScan'AllP3 L dev aenv combine+ , mkScan'Fill ptx aenv seed+ ]+ --+ | otherwise+ = (+++) <$> mkScan'Dim L dev aenv combine seed arr+ <*> mkScan'Fill ptx aenv seed+++-- 'Data.List.scanr' style right-to-left exclusive scan, but with the+-- restriction that the combination function must be associative to enable+-- efficient parallel implementation.+--+-- > scanr (+) 10 (use $ fromList (Z :. 10) [0..])+-- >+-- > ==> Array (Z :. 11) [55,55,54,52,49,45,40,34,27,19,10]+--+mkScanr+ :: forall aenv sh e. (Shape sh, Elt e)+ => PTX+ -> Gamma aenv+ -> IRFun2 PTX aenv (e -> e -> e)+ -> IRExp PTX aenv e+ -> IRDelayed PTX aenv (Array (sh:.Int) e)+ -> CodeGen (IROpenAcc PTX aenv (Array (sh:.Int) e))+mkScanr ptx@(deviceProperties . ptxContext -> dev) aenv combine seed arr+ | Just Refl <- matchShapeType (undefined::sh) (undefined::Z)+ = foldr1 (+++) <$> sequence [ mkScanAllP1 R dev aenv combine (Just seed) arr+ , mkScanAllP2 R dev aenv combine+ , mkScanAllP3 R dev aenv combine (Just seed)+ , mkScanFill ptx aenv seed+ ]+ --+ | otherwise+ = (+++) <$> mkScanDim R dev aenv combine (Just seed) arr+ <*> mkScanFill ptx aenv seed+++-- 'Data.List.scanr1' style right-to-left inclusive scan, but with the+-- restriction that the combination function must be associative to enable+-- efficient parallel implementation. The array must not be empty.+--+-- > scanr (+) 10 (use $ fromList (Z :. 10) [0..])+-- >+-- > ==> Array (Z :. 10) [45,45,44,42,39,35,30,24,17,9]+--+mkScanr1+ :: forall aenv sh e. (Shape sh, Elt e)+ => PTX+ -> Gamma aenv+ -> IRFun2 PTX aenv (e -> e -> e)+ -> IRDelayed PTX aenv (Array (sh:.Int) e)+ -> CodeGen (IROpenAcc PTX aenv (Array (sh:.Int) e))+mkScanr1 (deviceProperties . ptxContext -> dev) aenv combine arr+ | Just Refl <- matchShapeType (undefined::sh) (undefined::Z)+ = foldr1 (+++) <$> sequence [ mkScanAllP1 R dev aenv combine Nothing arr+ , mkScanAllP2 R dev aenv combine+ , mkScanAllP3 R dev aenv combine Nothing+ ]+ --+ | otherwise+ = mkScanDim R dev aenv combine Nothing arr+++-- Variant of 'scanr' where the final result is returned in a separate array.+--+-- > scanr' (+) 10 (use $ fromList (Z :. 10) [0..])+-- >+-- > ==> ( Array (Z :. 10) [55,54,52,49,45,40,34,27,19,10]+-- , Array Z [55]+-- )+--+mkScanr'+ :: forall aenv sh e. (Shape sh, Elt e)+ => PTX+ -> Gamma aenv+ -> IRFun2 PTX aenv (e -> e -> e)+ -> IRExp PTX aenv e+ -> IRDelayed PTX aenv (Array (sh:.Int) e)+ -> CodeGen (IROpenAcc PTX aenv (Array (sh:.Int) e, Array sh e))+mkScanr' ptx@(deviceProperties . ptxContext -> dev) aenv combine seed arr+ | Just Refl <- matchShapeType (undefined::sh) (undefined::Z)+ = foldr1 (+++) <$> sequence [ mkScan'AllP1 R dev aenv combine seed arr+ , mkScan'AllP2 R dev aenv combine+ , mkScan'AllP3 R dev aenv combine+ , mkScan'Fill ptx aenv seed+ ]+ --+ | otherwise+ = (+++) <$> mkScan'Dim R dev aenv combine seed arr+ <*> mkScan'Fill ptx aenv seed+++-- Device wide scans+-- -----------------+--+-- This is a classic two-pass algorithm which proceeds in two phases and+-- requires ~4n data movement to global memory. In future we would like to+-- replace this with a single pass algorithm.+--++-- Parallel scan, step 1.+--+-- Threads scan a stripe of the input into a temporary array, incorporating the+-- initial element and any fused functions on the way. The final reduction+-- result of this chunk is written to a separate array.+--+mkScanAllP1+ :: forall aenv e. Elt e+ => Direction+ -> DeviceProperties -- ^ properties of the target GPU+ -> Gamma aenv -- ^ array environment+ -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function+ -> Maybe (IRExp PTX aenv e) -- ^ seed element, if this is an exclusive scan+ -> IRDelayed PTX aenv (Vector e) -- ^ input data+ -> CodeGen (IROpenAcc PTX aenv (Vector e))+mkScanAllP1 dir dev aenv combine mseed IRDelayed{..} =+ let+ (start, end, paramGang) = gangParam+ (arrOut, paramOut) = mutableArray ("out" :: Name (Vector e))+ (arrTmp, paramTmp) = mutableArray ("tmp" :: Name (Vector e))+ paramEnv = envParam aenv+ --+ config = launchConfig dev (CUDA.incWarp dev) smem const+ smem n = warps * (1 + per_warp) * bytes+ where+ ws = CUDA.warpSize dev+ warps = n `div` ws+ per_warp = ws + ws `div` 2+ bytes = sizeOf (eltType (undefined :: e))+ in+ makeOpenAccWith config "scanP1" (paramGang ++ paramTmp ++ paramOut ++ paramEnv) $ do++ -- Size of the input array+ sz <- A.fromIntegral integralType numType . indexHead =<< delayedExtent++ -- A thread block scans a non-empty stripe of the input, storing the final+ -- block-wide aggregate into a separate array+ --+ -- For exclusive scans, thread 0 of segment 0 must incorporate the initial+ -- element into the input and output. Threads shuffle their indices+ -- appropriately.+ --+ bid <- blockIdx+ gd <- gridDim+ s0 <- A.add numType start bid++ -- iterating over thread-block-wide segments+ imapFromStepTo s0 gd end $ \chunk -> do++ bd <- blockDim+ inf <- A.mul numType chunk bd++ -- index i* is the index that this thread will read data from. Recall that+ -- the supremum index is exclusive+ tid <- threadIdx+ i0 <- case dir of+ L -> A.add numType inf tid+ R -> do x <- A.sub numType sz inf+ y <- A.sub numType x tid+ z <- A.sub numType y (lift 1)+ return z++ -- index j* is the index that we write to. Recall that for exclusive scans+ -- the output array is one larger than the input; the initial element will+ -- be written into this spot by thread 0 of the first thread block.+ j0 <- case mseed of+ Nothing -> return i0+ Just _ -> case dir of+ L -> A.add numType i0 (lift 1)+ R -> return i0++ -- If this thread has input, read data and participate in thread-block scan+ let valid i = case dir of+ L -> A.lt scalarType i sz+ R -> A.gte scalarType i (lift 0)++ when (valid i0) $ do+ x0 <- app1 delayedLinearIndex =<< A.fromIntegral integralType numType i0+ x1 <- case mseed of+ Nothing -> return x0+ Just seed ->+ if A.eq scalarType tid (lift 0) `A.land` A.eq scalarType chunk (lift 0)+ then do+ z <- seed+ case dir of+ L -> writeArray arrOut (lift 0 :: IR Int32) z >> app2 combine z x0+ R -> writeArray arrOut sz z >> app2 combine x0 z+ else+ return x0++ n <- A.sub numType sz inf+ x2 <- if A.gte scalarType n bd+ then scanBlockSMem dir dev combine Nothing x1+ else scanBlockSMem dir dev combine (Just n) x1++ -- Write this thread's scan result to memory+ writeArray arrOut j0 x2++ -- The last thread also writes its result---the aggregate for this+ -- thread block---to the temporary partial sums array. This is only+ -- necessary for full blocks in a multi-block scan; the final+ -- partially-full tile does not have a successor block.+ last <- A.sub numType bd (lift 1)+ when (A.gt scalarType gd (lift 1) `land` A.eq scalarType tid last) $+ case dir of+ L -> writeArray arrTmp chunk x2+ R -> do u <- A.sub numType end chunk+ v <- A.sub numType u (lift 1)+ writeArray arrTmp v x2++ return_+++-- Parallel scan, step 2+--+-- A single thread block performs a scan of the per-block aggregates computed in+-- step 1. This gives the per-block prefix which must be added to each element+-- in step 3.+--+mkScanAllP2+ :: forall aenv e. Elt e+ => Direction+ -> DeviceProperties -- ^ properties of the target GPU+ -> Gamma aenv -- ^ array environment+ -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function+ -> CodeGen (IROpenAcc PTX aenv (Vector e))+mkScanAllP2 dir dev aenv combine =+ let+ (start, end, paramGang) = gangParam+ (arrTmp, paramTmp) = mutableArray ("tmp" :: Name (Vector e))+ paramEnv = envParam aenv+ --+ config = launchConfig dev (CUDA.incWarp dev) smem grid+ grid _ _ = 1+ smem n = warps * (1 + per_warp) * bytes+ where+ ws = CUDA.warpSize dev+ warps = n `div` ws+ per_warp = ws + ws `div` 2+ bytes = sizeOf (eltType (undefined :: e))+ in+ makeOpenAccWith config "scanP2" (paramGang ++ paramTmp ++ paramEnv) $ do++ -- The first and last threads of the block need to communicate the+ -- block-wide aggregate as a carry-in value across iterations.+ --+ -- TODO: We could optimise this a bit if we can get access to the shared+ -- memory area used by 'scanBlockSMem', and from there directly read the+ -- value computed by the last thread.+ carry <- staticSharedMem 1++ bd <- blockDim+ imapFromStepTo start bd end $ \offset -> do++ -- Index of the partial sums array that this thread will process.+ tid <- threadIdx+ i0 <- case dir of+ L -> A.add numType offset tid+ R -> do x <- A.sub numType end offset+ y <- A.sub numType x tid+ z <- A.sub numType y (lift 1)+ return z++ let valid i = case dir of+ L -> A.lt scalarType i end+ R -> A.gte scalarType i start++ when (valid i0) $ do++ __syncthreads++ x0 <- readArray arrTmp i0+ x1 <- if A.gt scalarType offset (lift 0) `land` A.eq scalarType tid (lift 0)+ then do+ c <- readArray carry (lift 0 :: IR Int32)+ case dir of+ L -> app2 combine c x0+ R -> app2 combine x0 c+ else do+ return x0++ n <- A.sub numType end offset+ x2 <- if A.gte scalarType n bd+ then scanBlockSMem dir dev combine Nothing x1+ else scanBlockSMem dir dev combine (Just n) x1++ -- Update the temporary array with this thread's result+ writeArray arrTmp i0 x2++ -- The last thread writes the carry-out value. If the last thread is not+ -- active, then this must be the last stripe anyway.+ last <- A.sub numType bd (lift 1)+ when (A.eq scalarType tid last) $+ writeArray carry (lift 0 :: IR Int32) x2++ return_+++-- Parallel scan, step 3.+--+-- Threads combine every element of the partial block results with the carry-in+-- value computed in step 2.+--+mkScanAllP3+ :: forall aenv e. Elt e+ => Direction+ -> DeviceProperties -- ^ properties of the target GPU+ -> Gamma aenv -- ^ array environment+ -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function+ -> Maybe (IRExp PTX aenv e) -- ^ seed element, if this is an exclusive scan+ -> CodeGen (IROpenAcc PTX aenv (Vector e))+mkScanAllP3 dir dev aenv combine mseed =+ let+ (start, end, paramGang) = gangParam+ (arrOut, paramOut) = mutableArray ("out" :: Name (Vector e))+ (arrTmp, paramTmp) = mutableArray ("tmp" :: Name (Vector e))+ paramEnv = envParam aenv+ --+ stride = local scalarType ("ix.stride" :: Name Int32)+ paramStride = scalarParameter scalarType ("ix.stride" :: Name Int32)+ --+ config = launchConfig dev (CUDA.incWarp dev) (const 0) const+ in+ makeOpenAccWith config "scanP3" (paramGang ++ paramTmp ++ paramOut ++ paramStride : paramEnv) $ do++ sz <- A.fromIntegral integralType numType (indexHead (irArrayShape arrOut))+ tid <- threadIdx++ -- Threads that will never contribute can just exit immediately. The size of+ -- each chunk is set by the block dimension of the step 1 kernel, which may+ -- be different from the block size of this kernel.+ when (A.lt scalarType tid stride) $ do++ -- Iterate over the segments computed in phase 1. Note that we have one+ -- fewer chunk to process because the first has no carry-in.+ bid <- blockIdx+ gd <- gridDim+ c0 <- A.add numType start bid+ imapFromStepTo c0 gd end $ \chunk -> do++ -- Determine the start and end indicies of this chunk to which we will+ -- carry-in the value. Returned for left-to-right traversal.+ (inf,sup) <- case dir of+ L -> do+ a <- A.add numType chunk (lift 1)+ b <- A.mul numType stride a+ case mseed of+ Just{} -> do+ c <- A.add numType b (lift 1)+ d <- A.add numType c stride+ e <- A.min scalarType d sz+ return (c,e)+ Nothing -> do+ c <- A.add numType b stride+ d <- A.min scalarType c sz+ return (b,d)+ R -> do+ a <- A.sub numType end chunk+ b <- A.mul numType stride a+ c <- A.sub numType sz b+ case mseed of+ Just{} -> do+ d <- A.sub numType c (lift 1)+ e <- A.sub numType d stride+ f <- A.max scalarType e (lift 0)+ return (f,d)+ Nothing -> do+ d <- A.sub numType c stride+ e <- A.max scalarType d (lift 0)+ return (e,c)++ -- Read the carry-in value+ carry <- case dir of+ L -> readArray arrTmp chunk+ R -> do+ a <- A.add numType chunk (lift 1)+ b <- readArray arrTmp a+ return b++ -- Apply the carry-in value to each element in the chunk+ bd <- blockDim+ i0 <- A.add numType inf tid+ imapFromStepTo i0 bd sup $ \i -> do+ v <- readArray arrOut i+ u <- case dir of+ L -> app2 combine carry v+ R -> app2 combine v carry+ writeArray arrOut i u++ return_+++-- Parallel scan', step 1.+--+-- Similar to mkScanAllP1. Threads scan a stripe of the input into a temporary+-- array, incorporating the initial element and any fused functions on the way.+-- The final reduction result of this chunk is written to a separate array.+--+mkScan'AllP1+ :: forall aenv e. Elt e+ => Direction+ -> DeviceProperties+ -> Gamma aenv+ -> IRFun2 PTX aenv (e -> e -> e)+ -> IRExp PTX aenv e+ -> IRDelayed PTX aenv (Vector e)+ -> CodeGen (IROpenAcc PTX aenv (Vector e, Scalar e))+mkScan'AllP1 dir dev aenv combine seed IRDelayed{..} =+ let+ (start, end, paramGang) = gangParam+ (arrOut, paramOut) = mutableArray ("out" :: Name (Vector e))+ (arrTmp, paramTmp) = mutableArray ("tmp" :: Name (Vector e))+ paramEnv = envParam aenv+ --+ config = launchConfig dev (CUDA.incWarp dev) smem const+ smem n = warps * (1 + per_warp) * bytes+ where+ ws = CUDA.warpSize dev+ warps = n `div` ws+ per_warp = ws + ws `div` 2+ bytes = sizeOf (eltType (undefined :: e))+ in+ makeOpenAccWith config "scanP1" (paramGang ++ paramTmp ++ paramOut ++ paramEnv) $ do++ -- Size of the input array+ sz <- A.fromIntegral integralType numType . indexHead =<< delayedExtent++ -- A thread block scans a non-empty stripe of the input, storing the partial+ -- result and the final block-wide aggregate+ bid <- blockIdx+ gd <- gridDim+ s0 <- A.add numType start bid++ -- iterate over thread-block wide segments+ imapFromStepTo s0 gd end $ \seg -> do++ bd <- blockDim+ inf <- A.mul numType seg bd++ -- i* is the index that this thread will read data from+ tid <- threadIdx+ i0 <- case dir of+ L -> A.add numType inf tid+ R -> do x <- A.sub numType sz inf+ y <- A.sub numType x tid+ z <- A.sub numType y (lift 1)+ return z++ -- j* is the index this thread will write to. This is just shifted by one+ -- to make room for the initial element+ j0 <- case dir of+ L -> A.add numType i0 (lift 1)+ R -> A.sub numType i0 (lift 1)++ -- If this thread has input it participates in the scan+ let valid i = case dir of+ L -> A.lt scalarType i sz+ R -> A.gte scalarType i (lift 0)++ when (valid i0) $ do+ x0 <- app1 delayedLinearIndex =<< A.fromIntegral integralType numType i0++ -- Thread 0 of the first segment must also evaluate and store the+ -- initial element+ x1 <- if A.eq scalarType tid (lift 0) `A.land` A.eq scalarType seg (lift 0)+ then do+ z <- seed+ writeArray arrOut i0 z+ case dir of+ L -> app2 combine z x0+ R -> app2 combine x0 z+ else+ return x0++ -- Block-wide scan+ n <- A.sub numType sz inf+ x2 <- if A.gte scalarType n bd+ then scanBlockSMem dir dev combine Nothing x1+ else scanBlockSMem dir dev combine (Just n) x1++ -- Write this thread's scan result to memory. Recall that we had to make+ -- space for the initial element, so the very last thread does not store+ -- its result here.+ case dir of+ L -> when (A.lt scalarType j0 sz) $ writeArray arrOut j0 x2+ R -> when (A.gte scalarType j0 (lift 0)) $ writeArray arrOut j0 x2++ -- Last active thread writes its result to the partial sums array. These+ -- will be used to compute the carry-in value in step 2.+ m <- do x <- A.min scalarType n bd+ y <- A.sub numType x (lift 1)+ return y+ when (A.eq scalarType tid m) $+ case dir of+ L -> writeArray arrTmp seg x2+ R -> do x <- A.sub numType end seg+ y <- A.sub numType x (lift 1)+ writeArray arrTmp y x2++ return_+++-- Parallel scan', step 2+--+-- A single thread block performs an inclusive scan of the partial sums array to+-- compute the per-block carry-in values, as well as the final reduction result.+--+mkScan'AllP2+ :: forall aenv e. Elt e+ => Direction+ -> DeviceProperties+ -> Gamma aenv+ -> IRFun2 PTX aenv (e -> e -> e)+ -> CodeGen (IROpenAcc PTX aenv (Vector e, Scalar e))+mkScan'AllP2 dir dev aenv combine =+ let+ (start, end, paramGang) = gangParam+ (arrTmp, paramTmp) = mutableArray ("tmp" :: Name (Vector e))+ (arrSum, paramSum) = mutableArray ("sum" :: Name (Scalar e))+ paramEnv = envParam aenv+ --+ config = launchConfig dev (CUDA.incWarp dev) smem grid+ grid _ _ = 1+ smem n = warps * (1 + per_warp) * bytes+ where+ ws = CUDA.warpSize dev+ warps = n `div` ws+ per_warp = ws + ws `div` 2+ bytes = sizeOf (eltType (undefined :: e))+ in+ makeOpenAccWith config "scanP2" (paramGang ++ paramTmp ++ paramSum ++ paramEnv) $ do++ -- The first and last threads of the block need to communicate the+ -- block-wide aggregate as a carry-in value across iterations.+ carry <- staticSharedMem 1++ -- A single thread block iterates over the per-block partial results from+ -- step 1+ tid <- threadIdx+ bd <- blockDim+ imapFromStepTo start bd end $ \offset -> do++ i0 <- case dir of+ L -> A.add numType offset tid+ R -> do x <- A.sub numType end offset+ y <- A.sub numType x tid+ z <- A.sub numType y (lift 1)+ return z++ let valid i = case dir of+ L -> A.lt scalarType i end+ R -> A.gte scalarType i start++ when (valid i0) $ do++ -- wait for the carry-in value to be updated+ __syncthreads++ x0 <- readArray arrTmp i0+ x1 <- if A.gt scalarType offset (lift 0) `A.land` A.eq scalarType tid (lift 0)+ then do+ c <- readArray carry (lift 0 :: IR Int32)+ case dir of+ L -> app2 combine c x0+ R -> app2 combine x0 c+ else+ return x0++ n <- A.sub numType end offset+ x2 <- if A.gte scalarType n bd+ then scanBlockSMem dir dev combine Nothing x1+ else scanBlockSMem dir dev combine (Just n) x1++ -- Update the partial results array+ writeArray arrTmp i0 x2++ -- The last active thread saves its result as the carry-out value.+ m <- do x <- A.min scalarType bd n+ y <- A.sub numType x (lift 1)+ return y+ when (A.eq scalarType tid m) $+ writeArray carry (lift 0 :: IR Int32) x2++ -- First thread stores the final carry-out values at the final reduction+ -- result for the entire array+ __syncthreads++ when (A.eq scalarType tid (lift 0)) $+ writeArray arrSum (lift 0 :: IR Int32) =<< readArray carry (lift 0 :: IR Int32)++ return_+++-- Parallel scan', step 3.+--+-- Threads combine every element of the partial block results with the carry-in+-- value computed in step 2.+--+mkScan'AllP3+ :: forall aenv e. Elt e+ => Direction+ -> DeviceProperties -- ^ properties of the target GPU+ -> Gamma aenv -- ^ array environment+ -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function+ -> CodeGen (IROpenAcc PTX aenv (Vector e, Scalar e))+mkScan'AllP3 dir dev aenv combine =+ let+ (start, end, paramGang) = gangParam+ (arrOut, paramOut) = mutableArray ("out" :: Name (Vector e))+ (arrTmp, paramTmp) = mutableArray ("tmp" :: Name (Vector e))+ paramEnv = envParam aenv+ --+ stride = local scalarType ("ix.stride" :: Name Int32)+ paramStride = scalarParameter scalarType ("ix.stride" :: Name Int32)+ --+ config = launchConfig dev (CUDA.incWarp dev) (const 0) const+ in+ makeOpenAccWith config "scanP3" (paramGang ++ paramTmp ++ paramOut ++ paramStride : paramEnv) $ do++ sz <- A.fromIntegral integralType numType (indexHead (irArrayShape arrOut))+ tid <- threadIdx++ when (A.lt scalarType tid stride) $ do++ bid <- blockIdx+ gd <- gridDim+ c0 <- A.add numType start bid+ imapFromStepTo c0 gd end $ \chunk -> do++ (inf,sup) <- case dir of+ L -> do+ a <- A.add numType chunk (lift 1)+ b <- A.mul numType stride a+ c <- A.add numType b (lift 1)+ d <- A.add numType c stride+ e <- A.min scalarType d sz+ return (c,e)+ R -> do+ a <- A.sub numType end chunk+ b <- A.mul numType stride a+ c <- A.sub numType sz b+ d <- A.sub numType c (lift 1)+ e <- A.sub numType d stride+ f <- A.max scalarType e (lift 0)+ return (f,d)++ carry <- case dir of+ L -> readArray arrTmp chunk+ R -> do+ a <- A.add numType chunk (lift 1)+ b <- readArray arrTmp a+ return b++ -- Apply the carry-in value to each element in the chunk+ bd <- blockDim+ i0 <- A.add numType inf tid+ imapFromStepTo i0 bd sup $ \i -> do+ v <- readArray arrOut i+ u <- case dir of+ L -> app2 combine carry v+ R -> app2 combine v carry+ writeArray arrOut i u++ return_+++-- Multidimensional scans+-- ----------------------++-- Multidimensional scan along the innermost dimension+--+-- A thread block individually computes along each innermost dimension. This is+-- a single-pass operation.+--+-- * We can assume that the array is non-empty; exclusive scans with empty+-- innermost dimension will be instead filled with the seed element via+-- 'mkScanFill'.+--+-- * Small but non-empty innermost dimension arrays (size << thread+-- block size) will have many threads which do no work.+--+mkScanDim+ :: forall aenv sh e. (Shape sh, Elt e)+ => Direction+ -> DeviceProperties -- ^ properties of the target GPU+ -> Gamma aenv -- ^ array environment+ -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function+ -> Maybe (IRExp PTX aenv e) -- ^ seed element, if this is an exclusive scan+ -> IRDelayed PTX aenv (Array (sh:.Int) e) -- ^ input data+ -> CodeGen (IROpenAcc PTX aenv (Array (sh:.Int) e))+mkScanDim dir dev aenv combine mseed IRDelayed{..} =+ let+ (start, end, paramGang) = gangParam+ (arrOut, paramOut) = mutableArray ("out" :: Name (Array (sh:.Int) e))+ paramEnv = envParam aenv+ --+ config = launchConfig dev (CUDA.incWarp dev) smem const+ smem n = warps * (1 + per_warp) * bytes+ where+ ws = CUDA.warpSize dev+ warps = n `div` ws+ per_warp = ws + ws `div` 2+ bytes = sizeOf (eltType (undefined :: e))+ in+ makeOpenAccWith config "scan" (paramGang ++ paramOut ++ paramEnv) $ do++ -- The first and last threads of the block need to communicate the+ -- block-wide aggregate as a carry-in value across iterations.+ --+ -- TODO: we could optimise this a bit if we can get access to the shared+ -- memory area used by 'scanBlockSMem', and from there directly read the+ -- value computed by the last thread.+ carry <- staticSharedMem 1++ -- Size of the input array+ sz <- A.fromIntegral integralType numType . indexHead =<< delayedExtent++ -- Thread blocks iterate over the outer dimensions. Threads in a block+ -- cooperatively scan along one dimension, but thread blocks do not+ -- communicate with each other.+ --+ bid <- blockIdx+ gd <- gridDim+ s0 <- A.add numType start bid+ imapFromStepTo s0 gd end $ \seg -> do++ -- Index this thread reads from+ tid <- threadIdx+ i0 <- case dir of+ L -> do x <- A.mul numType seg sz+ y <- A.add numType x tid+ return y++ R -> do x <- A.add numType seg (lift 1)+ y <- A.mul numType x sz+ z <- A.sub numType y tid+ w <- A.sub numType z (lift 1)+ return w++ -- Index this thread writes to+ j0 <- case mseed of+ Nothing -> return i0+ Just{} -> do szp1 <- A.fromIntegral integralType numType (indexHead (irArrayShape arrOut))+ case dir of+ L -> do x <- A.mul numType seg szp1+ y <- A.add numType x tid+ return y++ R -> do x <- A.add numType seg (lift 1)+ y <- A.mul numType x szp1+ z <- A.sub numType y tid+ w <- A.sub numType z (lift 1)+ return w++ -- Stride indices by block dimension+ bd <- blockDim+ let next ix = case dir of+ L -> A.add numType ix bd+ R -> A.sub numType ix bd++ -- Initialise this scan segment+ --+ -- If this is an exclusive scan then the first thread just evaluates the+ -- seed element and stores this value into the carry-in slot. All threads+ -- shift their write-to index (j) by one, to make space for this element.+ --+ -- If this is an inclusive scan then do a block-wide scan. The last thread+ -- in the block writes the carry-in value.+ --+ r <-+ case mseed of+ Just seed -> do+ when (A.eq scalarType tid (lift 0)) $ do+ z <- seed+ writeArray arrOut j0 z+ writeArray carry (lift 0 :: IR Int32) z+ j1 <- case dir of+ L -> A.add numType j0 (lift 1)+ R -> A.sub numType j0 (lift 1)+ return $ A.trip sz i0 j1++ Nothing -> do+ when (A.lt scalarType tid sz) $ do+ x0 <- app1 delayedLinearIndex =<< A.fromIntegral integralType numType i0+ r0 <- if A.gte scalarType sz bd+ then scanBlockSMem dir dev combine Nothing x0+ else scanBlockSMem dir dev combine (Just sz) x0+ writeArray arrOut j0 r0++ ll <- A.sub numType bd (lift 1)+ when (A.eq scalarType tid ll) $+ writeArray carry (lift 0 :: IR Int32) r0++ n1 <- A.sub numType sz bd+ i1 <- next i0+ j1 <- next j0+ return $ A.trip n1 i1 j1++ -- Iterate over the remaining elements in this segment+ void $ while+ (\(A.fst3 -> n) -> A.gt scalarType n (lift 0))+ (\(A.untrip -> (n,i,j)) -> do++ -- Wait for the carry-in value from the previous iteration to be updated+ __syncthreads++ -- Compute and store the next element of the scan+ --+ -- NOTE: As with 'foldSeg' we require all threads to participate in+ -- every iteration of the loop otherwise they will die prematurely.+ -- Out-of-bounds threads return 'undef' at this point, which is really+ -- unfortunate ):+ --+ x <- if A.lt scalarType tid n+ then app1 delayedLinearIndex =<< A.fromIntegral integralType numType i+ else let+ go :: TupleType a -> Operands a+ go UnitTuple = OP_Unit+ go (PairTuple a b) = OP_Pair (go a) (go b)+ go (SingleTuple t) = ir' t (undef t)+ in+ return . IR $ go (eltType (undefined::e))++ -- Thread zero incorporates the carry-in element+ y <- if A.eq scalarType tid (lift 0)+ then do+ c <- readArray carry (lift 0 :: IR Int32)+ case dir of+ L -> app2 combine c x+ R -> app2 combine x c+ else+ return x++ -- Perform the scan and write the result to memory+ z <- if A.gte scalarType n bd+ then scanBlockSMem dir dev combine Nothing y+ else scanBlockSMem dir dev combine (Just n) y++ when (A.lt scalarType tid n) $ do+ writeArray arrOut j z++ -- The last thread of the block writes its result as the carry-out+ -- value. If this thread is not active then we are on the last+ -- iteration of the loop and it will not be needed.+ w <- A.sub numType bd (lift 1)+ when (A.eq scalarType tid w) $+ writeArray carry (lift 0 :: IR Int32) z++ -- Update indices for the next iteration+ n' <- A.sub numType n bd+ i' <- next i+ j' <- next j+ return $ A.trip n' i' j')+ r++ return_+++-- Multidimensional scan' along the innermost dimension+--+-- A thread block individually computes along each innermost dimension. This is+-- a single-pass operation.+--+-- * We can assume that the array is non-empty; exclusive scans with empty+-- innermost dimension will be instead filled with the seed element via+-- 'mkScan'Fill'.+--+-- * Small but non-empty innermost dimension arrays (size << thread+-- block size) will have many threads which do no work.+--+mkScan'Dim+ :: forall aenv sh e. (Shape sh, Elt e)+ => Direction+ -> DeviceProperties -- ^ properties of the target GPU+ -> Gamma aenv -- ^ array environment+ -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function+ -> IRExp PTX aenv e -- ^ seed element+ -> IRDelayed PTX aenv (Array (sh:.Int) e) -- ^ input data+ -> CodeGen (IROpenAcc PTX aenv (Array (sh:.Int) e, Array sh e))+mkScan'Dim dir dev aenv combine seed IRDelayed{..} =+ let+ (start, end, paramGang) = gangParam+ (arrOut, paramOut) = mutableArray ("out" :: Name (Array (sh:.Int) e))+ (arrSum, paramSum) = mutableArray ("sum" :: Name (Array sh e))+ paramEnv = envParam aenv+ --+ config = launchConfig dev (CUDA.incWarp dev) smem const+ smem n = warps * (1 + per_warp) * bytes+ where+ ws = CUDA.warpSize dev+ warps = n `div` ws+ per_warp = ws + ws `div` 2+ bytes = sizeOf (eltType (undefined :: e))+ in+ makeOpenAccWith config "scan" (paramGang ++ paramOut ++ paramSum ++ paramEnv) $ do++ -- The first and last threads of the block need to communicate the+ -- block-wide aggregate as a carry-in value across iterations.+ --+ -- TODO: we could optimise this a bit if we can get access to the shared+ -- memory area used by 'scanBlockSMem', and from there directly read the+ -- value computed by the last thread.+ carry <- staticSharedMem 1++ -- Size of the input array+ sz <- A.fromIntegral integralType numType . indexHead =<< delayedExtent++ -- If the innermost dimension is smaller than the number of threads in the+ -- block, those threads will never contribute to the output.+ tid <- threadIdx+ when (A.lte scalarType tid sz) $ do++ -- Thread blocks iterate over the outer dimensions, each thread block+ -- cooperatively scanning along each outermost index.+ bid <- blockIdx+ gd <- gridDim+ s0 <- A.add numType start bid+ imapFromStepTo s0 gd end $ \seg -> do++ -- Not necessary to wait for threads to catch up before starting this segment+ -- __syncthreads++ -- Linear index bounds for this segment+ inf <- A.mul numType seg sz+ sup <- A.add numType inf sz++ -- Index that this thread will read from. Recall that the supremum index+ -- is exclusive.+ i0 <- case dir of+ L -> A.add numType inf tid+ R -> do x <- A.sub numType sup tid+ y <- A.sub numType x (lift 1)+ return y++ -- The index that this thread will write to. This is just shifted along+ -- by one to make room for the initial element.+ j0 <- case dir of+ L -> A.add numType i0 (lift 1)+ R -> A.sub numType i0 (lift 1)++ -- Evaluate the initial element. Store it into the carry-in slot as well+ -- as to the array as the first element. This is always valid because if+ -- the input array is empty then we will be evaluating via mkScan'Fill.+ when (A.eq scalarType tid (lift 0)) $ do+ z <- seed+ writeArray arrOut i0 z+ writeArray carry (lift 0 :: IR Int32) z++ bd <- blockDim+ let next ix = case dir of+ L -> A.add numType ix bd+ R -> A.sub numType ix bd++ -- Now, threads iterate over the elements along the innermost dimension.+ -- At each iteration the first thread incorporates the carry-in value+ -- from the previous step.+ --+ -- The index tracks how many elements remain for the thread block, since+ -- indices i* and j* are local to each thread+ n0 <- A.sub numType sup inf+ void $ while+ (\(A.fst3 -> n) -> A.gt scalarType n (lift 0))+ (\(A.untrip -> (n,i,j)) -> do++ -- Wait for threads to catch up to ensure the carry-in value from+ -- the last iteration has been updated+ __syncthreads++ -- If all threads in the block will participate this round we can+ -- avoid (almost) all bounds checks.+ _ <- if A.gte scalarType n bd+ -- All threads participate. No bounds checks required but+ -- the last thread needs to update the carry-in value.+ then do+ x <- app1 delayedLinearIndex =<< A.fromIntegral integralType numType i+ y <- if A.eq scalarType tid (lift 0)+ then do+ c <- readArray carry (lift 0 :: IR Int32)+ case dir of+ L -> app2 combine c x+ R -> app2 combine x c+ else+ return x+ z <- scanBlockSMem dir dev combine Nothing y++ -- Write results to the output array. Note that if we+ -- align directly on the boundary of the array this is not+ -- valid for the last thread.+ case dir of+ L -> when (A.lt scalarType j sup) $ writeArray arrOut j z+ R -> when (A.gte scalarType j inf) $ writeArray arrOut j z++ -- Last thread of the block also saves its result as the+ -- carry-in value+ bd1 <- A.sub numType bd (lift 1)+ when (A.eq scalarType tid bd1) $+ writeArray carry (lift 0 :: IR Int32) z++ return (IR OP_Unit :: IR ())++ -- Only threads that are in bounds can participate. This is+ -- the last iteration of the loop. The last active thread+ -- still needs to store its value into the carry-in slot.+ else do+ when (A.lt scalarType tid n) $ do+ x <- app1 delayedLinearIndex =<< A.fromIntegral integralType numType i+ y <- if A.eq scalarType tid (lift 0)+ then do+ c <- readArray carry (lift 0 :: IR Int32)+ case dir of+ L -> app2 combine c x+ R -> app2 combine x c+ else+ return x+ z <- scanBlockSMem dir dev combine (Just n) y++ m <- A.sub numType n (lift 1)+ _ <- if A.lt scalarType tid m+ then writeArray arrOut j z >> return (IR OP_Unit :: IR ())+ else writeArray carry (lift 0 :: IR Int32) z >> return (IR OP_Unit :: IR ())++ return ()+ return (IR OP_Unit :: IR ())++ A.trip <$> A.sub numType n bd <*> next i <*> next j)+ (A.trip n0 i0 j0)++ -- Wait for the carry-in value to be updated+ __syncthreads++ -- Store the carry-in value to the separate final results array+ when (A.eq scalarType tid (lift 0)) $+ writeArray arrSum seg =<< readArray carry (lift 0 :: IR Int32)++ return_++++-- Parallel scan, auxiliary+--+-- If this is an exclusive scan of an empty array, we just fill the result with+-- the seed element.+--+mkScanFill+ :: (Shape sh, Elt e)+ => PTX+ -> Gamma aenv+ -> IRExp PTX aenv e+ -> CodeGen (IROpenAcc PTX aenv (Array sh e))+mkScanFill ptx aenv seed =+ mkGenerate ptx aenv (IRFun1 (const seed))++mkScan'Fill+ :: forall aenv sh e. (Shape sh, Elt e)+ => PTX+ -> Gamma aenv+ -> IRExp PTX aenv e+ -> CodeGen (IROpenAcc PTX aenv (Array (sh:.Int) e, Array sh e))+mkScan'Fill ptx aenv seed =+ Safe.coerce <$> (mkGenerate ptx aenv (IRFun1 (const seed)) :: CodeGen (IROpenAcc PTX aenv (Array sh e)))+++-- Block wide scan+-- ---------------++-- Efficient block-wide (inclusive) scan using the specified operator.+--+-- Each block requires (#warps * (1 + 1.5*warp size)) elements of dynamically+-- allocated shared memory.+--+-- Example: https://github.com/NVlabs/cub/blob/1.5.4/cub/block/specializations/block_scan_warp_scans.cuh+--+scanBlockSMem+ :: forall aenv e. Elt e+ => Direction+ -> DeviceProperties -- ^ properties of the target device+ -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function+ -> Maybe (IR Int32) -- ^ number of valid elements (may be less than block size)+ -> IR e -- ^ calling thread's input element+ -> CodeGen (IR e)+scanBlockSMem dir dev combine nelem = warpScan >=> warpPrefix+ where+ 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_bytes = warp_smem_elems * sizeOf (eltType (undefined::e))++ -- Step 1: Scan in every warp+ warpScan :: IR e -> CodeGen (IR e)+ warpScan input = do+ -- Allocate (1.5 * warpSize) elements of shared memory for each warp+ -- (individually addressable by each warp)+ wid <- warpId+ skip <- A.mul numType wid (int32 warp_smem_bytes)+ smem <- dynamicSharedMem (int32 warp_smem_elems) skip+ scanWarpSMem dir dev combine smem input++ -- Step 2: Collect the aggregate results of each warp to compute the prefix+ -- values for each warp and combine with the partial result to compute each+ -- thread's final value.+ warpPrefix :: IR e -> CodeGen (IR e)+ warpPrefix input = do+ -- Allocate #warps elements of shared memory+ bd <- blockDim+ warps <- A.quot integralType bd (int32 (CUDA.warpSize dev))+ skip <- A.mul numType warps (int32 warp_smem_bytes)+ smem <- dynamicSharedMem warps skip++ -- Share warp aggregates+ wid <- warpId+ lane <- laneId+ when (A.eq scalarType lane (int32 (CUDA.warpSize dev - 1))) $ do+ writeArray smem wid input++ -- Wait for each warp to finish its local scan and share the aggregate+ __syncthreads++ -- Compute the prefix value for this warp and add to the partial result.+ -- This step is not required for the first warp, which has no carry-in.+ if A.eq scalarType wid (lift 0)+ then return input+ else do+ -- Every thread sequentially scans the warp aggregates to compute+ -- their prefix value. We do this sequentially, but could also have+ -- warp 0 do it cooperatively if we limit thread block sizes to+ -- (warp size ^ 2).+ steps <- case nelem of+ Nothing -> return wid+ Just n -> A.min scalarType wid =<< A.quot integralType n (int32 (CUDA.warpSize dev))++ p0 <- readArray smem (lift 0 :: IR Int32)+ prefix <- iterFromStepTo (lift 1) (lift 1) steps p0 $ \step x -> do+ y <- readArray smem step+ case dir of+ L -> app2 combine x y+ R -> app2 combine y x++ case dir of+ L -> app2 combine prefix input+ R -> app2 combine input prefix+++-- Warp-wide scan+-- --------------++-- Efficient warp-wide (inclusive) scan using the specified operator.+--+-- Each warp requires 48 (1.5 x warp size) elements of shared memory. The+-- routine assumes that it is allocated individually per-warp (i.e. can be+-- indexed in the range [0, warp size)).+--+-- Example: https://github.com/NVlabs/cub/blob/1.5.4/cub/warp/specializations/warp_scan_smem.cuh+--+scanWarpSMem+ :: forall aenv e. Elt e+ => Direction+ -> DeviceProperties -- ^ properties of the target device+ -> IRFun2 PTX aenv (e -> e -> e) -- ^ combination function+ -> IRArray (Vector e) -- ^ temporary storage array in shared memory (1.5 x warp size elements)+ -> IR e -- ^ calling thread's input element+ -> CodeGen (IR e)+scanWarpSMem dir dev combine smem = scan 0+ where+ log2 :: Double -> Double+ log2 = P.logBase 2++ -- Number of steps required to scan warp+ steps = P.floor (log2 (P.fromIntegral (CUDA.warpSize dev)))+ halfWarp = P.fromIntegral (CUDA.warpSize dev `div` 2)++ -- Unfold the scan as a recursive code generation function+ scan :: Int -> IR e -> CodeGen (IR e)+ scan step x+ | step >= steps = return x+ | offset <- 1 `P.shiftL` step = do+ -- share partial result through shared memory buffer+ lane <- laneId+ i <- A.add numType lane (lift halfWarp)+ writeArray smem i x++ -- update partial result if in range+ x' <- if A.gte scalarType lane (lift offset)+ then do+ i' <- A.sub numType i (lift offset) -- lane + HALF_WARP - offset+ x' <- readArray smem i'+ case dir of+ L -> app2 combine x' x+ R -> app2 combine x x'++ else+ return x++ scan (step+1) x'+
+ Data/Array/Accelerate/LLVM/PTX/Compile.hs view
@@ -0,0 +1,289 @@+{-# LANGUAGE CPP #-}+{-# LANGUAGE OverloadedStrings #-}+{-# LANGUAGE RecordWildCards #-}+{-# LANGUAGE TemplateHaskell #-}+{-# LANGUAGE TypeFamilies #-}+{-# OPTIONS_GHC -fno-warn-orphans #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Compile+-- Copyright : [2014..2017] Trevor L. McDonell+-- [2014..2014] Vinod Grover (NVIDIA Corporation)+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Compile (++ module Data.Array.Accelerate.LLVM.Compile,+ ExecutableR(..), Kernel(..),++) 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++-- accelerate+import Data.Array.Accelerate.Error ( internalError )+import Data.Array.Accelerate.Lifetime+import Data.Array.Accelerate.Trafo ( DelayedOpenAcc )++import Data.Array.Accelerate.LLVM.CodeGen+import Data.Array.Accelerate.LLVM.CodeGen.Environment ( Gamma )+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.Foreign ( )+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+#ifdef ACCELERATE_USE_NVVM+import qualified Foreign.NVVM as NVVM+#endif++-- standard library+import Control.Monad.Except+import Control.Monad.State+import Data.ByteString ( ByteString )+import Data.List ( intercalate )+import Text.Printf ( printf )+import qualified Data.ByteString.Char8 as B+import qualified Data.Map as Map+import Prelude as P+++instance Compile PTX where+ data ExecutableR PTX = PTXR { ptxKernel :: ![Kernel]+ , ptxModule :: {-# UNPACK #-} !(Lifetime CUDA.Module)+ }+ compileForTarget = compileForPTX+++data Kernel = Kernel {+ kernelFun :: {-# UNPACK #-} !CUDA.Fun+ , kernelOccupancy :: {-# UNPACK #-} !CUDA.Occupancy+ , kernelSharedMemBytes :: {-# UNPACK #-} !Int+ , kernelThreadBlockSize :: {-# UNPACK #-} !Int+ , kernelThreadBlocks :: (Int -> Int)+ , kernelName :: String+ }++-- | 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.+--+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+ --+ 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'+++-- | Compile the LLVM module to produce a CUDA module.+--+-- * If we are using NVVM, this includes all LLVM optimisations plus some+-- sekrit optimisations.+--+-- * If we are just using the llvm ptx backend, we still need to run the+-- standard optimisations.+--+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+++#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.+--+compileModuleNVVM :: CUDA.DeviceProperties -> String -> [(String, ByteString)] -> LLVM.Module -> IO CUDA.Module+compileModuleNVVM dev name libdevice mdl = do+ _debug <- Debug.queryFlag Debug.debug_cc+ --+ let arch = CUDA.computeCapability dev+ verbose = if _debug then [ NVVM.GenerateDebugInfo ] else []+ flags = NVVM.Target arch : verbose++ -- Note: [NVVM and target datalayout]+ --+ -- The NVVM library does not correctly parse the target datalayout field,+ -- instead doing a (very dodgy) string compare against exactly two+ -- expected values. This means that it is sensitive to, e.g. the ordering+ -- of the fields, and changes to the representation in each LLVM release.+ --+ -- We get around this by only specifying the data layout in a separate+ -- (otherwise empty) module that we additionally link against.+ --+ header = case bitSize (undefined::Int) of+ 32 -> "target triple = \"nvptx-nvidia-cuda\"\ntarget datalayout = \"e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64\""+ 64 -> "target triple = \"nvptx64-nvidia-cuda\"\ntarget datalayout = \"e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64\""+ _ -> $internalError "compileModuleNVVM" "I don't know what architecture I am"++ 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++ -- Lower the generated module to bitcode, then compile and link together with+ -- the shim header and libdevice library (if necessary)+ bc <- LLVM.moduleBitcode mdl+ 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)++ -- Link into a new CUDA module in the current context+ linkPTX name (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 =+ 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++ LLVM.withPassManager pss $ \pm -> do+#ifdef ACCELERATE_INTERNAL_CHECKS+ runError $ 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++ -- 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+++{--+-- | Extract the names of the function definitions from the module.+--+-- Note: [Extracting global function names]+--+-- It is important to run this on the module given to us by code generation.+-- After combining modules with 'libdevice', extra function definitions,+-- corresponding to basic maths operations, will be added to the module. These+-- functions will not be callable as __global__ functions.+--+-- The list of names will be exported in the order that they appear in the+-- module.+--+globalFunctions :: [Definition] -> [String]+globalFunctions defs =+ [ n | GlobalDefinition Function{..} <- defs+ , not (null basicBlocks)+ , let Name n = name+ ]+--}+
+ Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice.hs view
@@ -0,0 +1,575 @@+{-# LANGUAGE FlexibleInstances #-}+{-# LANGUAGE OverloadedStrings #-}+{-# LANGUAGE RecordWildCards #-}+{-# LANGUAGE TemplateHaskell #-}+{-# LANGUAGE TupleSections #-}+{-# OPTIONS_GHC -fno-warn-orphans #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice+-- Copyright : [2014..2017] Trevor L. McDonell+-- [2014..2014] Vinod Grover (NVIDIA Corporation)+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice (++ nvvmReflect, libdevice,++) where++-- llvm-hs+import LLVM.Context+import LLVM.Module as LLVM+import LLVM.AST as AST ( Module(..), Definition(..) )+import LLVM.AST.Attribute+import LLVM.AST.Global as G+import qualified LLVM.AST.Name as AST++-- 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++-- cuda+import Foreign.CUDA.Analysis++-- standard library+import Control.Monad.Except+import Data.ByteString ( ByteString )+import Data.HashMap.Strict ( HashMap )+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+++-- 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.+--+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:+--+-- 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. The top-level+-- unsafePerformIO ensures that the data is read only once per program+-- execution.+--+{-# 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:+--+-- 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++ -- 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)+++-- 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 -> 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"++ path <- libdevicePath+ files <- getDirectoryContents path+ name <- maybe err return . listToMaybe . sortBy (flip compare) $ filter best files+ bc <- B.readFile (path </> name)++ return (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))+++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 String Label+libdeviceIndex =+ let nv base = (base, Label $ "__nv_" ++ base)+ in+ HashMap.fromList $ map nv+ [ "abs"+ , "acos"+ , "acosf"+ , "acosh"+ , "acoshf"+ , "asin"+ , "asinf"+ , "asinh"+ , "asinhf"+ , "atan"+ , "atan2"+ , "atan2f"+ , "atanf"+ , "atanh"+ , "atanhf"+ , "brev"+ , "brevll"+ , "byte_perm"+ , "cbrt"+ , "cbrtf"+ , "ceil"+ , "ceilf"+ , "clz"+ , "clzll"+ , "copysign"+ , "copysignf"+ , "cos"+ , "cosf"+ , "cosh"+ , "coshf"+ , "cospi"+ , "cospif"+ , "dadd_rd"+ , "dadd_rn"+ , "dadd_ru"+ , "dadd_rz"+ , "ddiv_rd"+ , "ddiv_rn"+ , "ddiv_ru"+ , "ddiv_rz"+ , "dmul_rd"+ , "dmul_rn"+ , "dmul_ru"+ , "dmul_rz"+ , "double2float_rd"+ , "double2float_rn"+ , "double2float_ru"+ , "double2float_rz"+ , "double2hiint"+ , "double2int_rd"+ , "double2int_rn"+ , "double2int_ru"+ , "double2int_rz"+ , "double2ll_rd"+ , "double2ll_rn"+ , "double2ll_ru"+ , "double2ll_rz"+ , "double2loint"+ , "double2uint_rd"+ , "double2uint_rn"+ , "double2uint_ru"+ , "double2uint_rz"+ , "double2ull_rd"+ , "double2ull_rn"+ , "double2ull_ru"+ , "double2ull_rz"+ , "double_as_longlong"+ , "drcp_rd"+ , "drcp_rn"+ , "drcp_ru"+ , "drcp_rz"+ , "dsqrt_rd"+ , "dsqrt_rn"+ , "dsqrt_ru"+ , "dsqrt_rz"+ , "erf"+ , "erfc"+ , "erfcf"+ , "erfcinv"+ , "erfcinvf"+ , "erfcx"+ , "erfcxf"+ , "erff"+ , "erfinv"+ , "erfinvf"+ , "exp"+ , "exp10"+ , "exp10f"+ , "exp2"+ , "exp2f"+ , "expf"+ , "expm1"+ , "expm1f"+ , "fabs"+ , "fabsf"+ , "fadd_rd"+ , "fadd_rn"+ , "fadd_ru"+ , "fadd_rz"+ , "fast_cosf"+ , "fast_exp10f"+ , "fast_expf"+ , "fast_fdividef"+ , "fast_log10f"+ , "fast_log2f"+ , "fast_logf"+ , "fast_powf"+ , "fast_sincosf"+ , "fast_sinf"+ , "fast_tanf"+ , "fdim"+ , "fdimf"+ , "fdiv_rd"+ , "fdiv_rn"+ , "fdiv_ru"+ , "fdiv_rz"+ , "ffs"+ , "ffsll"+ , "finitef"+ , "float2half_rn"+ , "float2int_rd"+ , "float2int_rn"+ , "float2int_ru"+ , "float2int_rz"+ , "float2ll_rd"+ , "float2ll_rn"+ , "float2ll_ru"+ , "float2ll_rz"+ , "float2uint_rd"+ , "float2uint_rn"+ , "float2uint_ru"+ , "float2uint_rz"+ , "float2ull_rd"+ , "float2ull_rn"+ , "float2ull_ru"+ , "float2ull_rz"+ , "float_as_int"+ , "floor"+ , "floorf"+ , "fma"+ , "fma_rd"+ , "fma_rn"+ , "fma_ru"+ , "fma_rz"+ , "fmaf"+ , "fmaf_rd"+ , "fmaf_rn"+ , "fmaf_ru"+ , "fmaf_rz"+ , "fmax"+ , "fmaxf"+ , "fmin"+ , "fminf"+ , "fmod"+ , "fmodf"+ , "fmul_rd"+ , "fmul_rn"+ , "fmul_ru"+ , "fmul_rz"+ , "frcp_rd"+ , "frcp_rn"+ , "frcp_ru"+ , "frcp_rz"+ , "frexp"+ , "frexpf"+ , "frsqrt_rn"+ , "fsqrt_rd"+ , "fsqrt_rn"+ , "fsqrt_ru"+ , "fsqrt_rz"+ , "fsub_rd"+ , "fsub_rn"+ , "fsub_ru"+ , "fsub_rz"+ , "hadd"+ , "half2float"+ , "hiloint2double"+ , "hypot"+ , "hypotf"+ , "ilogb"+ , "ilogbf"+ , "int2double_rn"+ , "int2float_rd"+ , "int2float_rn"+ , "int2float_ru"+ , "int2float_rz"+ , "int_as_float"+ , "isfinited"+ , "isinfd"+ , "isinff"+ , "isnand"+ , "isnanf"+ , "j0"+ , "j0f"+ , "j1"+ , "j1f"+ , "jn"+ , "jnf"+ , "ldexp"+ , "ldexpf"+ , "lgamma"+ , "lgammaf"+ , "ll2double_rd"+ , "ll2double_rn"+ , "ll2double_ru"+ , "ll2double_rz"+ , "ll2float_rd"+ , "ll2float_rn"+ , "ll2float_ru"+ , "ll2float_rz"+ , "llabs"+ , "llmax"+ , "llmin"+ , "llrint"+ , "llrintf"+ , "llround"+ , "llroundf"+ , "log"+ , "log10"+ , "log10f"+ , "log1p"+ , "log1pf"+ , "log2"+ , "log2f"+ , "logb"+ , "logbf"+ , "logf"+ , "longlong_as_double"+ , "max"+ , "min"+ , "modf"+ , "modff"+ , "mul24"+ , "mul64hi"+ , "mulhi"+ , "nan"+ , "nanf"+ , "nearbyint"+ , "nearbyintf"+ , "nextafter"+ , "nextafterf"+ , "normcdf"+ , "normcdff"+ , "normcdfinv"+ , "normcdfinvf"+ , "popc"+ , "popcll"+ , "pow"+ , "powf"+ , "powi"+ , "powif"+ , "rcbrt"+ , "rcbrtf"+ , "remainder"+ , "remainderf"+ , "remquo"+ , "remquof"+ , "rhadd"+ , "rint"+ , "rintf"+ , "round"+ , "roundf"+ , "rsqrt"+ , "rsqrtf"+ , "sad"+ , "saturatef"+ , "scalbn"+ , "scalbnf"+ , "signbitd"+ , "signbitf"+ , "sin"+ , "sincos"+ , "sincosf"+ , "sincospi"+ , "sincospif"+ , "sinf"+ , "sinh"+ , "sinhf"+ , "sinpi"+ , "sinpif"+ , "sqrt"+ , "sqrtf"+ , "tan"+ , "tanf"+ , "tanh"+ , "tanhf"+ , "tgamma"+ , "tgammaf"+ , "trunc"+ , "truncf"+ , "uhadd"+ , "uint2double_rn"+ , "uint2float_rd"+ , "uint2float_rn"+ , "uint2float_ru"+ , "uint2float_rz"+ , "ull2double_rd"+ , "ull2double_rn"+ , "ull2double_ru"+ , "ull2double_rz"+ , "ull2float_rd"+ , "ull2float_rn"+ , "ull2float_ru"+ , "ull2float_rz"+ , "ullmax"+ , "ullmin"+ , "umax"+ , "umin"+ , "umul24"+ , "umul64hi"+ , "umulhi"+ , "urhadd"+ , "usad"+ , "y0"+ , "y0f"+ , "y1"+ , "y1f"+ , "yn"+ , "ynf"+ ]+
+ Data/Array/Accelerate/LLVM/PTX/Compile/Link.hs view
@@ -0,0 +1,183 @@+{-# LANGUAGE OverloadedStrings #-}+{-# LANGUAGE RecordWildCards #-}+{-# LANGUAGE TemplateHaskell #-}+{-# LANGUAGE ViewPatterns #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Compile.Link+-- Copyright : [2014..2017] Trevor L. McDonell+-- [2014..2014] Vinod Grover (NVIDIA Corporation)+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Compile.Link (++ withLibdeviceNVVM,+ withLibdeviceNVPTX,++) where++-- llvm-hs+import LLVM.Context+import qualified LLVM.Module as LLVM++import LLVM.AST as AST+import LLVM.AST.Global as G+import LLVM.AST.Linkage++-- accelerate+import Data.Array.Accelerate.Error++import Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice+import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug++-- cuda+import Foreign.CUDA.Analysis++-- standard library+import Control.Monad.Except+import Data.ByteString ( ByteString )+import Data.HashSet ( HashSet )+import Data.List+import Data.Maybe+import Text.Printf+import qualified Data.HashSet as Set+++-- | Lower an LLVM AST to C++ objects and link it against the libdevice module,+-- iff any libdevice functions are referenced from the base module.+--+-- Note: [Linking with libdevice]+--+-- The CUDA toolkit comes with an LLVM bitcode library called 'libdevice' that+-- implements many common mathematical functions. The library can be used as a+-- high performance math library for targets of the LLVM NVPTX backend, such as+-- this one. To link a module 'foo' with libdevice, the following compilation+-- pipeline is recommended:+--+-- 1. Save all external functions in module 'foo'+--+-- 2. Link module 'foo' with the appropriate 'libdevice_compute_XX.YY.bc'+--+-- 3. Internalise all functions not in the list from (1)+--+-- 4. Eliminate all unused internal functions+--+-- 5. Run the NVVMReflect pass (see note: [NVVM Reflect Pass])+--+-- 6. Run the standard optimisation pipeline+--+withLibdeviceNVPTX+ :: DeviceProperties+ -> Context+ -> Module+ -> (LLVM.Module -> IO a)+ -> IO a+withLibdeviceNVPTX dev ctx ast next =+ case Set.null externs of+ True -> runError $ LLVM.withModuleFromAST ctx ast next+ False ->+ runError $ LLVM.withModuleFromAST ctx ast $ \mdl ->+ runError $ LLVM.withModuleFromAST ctx nvvmReflect $ \refl ->+ runError $ LLVM.withModuleFromAST ctx (internalise externs libdev) $ \libd -> do+ runError $ linkModules mdl refl+ runError $ linkModules mdl libd+ Debug.traceIO Debug.dump_cc msg+ next mdl+ where+ -- Replace the target triple and datalayout from the libdevice.bc module+ -- with those of the generated code. This avoids warnings such as "linking+ -- two modules of different target triples..."+ libdev = (libdevice arch) { moduleTargetTriple = moduleTargetTriple ast+ , moduleDataLayout = moduleDataLayout ast+ }+ externs = analyse ast+ arch = computeCapability dev+ runError = either ($internalError "withLibdeviceNVPTX") return <=< runExceptT++ msg = printf "cc: linking with libdevice: %s"+ $ intercalate ", " (Set.toList externs)+++-- | Link LLVM modules by copying parts of the second argument into the first.+--+linkModules+ :: LLVM.Module -- module into which to link (destination: contains all symbols)+ -> LLVM.Module -- module to copy into the other (this is destroyed in the process)+ -> ExceptT String IO ()+linkModules = LLVM.linkModules+++-- | Lower an LLVM AST to C++ objects and prepare it for linking against+-- libdevice using the nvvm bindings, iff any libdevice functions are referenced+-- from the base module.+--+-- Rather than internalise and strip any unused functions ourselves, allow the+-- nvvm library to do so when linking the two modules together.+--+-- TLM: This really should work with the above method, however for some reason+-- we get a "CUDA Exception: function named symbol not found" error, even though+-- the function is clearly visible in the generated code. hmm...+--+withLibdeviceNVVM+ :: DeviceProperties+ -> Context+ -> Module+ -> ([(String, ByteString)] -> LLVM.Module -> IO a)+ -> IO a+withLibdeviceNVVM dev ctx ast next =+ runError $ LLVM.withModuleFromAST ctx ast $ \mdl -> do+ when withlib $ Debug.traceIO Debug.dump_cc msg+ next lib mdl+ where+ externs = analyse ast+ withlib = not (Set.null externs)+ lib | withlib = [ nvvmReflect, libdevice arch ]+ | otherwise = []++ arch = computeCapability dev+ runError = either ($internalError "withLibdeviceNVPTX") return <=< runExceptT++ msg = printf "cc: linking with libdevice: %s"+ $ intercalate ", " (Set.toList externs)+++-- | Analyse the LLVM AST module and determine if any of the external+-- declarations are intrinsics implemented by libdevice. The set of such+-- functions is returned, and will be used when determining which functions from+-- libdevice to internalise.+--+analyse :: Module -> HashSet String+analyse Module{..} =+ let intrinsic (GlobalDefinition Function{..})+ | null basicBlocks+ , Name n <- name+ , "__nv_" <- take 5 n+ = Just n++ intrinsic _+ = Nothing+ in+ Set.fromList (mapMaybe intrinsic moduleDefinitions)+++-- | Mark all definitions in the module as internal linkage. This means that+-- unused definitions can be removed as dead code. Be careful to leave any+-- declarations as external.+--+internalise :: HashSet String -> Module -> Module+internalise externals Module{..} =+ let internal (GlobalDefinition Function{..})+ | Name n <- name+ , not (Set.member n externals) -- we don't call this function directly; and+ , not (null basicBlocks) -- it is not an external declaration+ = GlobalDefinition (Function { linkage=Internal, .. })++ internal x+ = x+ in+ Module { moduleDefinitions = map internal moduleDefinitions, .. }+
+ Data/Array/Accelerate/LLVM/PTX/Context.hs view
@@ -0,0 +1,147 @@+{-# LANGUAGE RecordWildCards #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Context+-- Copyright : [2014..2017] Trevor L. McDonell+-- [2014..2014] Vinod Grover (NVIDIA Corporation)+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Context (++ Context(..),+ new, raw, withContext,++) where++import Data.Array.Accelerate.Lifetime+import Data.Array.Accelerate.LLVM.PTX.Analysis.Device+import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug++import qualified Foreign.CUDA.Analysis as CUDA+import qualified Foreign.CUDA.Driver as CUDA+import qualified Foreign.CUDA.Driver.Device as CUDA++import Control.Exception+import Control.Monad+import Text.PrettyPrint+++-- | An execution context, which is tied to a specific device and CUDA execution+-- context.+--+data Context = Context {+ deviceProperties :: {-# UNPACK #-} !CUDA.DeviceProperties -- information on hardware resources+ , deviceContext :: {-# UNPACK #-} !(Lifetime CUDA.Context) -- device execution context+ }++instance Eq Context where+ c1 == c2 = deviceContext c1 == deviceContext c2+++-- | Create a new CUDA execution context+--+new :: CUDA.Device+ -> CUDA.DeviceProperties+ -> [CUDA.ContextFlag]+ -> IO Context+new dev prp flags = do+ ctx <- raw dev prp =<< CUDA.create dev flags+ _ <- CUDA.pop+ return ctx++-- | Wrap a raw CUDA execution context+--+raw :: CUDA.Device+ -> CUDA.DeviceProperties+ -> CUDA.Context+ -> IO Context+raw dev prp ctx = do+ lft <- newLifetime ctx+ addFinalizer lft $ do+ message $ "finalise context " ++ showContext ctx+ CUDA.destroy ctx++ -- The kernels don't use much shared memory, so for devices that support it+ -- prefer using those memory banks as an L1 cache.+ --+ -- TLM: Is this a good idea? For example, external libraries such as cuBLAS+ -- rely heavily on shared memory and thus this could adversely affect+ -- performance. Perhaps we should use 'setCacheConfigFun' for individual+ -- functions which might benefit from this.+ --+ when (CUDA.computeCapability prp >= CUDA.Compute 2 0)+ (CUDA.setCache CUDA.PreferL1)++ -- Display information about the selected device+ Debug.traceIO Debug.verbose (deviceInfo dev prp)++ return $! Context prp lft+++-- | Push the context onto the CPUs thread stack of current contexts and execute+-- some operation.+--+{-# INLINE withContext #-}+withContext :: Context -> IO a -> IO a+withContext Context{..} action =+ withLifetime deviceContext $ \ctx ->+ bracket_ (push ctx) pop action++{-# INLINE push #-}+push :: CUDA.Context -> IO ()+push ctx = do+ message $ "push context: " ++ showContext ctx+ CUDA.push ctx++{-# INLINE pop #-}+pop :: IO ()+pop = do+ ctx <- CUDA.pop+ message $ "pop context: " ++ showContext ctx+++-- Debugging+-- ---------++-- Nicely format a summary of the selected CUDA device, example:+--+-- Device 0: GeForce 9600M GT (compute capability 1.1)+-- 4 multiprocessors @ 1.25GHz (32 cores), 512MB global memory+--+deviceInfo :: CUDA.Device -> CUDA.DeviceProperties -> String+deviceInfo dev prp = render $ reset <>+ devID <> colon <+> vcat [ name <+> parens compute+ , processors <+> at <+> text clock <+> parens cores <> comma <+> memory+ ]+ where+ name = text (CUDA.deviceName prp)+ compute = text "compute capatability" <+> text (show $ CUDA.computeCapability prp)+ devID = text "Device" <+> int (fromIntegral $ CUDA.useDevice dev) -- hax+ processors = int (CUDA.multiProcessorCount prp) <+> text "multiprocessors"+ cores = int (CUDA.multiProcessorCount prp * coresPerMultiProcessor prp) <+> text "cores"+ memory = text mem <+> text "global memory"+ --+ clock = Debug.showFFloatSIBase (Just 2) 1000 (fromIntegral $ CUDA.clockRate prp * 1000 :: Double) "Hz"+ mem = Debug.showFFloatSIBase (Just 0) 1024 (fromIntegral $ CUDA.totalGlobalMem prp :: Double) "B"+ at = char '@'+ reset = zeroWidthText "\r"+++{-# INLINE trace #-}+trace :: String -> IO a -> IO a+trace msg next = do+ Debug.traceIO Debug.dump_gc ("gc: " ++ msg)+ next++{-# INLINE message #-}+message :: String -> IO ()+message s = s `trace` return ()++{-# INLINE showContext #-}+showContext :: CUDA.Context -> String+showContext (CUDA.Context c) = show c+
+ Data/Array/Accelerate/LLVM/PTX/Debug.hs view
@@ -0,0 +1,101 @@+{-# LANGUAGE CPP #-}+{-# LANGUAGE TypeOperators #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Debug+-- Copyright : [2014..2017] Trevor L. McDonell+-- [2014..2014] Vinod Grover (NVIDIA Corporation)+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Debug (++ module Data.Array.Accelerate.Debug,+ module Data.Array.Accelerate.LLVM.PTX.Debug,++) where++import Data.Array.Accelerate.Debug hiding ( timed, elapsed )++import Foreign.CUDA.Driver.Stream ( Stream )+import qualified Foreign.CUDA.Driver.Event as Event++import Control.Concurrent+import Data.Label+import Data.Time.Clock+import System.CPUTime+import Text.Printf++import GHC.Float+++-- | Execute an action and time the results. The second argument specifies how+-- to format the output string given elapsed GPU and CPU time respectively+--+timed+ :: (Flags :-> Bool)+ -> (Double -> Double -> Double -> String)+ -> Maybe Stream+ -> IO ()+ -> IO ()+{-# INLINE timed #-}+timed f msg =+ monitorProcTime (queryFlag f) (\t1 t2 t3 -> traceIO f (msg t1 t2 t3))++monitorProcTime+ :: IO Bool+ -> (Double -> Double -> Double -> IO ())+ -> Maybe Stream+ -> IO ()+ -> IO ()+{-# INLINE monitorProcTime #-}+#if ACCELERATE_DEBUG+monitorProcTime enabled display stream action = do+ yes <- enabled+ if yes+ then do+ gpuBegin <- Event.create []+ gpuEnd <- Event.create []+ wallBegin <- getCurrentTime+ cpuBegin <- getCPUTime+ Event.record gpuBegin stream+ action+ Event.record gpuEnd stream+ cpuEnd <- getCPUTime+ wallEnd <- getCurrentTime++ -- Wait for the GPU to finish executing then display the timing execution+ -- message. Do this in a separate thread so that the remaining kernels can+ -- be queued asynchronously.+ --+ _ <- forkIO $ do+ Event.block gpuEnd+ diff <- Event.elapsedTime gpuBegin gpuEnd+ let gpuTime = float2Double $ diff * 1E-3 -- milliseconds+ cpuTime = fromIntegral (cpuEnd - cpuBegin) * 1E-12 -- picoseconds+ wallTime = realToFrac (diffUTCTime wallEnd wallBegin)++ Event.destroy gpuBegin+ Event.destroy gpuEnd+ --+ display wallTime cpuTime gpuTime+ --+ return ()++ else+ action+#else+monitorProcTime _ _ _ action = action+#endif++{-# INLINE elapsed #-}+elapsed :: Double -> Double -> Double -> String+elapsed wallTime cpuTime gpuTime =+ printf "%s (wall), %s (cpu), %s (gpu)"+ (showFFloatSIBase (Just 3) 1000 wallTime "s")+ (showFFloatSIBase (Just 3) 1000 cpuTime "s")+ (showFFloatSIBase (Just 3) 1000 gpuTime "s")+
+ Data/Array/Accelerate/LLVM/PTX/Execute.hs view
@@ -0,0 +1,601 @@+{-# LANGUAGE CPP #-}+{-# LANGUAGE FlexibleContexts #-}+{-# LANGUAGE GADTs #-}+{-# LANGUAGE RecordWildCards #-}+{-# LANGUAGE ScopedTypeVariables #-}+{-# LANGUAGE TemplateHaskell #-}+{-# LANGUAGE TypeOperators #-}+{-# OPTIONS_GHC -fno-warn-orphans #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Execute+-- Copyright : [2014..2017] Trevor L. McDonell+-- [2014..2014] Vinod Grover (NVIDIA Corporation)+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Execute (++ executeAcc, executeAfun1,++) where++-- accelerate+import Data.Array.Accelerate.Analysis.Match+import Data.Array.Accelerate.Array.Sugar+import Data.Array.Accelerate.Error+import Data.Array.Accelerate.Lifetime++import Data.Array.Accelerate.LLVM.Analysis.Match+import Data.Array.Accelerate.LLVM.Execute+import Data.Array.Accelerate.LLVM.State++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.Target+import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug++import Data.Range.Range ( Range(..) )+import Control.Parallel.Meta ( runExecutable )++-- cuda+import qualified Foreign.CUDA.Driver as CUDA++-- library+import Control.Monad ( when )+import Control.Monad.State ( gets, liftIO )+import Data.Int ( Int32 )+import Data.List ( find )+import Data.Maybe ( fromMaybe )+import Data.Word ( Word32 )+import Text.Printf ( printf )+import Prelude hiding ( exp, map, sum, scanl, scanr )+import qualified Prelude as P+++-- Array expression evaluation+-- ---------------------------++-- Computations are evaluated by traversing the AST bottom up, and for each node+-- distinguishing between three cases:+--+-- 1. If it is a Use node, we return a reference to the array data. The data+-- will already have been copied to the device during compilation of the+-- kernels.+--+-- 2. If it is a non-skeleton node, such as a let binding or shape conversion,+-- then execute directly by updating the environment or similar.+--+-- 3. If it is a skeleton node, then we need to execute the generated LLVM+-- code.+--+instance Execute PTX where+ map = simpleOp+ generate = simpleOp+ transform = simpleOp+ backpermute = simpleOp+ fold = foldOp+ fold1 = fold1Op+ foldSeg = foldSegOp+ fold1Seg = foldSegOp+ scanl = scanOp+ scanl1 = scan1Op+ scanl' = scan'Op+ scanr = scanOp+ scanr1 = scan1Op+ scanr' = scan'Op+ permute = permuteOp+ stencil1 = stencil1Op+ stencil2 = stencil2Op+++-- Skeleton implementation+-- -----------------------++-- Simple kernels just need to know the shape of the output array+--+simpleOp+ :: (Shape sh, Elt e)+ => ExecutableR PTX+ -> Gamma aenv+ -> Aval aenv+ -> Stream+ -> sh+ -> LLVM PTX (Array sh e)+simpleOp exe gamma aenv stream sh = do+ let kernel = case ptxKernel exe of+ k:_ -> k+ _ -> $internalError "simpleOp" "no kernels found"+ --+ out <- allocateRemote sh+ ptx <- gets llvmTarget+ liftIO $ executeOp ptx kernel gamma aenv stream (IE 0 (size sh)) out+ return out++simpleNamed+ :: (Shape sh, Elt e)+ => String+ -> ExecutableR PTX+ -> Gamma aenv+ -> Aval aenv+ -> Stream+ -> sh+ -> LLVM PTX (Array sh e)+simpleNamed fun exe gamma aenv stream sh = do+ let kernel = fromMaybe ($internalError "simpleNamed" ("not found: " ++ fun))+ $ lookupKernel fun exe+ --+ out <- allocateRemote sh+ ptx <- gets llvmTarget+ liftIO $ executeOp ptx kernel gamma aenv stream (IE 0 (size sh)) out+ return out+++-- There are two flavours of fold operation:+--+-- 1. If we are collapsing to a single value, then multiple thread blocks are+-- working together. Since thread blocks synchronise with each other via+-- kernel launches, each block computes a partial sum and the kernel is+-- launched recursively until the final value is reached.+--+-- 2. If this is a multidimensional reduction, then each inner dimension is+-- handled by a single thread block, so no global communication is+-- necessary. Furthermore are two kernel flavours: each innermost dimension+-- can be cooperatively reduced by (a) a thread warp; or (b) a thread+-- block. Currently we always use the first, but require benchmarking to+-- determine when to select each.+--+fold1Op+ :: (Shape sh, Elt e)+ => ExecutableR PTX+ -> Gamma aenv+ -> Aval aenv+ -> Stream+ -> (sh :. Int)+ -> LLVM PTX (Array sh e)+fold1Op exe gamma aenv stream sh@(sx :. sz)+ = $boundsCheck "fold1" "empty array" (sz > 0)+ $ case size sh of+ 0 -> allocateRemote sx -- empty, but possibly with one or more non-zero dimensions+ _ -> foldCore exe gamma aenv stream sh++foldOp+ :: (Shape sh, Elt e)+ => ExecutableR PTX+ -> Gamma aenv+ -> Aval aenv+ -> Stream+ -> (sh :. Int)+ -> LLVM PTX (Array sh e)+foldOp exe gamma aenv stream sh@(sx :. _)+ = case size sh of+ 0 -> simpleNamed "generate" exe gamma aenv stream (listToShape (P.map (max 1) (shapeToList sx)))+ _ -> foldCore exe gamma aenv stream sh++foldCore+ :: forall aenv sh e. (Shape sh, Elt e)+ => ExecutableR PTX+ -> Gamma aenv+ -> Aval aenv+ -> Stream+ -> (sh :. Int)+ -> LLVM PTX (Array sh e)+foldCore exe gamma aenv stream sh+ | Just Refl <- matchShapeType (undefined::sh) (undefined::Z)+ = foldAllOp exe gamma aenv stream sh+ --+ | otherwise+ = foldDimOp exe gamma aenv stream sh+++foldAllOp+ :: forall aenv e. Elt e+ => ExecutableR PTX+ -> Gamma aenv+ -> Aval aenv+ -> Stream+ -> DIM1+ -> LLVM PTX (Scalar e)+foldAllOp exe gamma aenv stream (Z :. n) = 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)+ --+ 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 gamma aenv stream (IE 0 n) out+ return out++ else do+ -- Multi-kernel reduction to a single element. The first kernel integrates+ -- any delayed elements, and the second is called recursively until+ -- reaching a single element.+ let+ rec :: Vector e -> LLVM PTX (Scalar e)+ rec tmp@(Array ((),m) adata)+ | m <= 1 = return $ Array () adata+ | otherwise = do+ let s = m `multipleOf` kernelThreadBlockSize km2+ out <- allocateRemote (Z :. s)+ 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 gamma aenv stream (IE 0 s) tmp+ rec tmp+++foldDimOp+ :: forall aenv sh e. (Shape sh, Elt e)+ => ExecutableR PTX+ -> Gamma aenv+ -> Aval aenv+ -> Stream+ -> (sh :. Int)+ -> LLVM PTX (Array sh e)+foldDimOp exe gamma aenv stream (sh :. sz) = do+ let+ kernel = fromMaybe ($internalError "foldDim" "kernel not found")+ $ if sz > 0+ then lookupKernel "fold" exe+ else lookupKernel "generate" exe+ --+ out <- allocateRemote sh+ ptx <- gets llvmTarget+ liftIO $ executeOp ptx kernel gamma aenv stream (IE 0 (size sh)) out+ return out+++foldSegOp+ :: (Shape sh, Elt e)+ => ExecutableR PTX+ -> Gamma aenv+ -> Aval aenv+ -> Stream+ -> (sh :. Int)+ -> (Z :. Int)+ -> LLVM PTX (Array (sh :. Int) e)+foldSegOp exe gamma aenv stream (sh :. sz) (Z :. ss) = 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+ --+ out <- allocateRemote (sh :. n)+ ptx <- gets llvmTarget+ liftIO $ do+ executeOp ptx foldseg gamma aenv stream (IE 0 m) out+ return out+++scanOp+ :: (Shape sh, Elt e)+ => ExecutableR PTX+ -> Gamma aenv+ -> Aval aenv+ -> Stream+ -> sh :. Int+ -> LLVM PTX (Array (sh:.Int) e)+scanOp exe gamma aenv stream (sz :. n) =+ case n of+ 0 -> simpleNamed "generate" exe gamma aenv stream (sz :. 1)+ _ -> scanCore exe gamma aenv stream sz n (n+1)++scan1Op+ :: (Shape sh, Elt e)+ => ExecutableR PTX+ -> Gamma aenv+ -> Aval aenv+ -> Stream+ -> sh :. Int+ -> LLVM PTX (Array (sh:.Int) e)+scan1Op exe gamma aenv stream (sz :. n)+ = $boundsCheck "scan1" "empty array" (n > 0)+ $ scanCore exe gamma aenv stream sz n n++scanCore+ :: forall aenv sh e. (Shape sh, Elt e)+ => ExecutableR PTX+ -> Gamma aenv+ -> Aval aenv+ -> Stream+ -> sh+ -> Int -- input size+ -> Int -- output size+ -> LLVM PTX (Array (sh:.Int) e)+scanCore exe gamma aenv stream sz n m+ | Just Refl <- matchShapeType (undefined::sh) (undefined::Z)+ = scanAllOp exe gamma aenv stream n m+ --+ | otherwise+ = scanDimOp exe gamma aenv stream sz m+++scanAllOp+ :: forall aenv e. Elt e+ => ExecutableR PTX+ -> Gamma aenv+ -> Aval aenv+ -> Stream+ -> Int -- input size+ -> Int -- output size+ -> LLVM PTX (Vector e)+scanAllOp exe gamma aenv stream n m = 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)+ --+ c = kernelThreadBlockSize k1+ s = n `multipleOf` c+ --+ ptx <- gets llvmTarget+ out <- allocateRemote (Z :. m)++ -- Step 1: Independent thread-block-wide scans of the input. Small arrays+ -- 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 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 gamma aenv stream (IE 0 s) tmp+ liftIO $ executeOp ptx k3 gamma aenv stream (IE 0 (s-1)) (tmp, out, i32 c)++ return out+++scanDimOp+ :: forall aenv sh e. (Shape sh, Elt e)+ => ExecutableR PTX+ -> Gamma aenv+ -> Aval aenv+ -> Stream+ -> sh+ -> Int+ -> LLVM PTX (Array (sh:.Int) e)+scanDimOp exe gamma aenv stream sz m = do+ let+ kernel = fromMaybe ($internalError "scanDimOp" "kernel not found")+ $ lookupKernel "scan" exe+ --+ ptx <- gets llvmTarget+ out <- allocateRemote (sz :. m)+ liftIO $ executeOp ptx kernel gamma aenv stream (IE 0 (size sz)) out+ return out+++scan'Op+ :: forall aenv sh e. (Shape sh, Elt e)+ => ExecutableR PTX+ -> Gamma aenv+ -> Aval aenv+ -> Stream+ -> sh :. Int+ -> LLVM PTX (Array (sh:.Int) e, Array sh e)+scan'Op exe gamma aenv stream sh@(sz :. n) =+ case n of+ 0 -> do out <- allocateRemote (sz :. 0)+ sum <- simpleNamed "generate" exe gamma aenv stream sz+ return (out, sum)+ _ -> scan'Core exe gamma aenv stream sh++scan'Core+ :: forall aenv sh e. (Shape sh, Elt e)+ => ExecutableR PTX+ -> Gamma aenv+ -> Aval aenv+ -> Stream+ -> sh :. Int+ -> LLVM PTX (Array (sh:.Int) e, Array sh e)+scan'Core exe gamma aenv stream sh+ | Just Refl <- matchShapeType (undefined::sh) (undefined::Z)+ = scan'AllOp exe gamma aenv stream sh+ --+ | otherwise+ = scan'DimOp exe gamma aenv stream sh++scan'AllOp+ :: forall aenv e. Elt e+ => ExecutableR PTX+ -> Gamma aenv+ -> Aval aenv+ -> Stream+ -> DIM1+ -> LLVM PTX (Vector e, Scalar e)+scan'AllOp exe gamma aenv stream (Z :. n) = 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)+ --+ c = kernelThreadBlockSize k1+ s = n `multipleOf` c+ --+ ptx <- gets llvmTarget+ out <- allocateRemote (Z :. n)+ tmp <- allocateRemote (Z :. s) :: LLVM PTX (Vector e)++ -- Step 1: independent thread-block-wide scans. Each block stores its partial+ -- sum to a temporary array.+ 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+ -- to the partial results.+ if s == 1+ then case tmp of+ Array _ ad -> return (out, Array () ad)+ else do+ sum <- allocateRemote Z+ 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)+++scan'DimOp+ :: forall aenv sh e. (Shape sh, Elt e)+ => ExecutableR PTX+ -> Gamma aenv+ -> Aval aenv+ -> Stream+ -> sh :. Int+ -> LLVM PTX (Array (sh:.Int) e, Array sh e)+scan'DimOp exe gamma aenv stream sh@(sz :. _) = do+ let kernel = fromMaybe ($internalError "scan'DimOp" "kernel not found")+ $ lookupKernel "scan" exe+ --+ ptx <- gets llvmTarget+ out <- allocateRemote sh+ sum <- allocateRemote sz+ liftIO $ executeOp ptx kernel gamma aenv stream (IE 0 (size sz)) (out,sum)+ return (out,sum)+++permuteOp+ :: (Shape sh, Shape sh', Elt e)+ => ExecutableR PTX+ -> Gamma aenv+ -> Aval aenv+ -> Stream+ -> Bool+ -> sh+ -> Array sh' e+ -> LLVM PTX (Array sh' e)+permuteOp exe gamma aenv stream inplace shIn dfs = do+ let n = size shIn+ m = size (shape dfs)+ kernel = case ptxKernel exe of+ k:_ -> k+ _ -> $internalError "permute" "no kernels found"+ --+ ptx <- gets llvmTarget+ out <- if inplace+ then return dfs+ else cloneArrayAsync stream dfs+ --+ case kernelName kernel of+ "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 gamma aenv stream (IE 0 n) (out, barrier)+ _ -> $internalError "permute" "unexpected kernel image"+ --+ return out+++-- Using the defaulting instances for stencil operations (for now).+--+stencil1Op+ :: (Shape sh, Elt b)+ => ExecutableR PTX+ -> Gamma aenv+ -> Aval aenv+ -> Stream+ -> Array sh a+ -> LLVM PTX (Array sh b)+stencil1Op exe gamma aenv stream arr =+ simpleOp exe gamma aenv stream (shape arr)++stencil2Op+ :: (Shape sh, Elt c)+ => ExecutableR PTX+ -> Gamma aenv+ -> Aval aenv+ -> Stream+ -> Array sh a+ -> Array sh b+ -> LLVM PTX (Array sh c)+stencil2Op exe gamma aenv stream arr brr =+ simpleOp exe gamma aenv stream (shape arr `intersect` shape brr)+++-- Skeleton execution+-- ------------------++-- TODO: Calculate this from the device properties, say [a multiple of] the+-- maximum number of in-flight threads that the device supports.+--+defaultPPT :: Int+defaultPPT = 32768++{-# INLINE i32 #-}+i32 :: Int -> Int32+i32 = fromIntegral++-- | Retrieve the named kernel+--+lookupKernel :: String -> ExecutableR PTX -> Maybe Kernel+lookupKernel name exe =+ find (\k -> kernelName k == name) (ptxKernel exe)+++-- Execute the function implementing this kernel.+--+executeOp+ :: Marshalable args+ => PTX+ -> Kernel+ -> Gamma aenv+ -> Aval aenv+ -> Stream+ -> Range+ -> args+ -> IO ()+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 stream (end-start) argv+++-- Execute a device function with the given thread configuration and function+-- parameters.+--+launch :: Kernel -> Stream -> Int -> [CUDA.FunParam] -> IO ()+launch Kernel{..} stream n args =+ when (n > 0) $+ 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++ -- Debugging/monitoring support+ 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)+
+ Data/Array/Accelerate/LLVM/PTX/Execute/Async.hs view
@@ -0,0 +1,64 @@+{-# LANGUAGE RecordWildCards #-}+{-# LANGUAGE TypeFamilies #-}+{-# OPTIONS_GHC -fno-warn-orphans #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Execute.Async+-- Copyright : [2014..2017] Trevor L. McDonell+-- [2014..2014] Vinod Grover (NVIDIA Corporation)+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Execute.Async (++ Async, Stream, Event,+ module Data.Array.Accelerate.LLVM.Execute.Async,++) where++-- accelerate+import Data.Array.Accelerate.LLVM.Execute.Async hiding ( Async )+import qualified Data.Array.Accelerate.LLVM.Execute.Async as A++import Data.Array.Accelerate.LLVM.PTX.Target+import Data.Array.Accelerate.LLVM.PTX.Execute.Event ( Event )+import Data.Array.Accelerate.LLVM.PTX.Execute.Stream ( Stream )+import qualified Data.Array.Accelerate.LLVM.PTX.Execute.Event as Event+import qualified Data.Array.Accelerate.LLVM.PTX.Execute.Stream as Stream++-- standard library+import Control.Monad.State+++-- Asynchronous arrays in the CUDA backend are tagged with an Event that will be+-- filled once the kernel implementing that array has completed.+--+type Async a = AsyncR PTX a++instance A.Async PTX where+ type StreamR PTX = Stream+ type EventR PTX = Event++ {-# INLINEABLE fork #-}+ fork =+ Stream.create++ {-# INLINEABLE join #-}+ join stream =+ liftIO $! Stream.destroy stream++ {-# INLINEABLE checkpoint #-}+ checkpoint stream =+ Event.waypoint stream++ {-# INLINEABLE after #-}+ after stream event =+ liftIO $! Event.after event stream++ {-# INLINEABLE block #-}+ block event =+ liftIO $! Event.block event+
+ Data/Array/Accelerate/LLVM/PTX/Execute/Environment.hs view
@@ -0,0 +1,22 @@+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Execute.Environment+-- Copyright : [2014..2017] Trevor L. McDonell+-- [2014..2014] Vinod Grover (NVIDIA Corporation)+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Execute.Environment (++ Aval, aprj++) where++import Data.Array.Accelerate.LLVM.PTX.Target+import Data.Array.Accelerate.LLVM.Execute.Environment++type Aval = AvalR PTX+
+ Data/Array/Accelerate/LLVM/PTX/Execute/Event.hs view
@@ -0,0 +1,162 @@+{-# LANGUAGE RecordWildCards #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Execute.Event+-- Copyright : [2014..2017] Trevor L. McDonell+-- [2014..2014] Vinod Grover (NVIDIA Corporation)+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Execute.Event (++ Event,+ create, destroy, query, waypoint, after, block,++) where++-- accelerate+import Data.Array.Accelerate.Lifetime+import qualified Data.Array.Accelerate.Array.Remote.LRU as Remote++import Data.Array.Accelerate.LLVM.PTX.Array.Remote ( )+import Data.Array.Accelerate.LLVM.PTX.Target ( PTX(..) )+import Data.Array.Accelerate.LLVM.State+import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug+import {-# SOURCE #-} Data.Array.Accelerate.LLVM.PTX.Execute.Stream++-- cuda+import Foreign.CUDA.Driver.Error+import qualified Foreign.CUDA.Driver.Event as Event+import qualified Foreign.CUDA.Driver.Stream as Stream++import Control.Exception+import Control.Monad.State+++-- | Events can be used for efficient device-side synchronisation between+-- execution streams and between the host.+--+type Event = Lifetime Event.Event+++-- | Create a new event. It will not be automatically garbage collected, and is+-- not suitable for timing purposes.+--+{-# INLINEABLE create #-}+create :: LLVM PTX Event+create = do+ e <- create'+ event <- liftIO $ newLifetime e+ liftIO $ addFinalizer event $ do message $ "destroy " ++ showEvent e+ Event.destroy e+ return event++create' :: LLVM PTX Event.Event+create' = do+ PTX{..} <- gets llvmTarget+ me <- attempt "create/new" (liftIO . catchOOM $ Event.create [Event.DisableTiming])+ `orElse` do+ Remote.reclaim ptxMemoryTable+ liftIO $ do+ message "create/new: failed (purging)"+ catchOOM $ Event.create [Event.DisableTiming]+ case me of+ Just e -> return e+ Nothing -> liftIO $ do+ message "create/new: failed (non-recoverable)"+ throwIO (ExitCode OutOfMemory)++ where+ catchOOM :: IO a -> IO (Maybe a)+ catchOOM it =+ liftM Just it `catch` \e -> case e of+ ExitCode OutOfMemory -> return Nothing+ _ -> throwIO e++ attempt :: MonadIO m => String -> m (Maybe a) -> m (Maybe a)+ attempt msg ea = do+ ma <- ea+ case ma of+ Nothing -> return Nothing+ Just a -> do liftIO (message msg)+ return (Just a)++ orElse :: MonadIO m => m (Maybe a) -> m (Maybe a) -> m (Maybe a)+ orElse ea eb = do+ ma <- ea+ case ma of+ Just a -> return (Just a)+ Nothing -> eb+++-- | Delete an event+--+{-# INLINEABLE destroy #-}+destroy :: Event -> IO ()+destroy = finalize++-- | Create a new event marker that will be filled once execution in the+-- specified stream has completed all previously submitted work.+--+{-# INLINEABLE waypoint #-}+waypoint :: Stream -> LLVM PTX Event+waypoint stream = do+ event <- create+ liftIO $+ withLifetime stream $ \s -> do+ withLifetime event $ \e -> do+ message $ "add waypoint " ++ showEvent e ++ " in stream " ++ showStream s+ Event.record e (Just s)+ return event++-- | Make all future work submitted to the given stream wait until the event+-- reports completion before beginning execution.+--+{-# INLINEABLE after #-}+after :: Event -> Stream -> IO ()+after event stream =+ withLifetime stream $ \s ->+ withLifetime event $ \e -> do+ message $ "after " ++ showEvent e ++ " in stream " ++ showStream s+ Event.wait e (Just s) []++-- | Block the calling thread until the event is recorded+--+{-# INLINEABLE block #-}+block :: Event -> IO ()+block event =+ withLifetime event $ \e -> do+ message $ "blocked on event " ++ showEvent e+ Event.block e++-- | Test whether an event has completed+--+{-# INLINEABLE query #-}+query :: Event -> IO Bool+query event = withLifetime event Event.query+++-- Debug+-- -----++{-# INLINE trace #-}+trace :: String -> IO a -> IO a+trace msg next = do+ Debug.traceIO Debug.dump_sched ("event: " ++ msg)+ next++{-# INLINE message #-}+message :: String -> IO ()+message s = s `trace` return ()++{-# INLINE showEvent #-}+showEvent :: Event.Event -> String+showEvent (Event.Event e) = show e++{-# INLINE showStream #-}+showStream :: Stream.Stream -> String+showStream (Stream.Stream s) = show s+
+ Data/Array/Accelerate/LLVM/PTX/Execute/Event.hs-boot view
@@ -0,0 +1,26 @@+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Execute.Event-boot+-- Copyright : [2016..2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Execute.Event (++ Event,+ query, block++) where++import Data.Array.Accelerate.Lifetime+import qualified Foreign.CUDA.Driver.Event as Event+++type Event = Lifetime Event.Event++query :: Event -> IO Bool+block :: Event -> IO ()+
+ Data/Array/Accelerate/LLVM/PTX/Execute/Marshal.hs view
@@ -0,0 +1,148 @@+{-# LANGUAGE CPP #-}+{-# LANGUAGE BangPatterns #-}+{-# LANGUAGE ConstraintKinds #-}+{-# LANGUAGE FlexibleContexts #-}+{-# LANGUAGE FlexibleInstances #-}+{-# LANGUAGE GADTs #-}+{-# LANGUAGE MultiParamTypeClasses #-}+{-# LANGUAGE RecordWildCards #-}+{-# LANGUAGE ScopedTypeVariables #-}+{-# LANGUAGE TypeFamilies #-}+{-# OPTIONS_GHC -fno-warn-orphans #-}+#if __GLASGOW_HASKELL__ <= 708+{-# LANGUAGE OverlappingInstances #-}+{-# OPTIONS_GHC -fno-warn-unrecognised-pragmas #-}+#endif+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Execute.Marshal+-- Copyright : [2014..2017] Trevor L. McDonell+-- [2014..2014] Vinod Grover (NVIDIA Corporation)+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Execute.Marshal (++ Marshalable, M.marshal++) where++-- accelerate+import Data.Array.Accelerate.LLVM.CodeGen.Environment ( Gamma, Idx'(..) )+import qualified Data.Array.Accelerate.LLVM.Execute.Marshal as M++import Data.Array.Accelerate.LLVM.PTX.State+import Data.Array.Accelerate.LLVM.PTX.Target+import Data.Array.Accelerate.LLVM.PTX.Array.Data+import Data.Array.Accelerate.LLVM.PTX.Execute.Async ( Async, AsyncR(..) )+import Data.Array.Accelerate.LLVM.PTX.Execute.Event ( after )+import Data.Array.Accelerate.LLVM.PTX.Execute.Environment+import qualified Data.Array.Accelerate.LLVM.PTX.Array.Prim as Prim++-- cuda+import qualified Foreign.CUDA.Driver as CUDA++-- libraries+import Control.Monad+import Data.Int+import Data.DList ( DList )+import Data.Typeable+import Foreign.Ptr+import Foreign.Storable ( Storable )+import qualified Data.DList as DL+import qualified Data.IntMap as IM+++-- Instances for the PTX backend+--+type Marshalable args = M.Marshalable PTX args+type instance M.ArgR PTX = CUDA.FunParam+++-- Instances for handling concrete types in this backend, namely shapes and+-- array data.+--+instance M.Marshalable PTX Int where+ marshal' _ _ x = return $ DL.singleton (CUDA.VArg x)++instance M.Marshalable PTX Int32 where+ marshal' _ _ x = return $ DL.singleton (CUDA.VArg x)++instance {-# OVERLAPS #-} M.Marshalable PTX (Gamma aenv, Aval aenv) where+ marshal' ptx stream (gamma, aenv)+ = fmap DL.concat+ $ mapM (\(_, Idx' idx) -> M.marshal' ptx stream =<< sync (aprj idx aenv)) (IM.elems gamma)+ where+ -- HAXORZ~ D:+ --+ -- The 'Async' class functions need to run in the LLVM monad, but the+ -- marshalling functions must run in IO because they will be executed in+ -- the lower-level scheduling code.+ --+ -- We hack around this impedance mismatch by calling the 'after'+ -- implementation directly.+ --+ sync :: Async a -> IO a+ sync (AsyncR event arr) = after event stream >> return arr++instance ArrayElt e => M.Marshalable PTX (ArrayData e) where+ marshal' ptx _ adata = do+ let marshalP :: forall e' a. (ArrayElt e', ArrayPtrs e' ~ Ptr a, Typeable e', Typeable a, Storable a)+ => ArrayData e'+ -> IO (DList CUDA.FunParam)+ marshalP ad =+ fmap (DL.singleton . CUDA.VArg)+ (unsafeGetDevicePtr ptx ad :: IO (CUDA.DevicePtr a))++ marshalR :: ArrayEltR e' -> ArrayData e' -> IO (DList CUDA.FunParam)+ marshalR ArrayEltRunit _ = return DL.empty+ marshalR (ArrayEltRpair aeR1 aeR2) ad =+ return DL.append `ap` marshalR aeR1 (fstArrayData ad)+ `ap` marshalR aeR2 (sndArrayData ad)+ marshalR ArrayEltRint ad = marshalP ad+ marshalR ArrayEltRint8 ad = marshalP ad+ marshalR ArrayEltRint16 ad = marshalP ad+ marshalR ArrayEltRint32 ad = marshalP ad+ marshalR ArrayEltRint64 ad = marshalP ad+ marshalR ArrayEltRword ad = marshalP ad+ marshalR ArrayEltRword8 ad = marshalP ad+ marshalR ArrayEltRword16 ad = marshalP ad+ marshalR ArrayEltRword32 ad = marshalP ad+ marshalR ArrayEltRword64 ad = marshalP ad+ marshalR ArrayEltRfloat ad = marshalP ad+ marshalR ArrayEltRdouble ad = marshalP ad+ marshalR ArrayEltRchar ad = marshalP ad+ marshalR ArrayEltRcshort ad = marshalP ad+ marshalR ArrayEltRcushort ad = marshalP ad+ marshalR ArrayEltRcint ad = marshalP ad+ marshalR ArrayEltRcuint ad = marshalP ad+ marshalR ArrayEltRclong ad = marshalP ad+ marshalR ArrayEltRculong ad = marshalP ad+ marshalR ArrayEltRcllong ad = marshalP ad+ marshalR ArrayEltRcullong ad = marshalP ad+ marshalR ArrayEltRcchar ad = marshalP ad+ marshalR ArrayEltRcschar ad = marshalP ad+ marshalR ArrayEltRcuchar ad = marshalP ad+ marshalR ArrayEltRcfloat ad = marshalP ad+ marshalR ArrayEltRcdouble ad = marshalP ad+ marshalR ArrayEltRbool ad = marshalP ad++ marshalR arrayElt adata+++-- TODO FIXME !!!+--+-- We will probably need to change marshal to be a bracketed function. We may+-- also want to reconsider whether to continue to restrict it to IO.+--+unsafeGetDevicePtr+ :: (ArrayElt e, ArrayPtrs e ~ Ptr a, Typeable e, Typeable a, Storable a)+ => PTX+ -> ArrayData e+ -> IO (CUDA.DevicePtr a)+unsafeGetDevicePtr !ptx !ad =+ evalPTX ptx $ Prim.withDevicePtr ad (\p -> return (Nothing,p))+
+ Data/Array/Accelerate/LLVM/PTX/Execute/Stream.hs view
@@ -0,0 +1,181 @@+{-# LANGUAGE BangPatterns #-}+{-# LANGUAGE MagicHash #-}+{-# LANGUAGE RecordWildCards #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Execute.Stream+-- Copyright : [2014..2017] Trevor L. McDonell+-- [2014..2014] Vinod Grover (NVIDIA Corporation)+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Execute.Stream (++ Reservoir, new,+ Stream, create, destroy, streaming,++) where++-- accelerate+import Data.Array.Accelerate.Lifetime+import qualified Data.Array.Accelerate.Array.Remote.LRU as Remote++import Data.Array.Accelerate.LLVM.PTX.Array.Remote ( )+import Data.Array.Accelerate.LLVM.PTX.Execute.Event ( Event )+import Data.Array.Accelerate.LLVM.PTX.Target ( PTX(..) )+import Data.Array.Accelerate.LLVM.State+import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug+import qualified Data.Array.Accelerate.LLVM.PTX.Execute.Event as Event+import Data.Array.Accelerate.LLVM.PTX.Execute.Stream.Reservoir as RSV++-- cuda+import Foreign.CUDA.Driver.Error+import qualified Foreign.CUDA.Driver.Stream as Stream++-- standard library+import Control.Exception+import Control.Monad.State+++-- | A 'Stream' represents an independent sequence of computations executed on+-- the GPU. Operations in different streams may be executed concurrently with+-- each other, but operations in the same stream can never overlap.+-- 'Data.Array.Accelerate.LLVM.PTX.Execute.Event.Event's can be used for+-- efficient cross-stream synchronisation.+--+type Stream = Lifetime Stream.Stream+++-- Executing operations in streams+-- -------------------------------++-- | Execute an operation in a unique execution stream. The (asynchronous)+-- result is passed to a second operation together with an event that will be+-- signalled once the operation is complete. The stream and event are released+-- after the second operation completes.+--+{-# INLINEABLE streaming #-}+streaming+ :: (Stream -> LLVM PTX a)+ -> (Event -> a -> LLVM PTX b)+ -> LLVM PTX b+streaming !action !after = do+ PTX{..} <- gets llvmTarget+ stream <- create+ first <- action stream+ end <- Event.waypoint stream+ final <- after end first+ liftIO $ do+ destroy stream+ Event.destroy end+ return final+++-- Primitive operations+-- --------------------++{--+-- | Delete all execution streams from the reservoir+--+{-# INLINEABLE flush #-}+flush :: Context -> Reservoir -> IO ()+flush !Context{..} !ref = do+ mc <- deRefWeak weakContext+ case mc of+ Nothing -> message "delete reservoir/dead context"+ Just ctx -> do+ message "flush reservoir"+ old <- swapMVar ref Seq.empty+ bracket_ (CUDA.push ctx) CUDA.pop $ Seq.mapM_ Stream.destroy old+--}+++-- | Create a CUDA execution stream. If an inactive stream is available for use,+-- use that, otherwise generate a fresh stream.+--+-- Note: [Finalising execution streams]+--+-- We don't actually ensure that the stream has executed all of its operations+-- to completion before attempting to return it to the reservoir for reuse.+-- Doing so increases overhead of the LLVM RTS due to 'forkIO', and consumes CPU+-- time as 'Stream.block' busy-waits for the stream to complete. It is quicker+-- to optimistically return the streams to the end of the reservoir immediately,+-- and just check whether the stream is done before reusing it.+--+-- > void . forkIO $ do+-- > Stream.block stream+-- > modifyMVar_ ref $ \rsv -> return (rsv Seq.|> stream)+--+{-# INLINEABLE create #-}+create :: LLVM PTX Stream+create = do+ PTX{..} <- gets llvmTarget+ s <- create'+ stream <- liftIO $ newLifetime s+ liftIO $ addFinalizer stream (RSV.insert ptxStreamReservoir s)+ return stream++create' :: LLVM PTX Stream.Stream+create' = do+ PTX{..} <- gets llvmTarget+ ms <- attempt "create/reservoir" (liftIO $ RSV.malloc ptxStreamReservoir)+ `orElse`+ attempt "create/new" (liftIO . catchOOM $ Stream.create [])+ `orElse` do+ Remote.reclaim ptxMemoryTable+ liftIO $ do+ message "create/new: failed (purging)"+ catchOOM $ Stream.create []+ case ms of+ Just s -> return s+ Nothing -> liftIO $ do+ message "create/new: failed (non-recoverable)"+ throwIO (ExitCode OutOfMemory)++ where+ catchOOM :: IO a -> IO (Maybe a)+ catchOOM it =+ liftM Just it `catch` \e -> case e of+ ExitCode OutOfMemory -> return Nothing+ _ -> throwIO e++ attempt :: MonadIO m => String -> m (Maybe a) -> m (Maybe a)+ attempt msg ea = do+ ma <- ea+ case ma of+ Nothing -> return Nothing+ Just a -> do liftIO (message msg)+ return (Just a)++ orElse :: MonadIO m => m (Maybe a) -> m (Maybe a) -> m (Maybe a)+ orElse ea eb = do+ ma <- ea+ case ma of+ Just a -> return (Just a)+ Nothing -> eb+++-- | Merge a stream back into the reservoir. This must only be done once all+-- pending operations in the stream have completed.+--+{-# INLINEABLE destroy #-}+destroy :: Stream -> IO ()+destroy = finalize+++-- Debug+-- -----++{-# INLINE trace #-}+trace :: String -> IO a -> IO a+trace msg next = do+ Debug.traceIO Debug.dump_sched ("stream: " ++ msg)+ next++{-# INLINE message #-}+message :: String -> IO ()+message s = s `trace` return ()+
+ Data/Array/Accelerate/LLVM/PTX/Execute/Stream.hs-boot view
@@ -0,0 +1,32 @@+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Execute.Stream-boot+-- Copyright : [2016..2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Execute.Stream (++ Stream,+ streaming,++) where++import Data.Array.Accelerate.Lifetime ( Lifetime )+import Data.Array.Accelerate.LLVM.State ( LLVM )+import Data.Array.Accelerate.LLVM.PTX.Target ( PTX )+import {-# SOURCE #-} Data.Array.Accelerate.LLVM.PTX.Execute.Event++import qualified Foreign.CUDA.Driver.Stream as Stream+++type Stream = Lifetime Stream.Stream++streaming+ :: (Stream -> LLVM PTX a)+ -> (Event -> a -> LLVM PTX b)+ -> LLVM PTX b+
+ Data/Array/Accelerate/LLVM/PTX/Execute/Stream/Reservoir.hs view
@@ -0,0 +1,102 @@+{-# LANGUAGE BangPatterns #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Execute.Stream.Reservoir+-- Copyright : [2016..2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Execute.Stream.Reservoir (++ Reservoir,+ new, malloc, insert,++) where++import Data.Array.Accelerate.LLVM.PTX.Context ( Context )+import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug++import Control.Concurrent.MVar+import Data.Sequence ( Seq )+import qualified Data.Sequence as Seq+import qualified Foreign.CUDA.Driver.Stream as Stream+++-- | The reservoir is a place to store CUDA execution streams that are currently+-- inactive. When a new stream is requested one is provided from the reservoir+-- if available, otherwise a fresh execution stream is created.+--+type Reservoir = MVar (Seq Stream.Stream)+++-- | Generate a new empty reservoir. It is not necessary to pre-populate it with+-- any streams because stream creation does not cause a device synchronisation.+--+-- Additionally, we do not need to finalise any of the streams. A reservoir is+-- tied to a specific execution context, so when the reservoir dies it is+-- because the PTX state and contained CUDA context have died, so there is+-- nothing more to do.+--+{-# INLINEABLE new #-}+new :: Context -> IO Reservoir+new _ctx = newMVar ( Seq.empty )+++-- | Retrieve an execution stream from the reservoir, if one is available.+--+-- Since we put streams back onto the reservoir once we have finished adding+-- work to them, not once they have completed execution of the tasks, we must+-- check for one which has actually completed.+--+-- See note: [Finalising execution streams]+--+{-# INLINEABLE malloc #-}+malloc :: Reservoir -> IO (Maybe Stream.Stream)+malloc !ref =+ modifyMVar ref (search Seq.empty)+ where+ -- scan through the streams in the reservoir looking for the first inactive+ -- one. Optimistically adding the streams to the end of the reservoir as+ -- soon as we stop assigning new work to them (c.f. async), and just+ -- checking they have completed before reusing them, is quicker than having+ -- a finaliser thread block until completion before retiring them.+ --+ search !acc !rsv =+ case Seq.viewl rsv of+ Seq.EmptyL -> return (acc, Nothing)+ s Seq.:< ss -> do+ done <- Stream.finished s+ case done of+ True -> return (acc Seq.>< ss, Just s)+ False -> search (acc Seq.|> s) ss+++-- | Add a stream to the reservoir+--+{-# INLINEABLE insert #-}+insert :: Reservoir -> Stream.Stream -> IO ()+insert !ref !stream = do+ message ("stash stream " ++ showStream stream)+ modifyMVar_ ref $ \rsv -> return (rsv Seq.|> stream)+++-- Debug+-- -----++{-# INLINE trace #-}+trace :: String -> IO a -> IO a+trace msg next = do+ Debug.traceIO Debug.dump_sched ("stream: " ++ msg)+ next++{-# INLINE message #-}+message :: String -> IO ()+message s = s `trace` return ()++{-# INLINE showStream #-}+showStream :: Stream.Stream -> String+showStream (Stream.Stream s) = show s+
+ Data/Array/Accelerate/LLVM/PTX/Foreign.hs view
@@ -0,0 +1,79 @@+{-# LANGUAGE DeriveDataTypeable #-}+{-# LANGUAGE GADTs #-}+{-# LANGUAGE ScopedTypeVariables #-}+{-# LANGUAGE StandaloneDeriving #-}+{-# OPTIONS_GHC -fno-warn-orphans #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Foreign+-- Copyright : [2016..2017] Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Foreign (++ -- Foreign functions+ ForeignAcc(..),+ ForeignExp(..),++ -- useful re-exports+ LLVM,+ PTX,+ liftIO,+ withDevicePtr,+ module Data.Array.Accelerate.LLVM.PTX.Array.Data,+ module Data.Array.Accelerate.LLVM.PTX.Execute.Async,++) where++import qualified Data.Array.Accelerate.Array.Sugar as S++import Data.Array.Accelerate.LLVM.State+import Data.Array.Accelerate.LLVM.CodeGen.Sugar++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.Execute.Async+import Data.Array.Accelerate.LLVM.PTX.Target ( PTX )++import Control.Monad.State+import Data.Typeable+++instance Foreign PTX where+ foreignAcc _ (ff :: asm (a -> b))+ | Just (ForeignAcc _ asm :: ForeignAcc (a -> b)) <- cast ff = Just asm+ | otherwise = Nothing++ foreignExp _ (ff :: asm (x -> y))+ | Just (ForeignExp _ asm :: ForeignExp (x -> y)) <- cast ff = Just asm+ | otherwise = Nothing++instance S.Foreign ForeignAcc where+ strForeign (ForeignAcc s _) = s++instance S.Foreign ForeignExp where+ strForeign (ForeignExp s _) = s+++-- Foreign functions in the PTX backend.+--+data ForeignAcc f where+ ForeignAcc :: String+ -> (Stream -> a -> LLVM PTX b)+ -> ForeignAcc (a -> b)++-- Foreign expressions in the PTX backend.+--+data ForeignExp f where+ ForeignExp :: String+ -> IRFun1 PTX () (x -> y)+ -> ForeignExp (x -> y)++deriving instance Typeable ForeignAcc+deriving instance Typeable ForeignExp+
+ Data/Array/Accelerate/LLVM/PTX/State.hs view
@@ -0,0 +1,108 @@+{-# LANGUAGE CPP #-}+{-# LANGUAGE TemplateHaskell #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.State+-- Copyright : [2014..2017] Trevor L. McDonell+-- [2014..2014] Vinod Grover (NVIDIA Corporation)+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.State (++ evalPTX,+ createTargetForDevice, createTargetFromContext, defaultTarget++) where++-- 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 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.Debug as Debug++import Data.Range.Range ( Range(..) )+import Control.Parallel.Meta ( Executable(..) )++-- standard library+import Control.Concurrent ( runInBoundThread )+import Control.Exception ( catch )+import System.IO.Unsafe ( unsafePerformIO )+import Foreign.CUDA.Driver.Error+import qualified Foreign.CUDA.Driver as CUDA+import qualified Foreign.CUDA.Driver.Context as Context+++-- | Execute a PTX computation+--+evalPTX :: PTX -> LLVM PTX a -> IO a+evalPTX ptx acc =+ runInBoundThread (CT.withContext (ptxContext ptx) (evalLLVM ptx acc))+ `catch`+ \e -> $internalError "unhandled" (show (e :: CUDAException))+++-- | Create a new PTX execution target for the given device+--+createTargetForDevice+ :: CUDA.Device+ -> CUDA.DeviceProperties+ -> [CUDA.ContextFlag]+ -> IO PTX+createTargetForDevice dev prp flags = do+ ctx <- CT.new dev prp flags+ mt <- MT.new ctx+ st <- ST.new ctx+ return $! PTX ctx mt st simpleIO+++-- | Create a PTX execute target for the given device context+--+createTargetFromContext+ :: CUDA.Context+ -> IO PTX+createTargetFromContext ctx' = do+ dev <- Context.device+ prp <- CUDA.props dev+ ctx <- CT.raw dev prp ctx'+ mt <- MT.new ctx+ st <- ST.new ctx+ return $! PTX ctx mt st simpleIO+++{-# INLINE simpleIO #-}+simpleIO :: Executable+simpleIO = Executable $ \_name _ppt range action ->+ case range of+ Empty -> return ()+ IE u v -> action u v 0+++-- Top-level mutable state+-- -----------------------+--+-- It is important to keep some information alive for the entire run of the+-- program, not just a single execution. These tokens use 'unsafePerformIO' to+-- ensure they are executed only once, and reused for subsequent invocations.+--++-- | Select and initialise the default CUDA device, and create a new target+-- context. The device is selected based on compute capability and estimated+-- maximum throughput.+--+{-# NOINLINE defaultTarget #-}+defaultTarget :: PTX+defaultTarget = unsafePerformIO $ do+ Debug.traceIO Debug.dump_gc "gc: initialise default PTX target"+ CUDA.initialise []+ (dev,prp) <- selectBestDevice+ createTargetForDevice dev prp [CUDA.SchedAuto]+
+ Data/Array/Accelerate/LLVM/PTX/Target.hs view
@@ -0,0 +1,176 @@+{-# LANGUAGE CPP #-}+{-# LANGUAGE EmptyDataDecls #-}+{-# LANGUAGE TemplateHaskell #-}+{-# LANGUAGE TypeFamilies #-}+-- |+-- Module : Data.Array.Accelerate.LLVM.PTX.Target+-- Copyright : [2014..2017] Trevor L. McDonell+-- [2014..2014] Vinod Grover (NVIDIA Corporation)+-- License : BSD3+--+-- Maintainer : Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.LLVM.PTX.Target (++ module Data.Array.Accelerate.LLVM.Target,+ module Data.Array.Accelerate.LLVM.PTX.Target,++) where++-- llvm-general+import LLVM.AST.AddrSpace+import LLVM.AST.DataLayout+import LLVM.Target hiding ( Target )+import qualified LLVM.Target as LLVM+import qualified LLVM.Relocation as R+import qualified LLVM.CodeModel as CM+import qualified LLVM.CodeGenOpt as CGO++-- accelerate+import Data.Array.Accelerate.Error++import Data.Array.Accelerate.LLVM.Target+import Data.Array.Accelerate.LLVM.Util++import Control.Parallel.Meta ( Executable )+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 )++-- CUDA+import qualified Foreign.CUDA.Driver as CUDA++-- standard library+import Control.Monad.Except+import System.IO.Unsafe+import Text.Printf+import qualified Data.Map as Map+import qualified Data.Set as Set+++-- | The PTX execution target for NVIDIA GPUs.+--+-- The execution target carries state specific for the current execution+-- context. The data here --- device memory and execution streams --- are+-- implicitly tied to this CUDA execution context.+--+-- Don't store anything here that is independent of the context, for example+-- state related to [persistent] kernel caching should _not_ go here.+--+data PTX = PTX {+ ptxContext :: {-# UNPACK #-} !Context+ , ptxMemoryTable :: {-# UNPACK #-} !MemoryTable+ , ptxStreamReservoir :: {-# UNPACK #-} !Reservoir+ , fillP :: {-# UNPACK #-} !Executable+ }++instance Target PTX where+ targetTriple _ = Just ptxTargetTriple+#if ACCELERATE_USE_NVVM+ targetDataLayout _ = Nothing -- see note: [NVVM and target data layout]+#else+ targetDataLayout _ = Just ptxDataLayout+#endif+++-- | Extract the properties of the device the current PTX execution state is+-- executing on.+--+ptxDeviceProperties :: PTX -> CUDA.DeviceProperties+ptxDeviceProperties = deviceProperties . ptxContext+++-- | A description of the various data layout properties that may be used during+-- optimisation. For CUDA the following data layouts are supported:+--+-- 32-bit:+-- e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64+--+-- 64-bit:+-- e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64+--+-- Thus, only the size of the pointer layout changes depending on the host+-- architecture.+--+ptxDataLayout :: DataLayout+ptxDataLayout = DataLayout+ { endianness = LittleEndian+ , mangling = Nothing+ , aggregateLayout = AlignmentInfo 0 (Just 64)+ , stackAlignment = Nothing+ , pointerLayouts = Map.fromList+ [ (AddrSpace 0, (wordSize, AlignmentInfo wordSize (Just 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] ]+ , nativeSizes = Just $ Set.fromList [ 16,32,64 ]+ }+ where+ wordSize = bitSize (undefined :: Int)+++-- | String that describes the target host.+--+ptxTargetTriple :: String+ptxTargetTriple =+ case bitSize (undefined::Int) of+ 32 -> "nvptx-nvidia-cuda"+ 64 -> "nvptx64-nvidia-cuda"+ _ -> $internalError "ptxTargetTriple" "I don't know what architecture I am"+++-- | Bracket creation and destruction of the NVVM TargetMachine.+--+withPTXTargetMachine+ :: CUDA.DeviceProperties+ -> (TargetMachine -> IO a)+ -> IO a+withPTXTargetMachine dev go =+ let CUDA.Compute m n = CUDA.computeCapability dev+ isa = ptxISAVersion m n+ sm = printf "sm_%d%d" m n+ in+ withTargetOptions $ \options -> do+ withTargetMachine+ ptxTarget+ ptxTargetTriple+ sm+ (Map.singleton isa True) -- CPU features+ options -- target options+ R.Default -- relocation model+ CM.Default -- code model+ CGO.Default -- optimisation level+ go++-- Some libdevice functions require at least ptx40, even though devices at+-- that compute capability also accept older ISA versions.+--+-- 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"+++-- | The NVPTX target for this host.+--+-- The top-level 'unsafePerformIO' is so that 'initializeAllTargets' is run once+-- per program execution (although that might not be necessary?)+--+{-# NOINLINE ptxTarget #-}+ptxTarget :: LLVM.Target+ptxTarget = unsafePerformIO $ do+ initializeAllTargets+ either error fst `fmap` runExceptT (lookupTarget Nothing ptxTargetTriple)+
+ LICENSE view
@@ -0,0 +1,23 @@+Copyright (c) [2014..2017] The Accelerate Team. All rights reserved.++Redistribution and use in source and binary forms, with or without+modification, are permitted provided that the following conditions are met:+ * Redistributions of source code must retain the above copyright+ notice, this list of conditions and the following disclaimer.+ * Redistributions in binary form must reproduce the above copyright+ notice, this list of conditions and the following disclaimer in the+ documentation and/or other materials provided with the distribution.+ * Neither the names of the contributors nor of their affiliations may+ be used to endorse or promote products derived from this software+ without specific prior written permission.++THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ''AS IS'' AND ANY+EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED+WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE+DISCLAIMED. IN NO EVENT SHALL COPYRIGHT HOLDERS BE LIABLE FOR ANY+DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES+(INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;+LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND+ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT+(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS+SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ Setup.hs view
@@ -0,0 +1,2 @@+import Distribution.Simple+main = defaultMain
+ accelerate-llvm-ptx.cabal view
@@ -0,0 +1,147 @@+name: accelerate-llvm-ptx+version: 1.0.0.0+cabal-version: >= 1.10+tested-with: GHC == 7.8.*+build-type: Simple++synopsis: Accelerate backend generating LLVM+description:+ This library implements a backend for the /Accelerate/ language which+ generates LLVM-IR targeting CUDA capable GPUs. For further information,+ refer to the main /Accelerate/ package:+ <http://hackage.haskell.org/package/accelerate>++license: BSD3+license-file: LICENSE+author: Trevor L. McDonell+maintainer: Trevor L. McDonell <tmcdonell@cse.unsw.edu.au>+bug-reports: https://github.com/AccelerateHS/accelerate/issues+category: Compilers/Interpreters, Concurrency, Data, Parallelism+++-- Configuration flags+-- -------------------++Flag nvvm+ Default: False+ Description: Use the NVVM library to generate optimised PTX++Flag debug+ Default: False+ Description:+ Enable debug tracing message flags. Note that 'debug' must be enabled in the+ base 'accelerate' package as well. See the 'accelerate' package for usage+ and available options.++Flag bounds-checks+ Default: True+ Description: Enable bounds checking++Flag unsafe-checks+ Default: False+ Description: Enable bounds checking in unsafe operations++Flag internal-checks+ Default: False+ Description: Enable internal consistency checks+++-- Build configuration+-- -------------------++Library+ exposed-modules:+ Data.Array.Accelerate.LLVM.PTX+ Data.Array.Accelerate.LLVM.PTX.Foreign++ other-modules:+ Data.Array.Accelerate.LLVM.PTX.Analysis.Device+ Data.Array.Accelerate.LLVM.PTX.Analysis.Launch+ Data.Array.Accelerate.LLVM.PTX.Array.Data+ Data.Array.Accelerate.LLVM.PTX.Array.Prim+ Data.Array.Accelerate.LLVM.PTX.Array.Remote+ Data.Array.Accelerate.LLVM.PTX.Array.Table+ Data.Array.Accelerate.LLVM.PTX.Context+ Data.Array.Accelerate.LLVM.PTX.Debug+ Data.Array.Accelerate.LLVM.PTX.State+ Data.Array.Accelerate.LLVM.PTX.Target++ Data.Array.Accelerate.LLVM.PTX.Execute+ Data.Array.Accelerate.LLVM.PTX.Execute.Async+ Data.Array.Accelerate.LLVM.PTX.Execute.Stream+ Data.Array.Accelerate.LLVM.PTX.Execute.Stream.Reservoir++ Data.Array.Accelerate.LLVM.PTX.CodeGen+ Data.Array.Accelerate.LLVM.PTX.CodeGen.Base+ 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.Loop+ Data.Array.Accelerate.LLVM.PTX.CodeGen.Map+ Data.Array.Accelerate.LLVM.PTX.CodeGen.Permute+ Data.Array.Accelerate.LLVM.PTX.CodeGen.Queue+ 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.Libdevice++ Data.Array.Accelerate.LLVM.PTX.Execute.Environment+ Data.Array.Accelerate.LLVM.PTX.Execute.Event+ Data.Array.Accelerate.LLVM.PTX.Execute.Marshal++ build-depends:+ base >= 4.7 && < 4.10+ , accelerate == 1.0.*+ , accelerate-llvm == 1.0.*+ , bytestring >= 0.9+ , containers >= 0.5 && <0.6+ , cuda >= 0.7+ , directory >= 1.0+ , dlist >= 0.6+ , fclabels >= 2.0+ , filepath >= 1.0+ , hashable >= 1.2+ , llvm-hs >= 3.9+ , llvm-hs-pure >= 3.9+ , mtl >= 2.2.1+ , pretty >= 1.1+ , time >= 1.4+ , unordered-containers >= 0.2++ default-language:+ Haskell2010++ ghc-options: -O2 -Wall -fwarn-tabs++ if impl(ghc >= 8.0)+ ghc-options: -Wmissed-specialisations++ if flag(nvvm)+ cpp-options: -DACCELERATE_USE_NVVM+ build-depends:+ nvvm >= 0.7.5++ if flag(debug)+ cpp-options: -DACCELERATE_DEBUG++ if flag(bounds-checks)+ cpp-options: -DACCELERATE_BOUNDS_CHECKS++ if flag(unsafe-checks)+ cpp-options: -DACCELERATE_UNSAFE_CHECKS++ if flag(internal-checks)+ cpp-options: -DACCELERATE_INTERNAL_CHECKS+++source-repository head+ type: git+ location: https://github.com/AccelerateHS/accelerate-llvm.git++source-repository this+ type: git+ tag: 1.0.0.0+ location: https://github.com/AccelerateHS/accelerate-llvm.git++-- vim: nospell