diff --git a/Data/Array/Accelerate/LLVM/PTX.hs b/Data/Array/Accelerate/LLVM/PTX.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX.hs @@ -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 +-- 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 + diff --git a/Data/Array/Accelerate/LLVM/PTX/Analysis/Device.hs b/Data/Array/Accelerate/LLVM/PTX/Analysis/Device.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Analysis/Device.hs @@ -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 +-- 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 + diff --git a/Data/Array/Accelerate/LLVM/PTX/Analysis/Launch.hs b/Data/Array/Accelerate/LLVM/PTX/Analysis/Launch.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Analysis/Launch.hs @@ -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 +-- 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) + diff --git a/Data/Array/Accelerate/LLVM/PTX/Array/Data.hs b/Data/Array/Accelerate/LLVM/PTX/Array/Data.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Array/Data.hs @@ -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 +-- 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 + diff --git a/Data/Array/Accelerate/LLVM/PTX/Array/Prim.hs b/Data/Array/Accelerate/LLVM/PTX/Array/Prim.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Array/Prim.hs @@ -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 +-- 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) + diff --git a/Data/Array/Accelerate/LLVM/PTX/Array/Remote.hs b/Data/Array/Accelerate/LLVM/PTX/Array/Remote.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Array/Remote.hs @@ -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 +-- 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 + diff --git a/Data/Array/Accelerate/LLVM/PTX/Array/Table.hs b/Data/Array/Accelerate/LLVM/PTX/Array/Table.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Array/Table.hs @@ -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 +-- 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 () + + diff --git a/Data/Array/Accelerate/LLVM/PTX/CodeGen.hs b/Data/Array/Accelerate/LLVM/PTX/CodeGen.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/CodeGen.hs @@ -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 +-- 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 + diff --git a/Data/Array/Accelerate/LLVM/PTX/CodeGen/Base.hs b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Base.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Base.hs @@ -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 +-- 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 +-- +-- +-- +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. +-- +-- +-- +__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. +-- +-- +-- +__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. +-- +-- +-- +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 + } + } + diff --git a/Data/Array/Accelerate/LLVM/PTX/CodeGen/Fold.hs b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Fold.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Fold.hs @@ -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 +-- 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 + diff --git a/Data/Array/Accelerate/LLVM/PTX/CodeGen/FoldSeg.hs b/Data/Array/Accelerate/LLVM/PTX/CodeGen/FoldSeg.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/CodeGen/FoldSeg.hs @@ -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 +-- 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 + diff --git a/Data/Array/Accelerate/LLVM/PTX/CodeGen/Generate.hs b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Generate.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Generate.hs @@ -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 +-- 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_ + + diff --git a/Data/Array/Accelerate/LLVM/PTX/CodeGen/Loop.hs b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Loop.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Loop.hs @@ -0,0 +1,47 @@ +-- | +-- Module : Data.Array.Accelerate.LLVM.PTX.CodeGen.Loop +-- Copyright : [2015..2017] Trevor L. McDonell +-- License : BSD3 +-- +-- Maintainer : Trevor L. McDonell +-- 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 + diff --git a/Data/Array/Accelerate/LLVM/PTX/CodeGen/Map.hs b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Map.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Map.hs @@ -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 +-- 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_ + diff --git a/Data/Array/Accelerate/LLVM/PTX/CodeGen/Permute.hs b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Permute.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Permute.hs @@ -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 +-- 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. + -- + -- 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)) + diff --git a/Data/Array/Accelerate/LLVM/PTX/CodeGen/Queue.hs b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Queue.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Queue.hs @@ -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 +-- 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_ + diff --git a/Data/Array/Accelerate/LLVM/PTX/CodeGen/Scan.hs b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Scan.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/CodeGen/Scan.hs @@ -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 +-- 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' + diff --git a/Data/Array/Accelerate/LLVM/PTX/Compile.hs b/Data/Array/Accelerate/LLVM/PTX/Compile.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Compile.hs @@ -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 +-- 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 + ] +--} + diff --git a/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice.hs b/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Compile/Libdevice.hs @@ -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 +-- 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" + ] + diff --git a/Data/Array/Accelerate/LLVM/PTX/Compile/Link.hs b/Data/Array/Accelerate/LLVM/PTX/Compile/Link.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Compile/Link.hs @@ -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 +-- Stability : experimental +-- Portability : non-portable (GHC extensions) +-- + +module Data.Array.Accelerate.LLVM.PTX.Compile.Link ( + + withLibdeviceNVVM, + withLibdeviceNVPTX, + +) where + +-- llvm-hs +import LLVM.Context +import qualified LLVM.Module as LLVM + +import LLVM.AST as AST +import LLVM.AST.Global as G +import LLVM.AST.Linkage + +-- accelerate +import Data.Array.Accelerate.Error + +import Data.Array.Accelerate.LLVM.PTX.Compile.Libdevice +import qualified Data.Array.Accelerate.LLVM.PTX.Debug as Debug + +-- cuda +import Foreign.CUDA.Analysis + +-- standard library +import Control.Monad.Except +import Data.ByteString ( ByteString ) +import Data.HashSet ( HashSet ) +import Data.List +import Data.Maybe +import Text.Printf +import qualified Data.HashSet as Set + + +-- | Lower an LLVM AST to C++ objects and link it against the libdevice module, +-- iff any libdevice functions are referenced from the base module. +-- +-- Note: [Linking with libdevice] +-- +-- The CUDA toolkit comes with an LLVM bitcode library called 'libdevice' that +-- implements many common mathematical functions. The library can be used as a +-- high performance math library for targets of the LLVM NVPTX backend, such as +-- this one. To link a module 'foo' with libdevice, the following compilation +-- pipeline is recommended: +-- +-- 1. Save all external functions in module 'foo' +-- +-- 2. Link module 'foo' with the appropriate 'libdevice_compute_XX.YY.bc' +-- +-- 3. Internalise all functions not in the list from (1) +-- +-- 4. Eliminate all unused internal functions +-- +-- 5. Run the NVVMReflect pass (see note: [NVVM Reflect Pass]) +-- +-- 6. Run the standard optimisation pipeline +-- +withLibdeviceNVPTX + :: DeviceProperties + -> Context + -> Module + -> (LLVM.Module -> IO a) + -> IO a +withLibdeviceNVPTX dev ctx ast next = + case Set.null externs of + True -> runError $ LLVM.withModuleFromAST ctx ast next + False -> + runError $ LLVM.withModuleFromAST ctx ast $ \mdl -> + runError $ LLVM.withModuleFromAST ctx nvvmReflect $ \refl -> + runError $ LLVM.withModuleFromAST ctx (internalise externs libdev) $ \libd -> do + runError $ linkModules mdl refl + runError $ linkModules mdl libd + Debug.traceIO Debug.dump_cc msg + next mdl + where + -- Replace the target triple and datalayout from the libdevice.bc module + -- with those of the generated code. This avoids warnings such as "linking + -- two modules of different target triples..." + libdev = (libdevice arch) { moduleTargetTriple = moduleTargetTriple ast + , moduleDataLayout = moduleDataLayout ast + } + externs = analyse ast + arch = computeCapability dev + runError = either ($internalError "withLibdeviceNVPTX") return <=< runExceptT + + msg = printf "cc: linking with libdevice: %s" + $ intercalate ", " (Set.toList externs) + + +-- | Link LLVM modules by copying parts of the second argument into the first. +-- +linkModules + :: LLVM.Module -- module into which to link (destination: contains all symbols) + -> LLVM.Module -- module to copy into the other (this is destroyed in the process) + -> ExceptT String IO () +linkModules = LLVM.linkModules + + +-- | Lower an LLVM AST to C++ objects and prepare it for linking against +-- libdevice using the nvvm bindings, iff any libdevice functions are referenced +-- from the base module. +-- +-- Rather than internalise and strip any unused functions ourselves, allow the +-- nvvm library to do so when linking the two modules together. +-- +-- TLM: This really should work with the above method, however for some reason +-- we get a "CUDA Exception: function named symbol not found" error, even though +-- the function is clearly visible in the generated code. hmm... +-- +withLibdeviceNVVM + :: DeviceProperties + -> Context + -> Module + -> ([(String, ByteString)] -> LLVM.Module -> IO a) + -> IO a +withLibdeviceNVVM dev ctx ast next = + runError $ LLVM.withModuleFromAST ctx ast $ \mdl -> do + when withlib $ Debug.traceIO Debug.dump_cc msg + next lib mdl + where + externs = analyse ast + withlib = not (Set.null externs) + lib | withlib = [ nvvmReflect, libdevice arch ] + | otherwise = [] + + arch = computeCapability dev + runError = either ($internalError "withLibdeviceNVPTX") return <=< runExceptT + + msg = printf "cc: linking with libdevice: %s" + $ intercalate ", " (Set.toList externs) + + +-- | Analyse the LLVM AST module and determine if any of the external +-- declarations are intrinsics implemented by libdevice. The set of such +-- functions is returned, and will be used when determining which functions from +-- libdevice to internalise. +-- +analyse :: Module -> HashSet String +analyse Module{..} = + let intrinsic (GlobalDefinition Function{..}) + | null basicBlocks + , Name n <- name + , "__nv_" <- take 5 n + = Just n + + intrinsic _ + = Nothing + in + Set.fromList (mapMaybe intrinsic moduleDefinitions) + + +-- | Mark all definitions in the module as internal linkage. This means that +-- unused definitions can be removed as dead code. Be careful to leave any +-- declarations as external. +-- +internalise :: HashSet String -> Module -> Module +internalise externals Module{..} = + let internal (GlobalDefinition Function{..}) + | Name n <- name + , not (Set.member n externals) -- we don't call this function directly; and + , not (null basicBlocks) -- it is not an external declaration + = GlobalDefinition (Function { linkage=Internal, .. }) + + internal x + = x + in + Module { moduleDefinitions = map internal moduleDefinitions, .. } + diff --git a/Data/Array/Accelerate/LLVM/PTX/Context.hs b/Data/Array/Accelerate/LLVM/PTX/Context.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Context.hs @@ -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 +-- 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 + diff --git a/Data/Array/Accelerate/LLVM/PTX/Debug.hs b/Data/Array/Accelerate/LLVM/PTX/Debug.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Debug.hs @@ -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 +-- 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") + diff --git a/Data/Array/Accelerate/LLVM/PTX/Execute.hs b/Data/Array/Accelerate/LLVM/PTX/Execute.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Execute.hs @@ -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 +-- 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) + diff --git a/Data/Array/Accelerate/LLVM/PTX/Execute/Async.hs b/Data/Array/Accelerate/LLVM/PTX/Execute/Async.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Execute/Async.hs @@ -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 +-- 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 + diff --git a/Data/Array/Accelerate/LLVM/PTX/Execute/Environment.hs b/Data/Array/Accelerate/LLVM/PTX/Execute/Environment.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Execute/Environment.hs @@ -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 +-- 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 + diff --git a/Data/Array/Accelerate/LLVM/PTX/Execute/Event.hs b/Data/Array/Accelerate/LLVM/PTX/Execute/Event.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Execute/Event.hs @@ -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 +-- 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 + diff --git a/Data/Array/Accelerate/LLVM/PTX/Execute/Event.hs-boot b/Data/Array/Accelerate/LLVM/PTX/Execute/Event.hs-boot new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Execute/Event.hs-boot @@ -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 +-- 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 () + diff --git a/Data/Array/Accelerate/LLVM/PTX/Execute/Marshal.hs b/Data/Array/Accelerate/LLVM/PTX/Execute/Marshal.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Execute/Marshal.hs @@ -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 +-- 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)) + diff --git a/Data/Array/Accelerate/LLVM/PTX/Execute/Stream.hs b/Data/Array/Accelerate/LLVM/PTX/Execute/Stream.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Execute/Stream.hs @@ -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 +-- 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 () + diff --git a/Data/Array/Accelerate/LLVM/PTX/Execute/Stream.hs-boot b/Data/Array/Accelerate/LLVM/PTX/Execute/Stream.hs-boot new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Execute/Stream.hs-boot @@ -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 +-- 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 + diff --git a/Data/Array/Accelerate/LLVM/PTX/Execute/Stream/Reservoir.hs b/Data/Array/Accelerate/LLVM/PTX/Execute/Stream/Reservoir.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Execute/Stream/Reservoir.hs @@ -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 +-- 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 + diff --git a/Data/Array/Accelerate/LLVM/PTX/Foreign.hs b/Data/Array/Accelerate/LLVM/PTX/Foreign.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Foreign.hs @@ -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 +-- 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 + diff --git a/Data/Array/Accelerate/LLVM/PTX/State.hs b/Data/Array/Accelerate/LLVM/PTX/State.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/State.hs @@ -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 +-- 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] + diff --git a/Data/Array/Accelerate/LLVM/PTX/Target.hs b/Data/Array/Accelerate/LLVM/PTX/Target.hs new file mode 100644 --- /dev/null +++ b/Data/Array/Accelerate/LLVM/PTX/Target.hs @@ -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 +-- 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) + diff --git a/LICENSE b/LICENSE new file mode 100644 --- /dev/null +++ b/LICENSE @@ -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. diff --git a/Setup.hs b/Setup.hs new file mode 100644 --- /dev/null +++ b/Setup.hs @@ -0,0 +1,2 @@ +import Distribution.Simple +main = defaultMain diff --git a/accelerate-llvm-ptx.cabal b/accelerate-llvm-ptx.cabal new file mode 100644 --- /dev/null +++ b/accelerate-llvm-ptx.cabal @@ -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: + + +license: BSD3 +license-file: LICENSE +author: Trevor L. McDonell +maintainer: Trevor L. McDonell +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