Obsidian 0.1.0.0 → 0.4.0.0
raw patch · 22 files changed
+2625/−1165 lines, 22 filesdep ~cudadep ~language-c-quote
Dependency ranges changed: cuda, language-c-quote
Files
- Obsidian.cabal +10/−7
- Obsidian.hs +14/−12
- Obsidian/Array.hs +303/−101
- Obsidian/Atomic.hs +1/−0
- Obsidian/CodeGen/CUDA.hs +41/−7
- Obsidian/CodeGen/CompileIM.hs +239/−234
- Obsidian/CodeGen/Liveness.hs +13/−2
- Obsidian/CodeGen/Memory.hs +0/−249
- Obsidian/CodeGen/Memory2.hs +457/−0
- Obsidian/CodeGen/Program.hs +70/−65
- Obsidian/CodeGen/Reify.hs +17/−9
- Obsidian/Data.hs +18/−0
- Obsidian/Exp.hs +227/−69
- Obsidian/Force.hs +106/−69
- Obsidian/Library.hs +431/−123
- Obsidian/Memory.hs +148/−14
- Obsidian/Mutable.hs +41/−27
- Obsidian/Names.hs +6/−0
- Obsidian/Program.hs +86/−108
- Obsidian/Run/CUDA/Exec.hs +289/−52
- Obsidian/SeqLoop.hs +82/−13
- Obsidian/Types.hs +26/−4
Obsidian.cabal view
@@ -1,5 +1,5 @@ Name: Obsidian-Version: 0.1.0.0+Version: 0.4.0.0 License: BSD3@@ -8,7 +8,7 @@ Maintainer: Joel Svensson<svenssonjoel@yahoo.se> Author: Joel Svensson<svenssonjoel@yahoo.se> -Copyright: Copyright (c) 2011-2014 Joel Svensson +Copyright: Copyright (c) 2011-2015 Joel Svensson Synopsis: Embedded language for GPU Programming HomePage: https://github.com/svenssonjoel/Obsidian@@ -18,7 +18,7 @@ Category: Language Cabal-Version: >=1.8-Tested-With: GHC == 7.6.1+Tested-With: GHC == 7.8.3 build-type: Simple @@ -38,14 +38,16 @@ , rdtsc == 1.3.0.0 -- Less-stable packages, upper bounds on next major version: build-depends: value-supply >= 0.6 && < 0.7- , language-c-quote >= 0.7.2 && < 0.8+ , language-c-quote >= 0.10.1.3 && < 0.11+-- >= 0.7.2 && < 0.9 , mainland-pretty >= 0.2.6 && < 0.3 , mwc-random >= 0.13.1.1 && < 0.14- , cuda >= 0.5.1.1 && < 0.6+ , cuda >= 0.5.1.1 && < 0.7 exposed-modules: Obsidian , Obsidian.CodeGen.CUDA- , Obsidian.Run.CUDA.Exec + , Obsidian.Run.CUDA.Exec+ other-modules: Obsidian.Array , Obsidian.Atomic@@ -60,9 +62,10 @@ , Obsidian.Program , Obsidian.SeqLoop , Obsidian.Types + , Obsidian.Data , Obsidian.CodeGen.CompileIM , Obsidian.CodeGen.Liveness- , Obsidian.CodeGen.Memory+ , Obsidian.CodeGen.Memory2 , Obsidian.CodeGen.Program , Obsidian.CodeGen.Reify
Obsidian.hs view
@@ -1,32 +1,34 @@ module Obsidian (module Obsidian.Array, module Obsidian.Program, module Obsidian.Exp, - module Obsidian.Types, module Obsidian.Force, module Obsidian.Library,- module Obsidian.CodeGen.Reify,- module Obsidian.CodeGen.CompileIM,- module Obsidian.CodeGen.CUDA, module Obsidian.Atomic, module Obsidian.SeqLoop, module Obsidian.Memory, - module Obsidian.Names,- module Obsidian.Mutable) where-+ module Obsidian.Mutable,+ module Obsidian.Data) where +-- These are internal. +-- module Obsidian.Types, +-- module Obsidian.CodeGen.Reify,+-- module Obsidian.CodeGen.CompileIM,+-- module Obsidian.CodeGen.CUDA,+-- module Obsidian.Names, import Obsidian.Program import Obsidian.Exp-import Obsidian.Types import Obsidian.Array import Obsidian.Library import Obsidian.Force-import Obsidian.CodeGen.Reify-import Obsidian.CodeGen.CompileIM-import Obsidian.CodeGen.CUDA import Obsidian.Atomic import Obsidian.SeqLoop import Obsidian.Memory-import Obsidian.Names import Obsidian.Mutable+import Obsidian.Data +--import Obsidian.Types+--import Obsidian.CodeGen.Reify+--import Obsidian.CodeGen.CompileIM+--import Obsidian.CodeGen.CUDA+--import Obsidian.Names
Obsidian/Array.hs view
@@ -1,41 +1,45 @@ {-# LANGUAGE MultiParamTypeClasses, FlexibleInstances,- GADTs #-} + GADTs #-}+{-# LANGUAGE TypeOperators #-}+{-# LANGUAGE FlexibleContexts #-} -{- Joel Svensson 2012+{- Joel Svensson 2012..2015 Notes:- 2014-04-08: Experimenting with API + 2014-04-08: Experimenting with API ---- OUTDATED ---- 2013-08-26: Experimenting with warp programs. These do not fit that well in established Idioms!- TODO: Improve this situation. + TODO: Improve this situation. ---- OUTDATED ---- 2013-01-08: Removed number-of-blocks field from Distribs- 2012-12-10: Drastically shortened. + 2012-12-10: Drastically shortened. -} -module Obsidian.Array (Pull, Push, SPull, DPull, SPush, DPush,- Pushable, - mkPull,- mkPush,- push,- setSize,- (!),- (<:),- Array(..),- ArrayLength(..),- ASize(..),- namedGlobal,- undefinedGlobal) where+module Obsidian.Array (Pull, Push, SPull, DPull, SPush, DPush+ , pushApp+ , mkPull+ , mkPush+ , push+ , pushThread+ , pushWarp+ , pushBlock+ , setSize+ , (!)+ , (<:)+ , Array(..)+ , ArrayLength(..)+ , ASize(..)+ , namedGlobal+ , undefinedGlobal+ ) where -import Obsidian.Exp -import Obsidian.Types-import Obsidian.Globs+import Obsidian.Exp as E import Obsidian.Program+import Obsidian.Globs -import Prelude hiding (replicate) -import Data.List hiding (replicate) +import Prelude hiding (replicate) import Data.Word ---------------------------------------------------------------------------@@ -45,13 +49,15 @@ type DPull = Pull EWord32 type SPush t a = Push t Word32 a-type DPush t a = Push t EWord32 a +type DPush t a = Push t EWord32 a --------------------------------------------------------------------------- -- Create arrays --------------------------------------------------------------------------- -- | An undefined array. Use as placeholder when generating code-undefinedGlobal n = Pull n $ \gix -> undefined--- | A named global array. +undefinedGlobal :: ASize s => s -> Pull s a+undefinedGlobal n = Pull n $ \_ -> undefined+-- | A named global array.+namedGlobal :: (ASize s, Scalar a) => Name -> s -> Pull s (Exp a) namedGlobal name n = Pull n $ \gix -> index name gix -- namedPull name n = Pull n $ \gix -> index name gix @@ -66,176 +72,372 @@ sizeConv = fromIntegral instance ASize (Exp Word32) where- sizeConv = id + sizeConv = id --------------------------------------------------------------------------- -- Push and Pull arrays --------------------------------------------------------------------------- -- | Push array. Parameterised over Program type and size type.-data Push p s a =- Push s ((a -> EWord32 -> TProgram ()) -> Program p ())+data Push t s a =+ Push s (PushFun t a) +type PushFun t a = Writer a -> Program t ()+type Writer a = a -> EWord32 -> TProgram ()+ -- | Pull array.-data Pull s a = Pull {pullLen :: s, +data Pull s a = Pull {pullLen :: s, pullFun :: EWord32 -> a} --- | Create a push array. +-- | Create a push array. mkPush :: s -> ((a -> EWord32 -> TProgram ()) -> Program t ()) -> Push t s a-mkPush n p = Push n p +mkPush n p = Push n p --- | Create a pull array. -mkPull n p = Pull n p +-- | Create a pull array.+mkPull :: s -> (EWord32 -> a) -> Pull s a+mkPull n p = Pull n p -- Fix this. -- * you cannot safely resize either push or pull arrays--- * you can shorten pull arrays safely. -setSize :: l -> Pull l a -> Pull l a+-- * you can shorten pull arrays safely.+setSize :: ASize l => l -> Pull l a -> Pull l a setSize n (Pull _ ixf) = mkPull n ixf ------------------------------------------------------------------------------- Array Class + ---------------------------------------------------------------------------+-- Array Class+---------------------------------------------------------------------------x class ArrayLength a where -- | Get the length of an array. len :: a s e -> s instance ArrayLength Pull where- len (Pull n ixf) = n+ len arr = pullLen arr instance ArrayLength (Push t) where- len (Push s p) = s+ len (Push s _) = s class Array a where -- | Array of consecutive integers iota :: ASize s => s -> a s EWord32- -- | Create an array by replicating an element. - replicate :: ASize s => s -> e -> a s e + -- | Create an array by replicating an element.+ replicate :: ASize s => s -> e -> a s e - -- | Map a function over an array. + -- | Map a function over an array. aMap :: (e -> e') -> a s e -> a s e'- -- | Perform arbitrary permutations (dangerous). + -- | Perform arbitrary permutations (dangerous). ixMap :: (EWord32 -> EWord32) -> a s e -> a s e- -- -- | Reduce an array using a provided operator. - -- fold1 :: (e -> e -> e) -> a Word32 e -> a Word32 e + -- requires Choice !+ -- Simply because the pull array implementation of it does.+ -- | Append two arrays.+ append :: (ASize s, Choice e) => a s e -> a s e -> a s e - -- would require Choice !- -- | Append two arrays. - append :: (ASize s, Choice e) => a s e -> a s e -> a s e - -- technicalities -- | Statically sized array to dynamically sized array. toDyn :: a Word32 e -> a EW32 e- -- | Dynamically sized array to statically sized array. - fromDyn :: Word32 -> a EW32 e -> a Word32 e - + -- | Dynamically sized array to statically sized array.+ fromDyn :: Word32 -> a EW32 e -> a Word32 e+ instance Array Pull where- iota s = Pull s $ \ix -> ix + iota s = Pull s $ \ix -> ix replicate s e = Pull s $ \_ -> e aMap f (Pull n ixf) = Pull n (f . ixf)- ixMap f (Pull n ixf) = Pull n (ixf . f) + ixMap f (Pull n ixf) = Pull n (ixf . f) append a1 a2 = Pull (n1+n2)- $ \ix -> ifThenElse (ix <* (sizeConv n1)) - (a1 ! ix) - (a2 ! (ix - (sizeConv n1)))- where + $ \ix -> ifThenElse (ix E.<* sizeConv n1)+ (a1 ! ix)+ (a2 ! (ix - sizeConv n1))+ where n1 = len a1- n2 = len a2 + n2 = len a2 -- technicalities toDyn (Pull n ixf) = Pull (fromIntegral n) ixf- fromDyn n (Pull _ ixf) = Pull n ixf - - -instance Array (Push t) where+ fromDyn n (Pull _ ixf) = Pull n ixf+++instance Array (Push Thread) where iota s = Push s $ \wf -> do- forAll (sizeConv s) $ \ix -> wf ix ix + forAll (sizeConv s) $ \ix -> wf ix ix replicate s e = Push s $ \wf -> do- forAll (sizeConv s) $ \ix -> wf e ix + forAll (sizeConv s) $ \ix -> wf e ix aMap f (Push s p) = Push s $ \wf -> p (\e ix -> wf (f e) ix) ixMap f (Push s p) = Push s $ \wf -> p (\e ix -> wf e (f ix)) - -- unfortunately a Choice constraint. + -- unfortunately a Choice constraint. append p1 p2 = Push (n1 + n2) $ \wf -> do p1 <: wf- p2 <: \a i -> wf a (sizeConv n1 + i) - where + p2 <: \a i -> wf a (sizeConv n1 + i)+ where n1 = len p1- n2 = len p2 + n2 = len p2 -- technicalities- toDyn (Push n p) = Push (fromIntegral n) p - fromDyn n (Push _ p) = Push n p - + toDyn (Push n p) = Push (fromIntegral n) p+ fromDyn n (Push _ p) = Push n p +instance Array (Push Warp) where+ iota s = Push s $ \wf ->+ do+ forAll (sizeConv s) $ \ix -> wf ix ix+ replicate s e = Push s $ \wf ->+ do+ forAll (sizeConv s) $ \ix -> wf e ix+ aMap f (Push s p) = Push s $ \wf -> p (\e ix -> wf (f e) ix)+ ixMap f (Push s p) = Push s $ \wf -> p (\e ix -> wf e (f ix))++ -- unfortunately a Choice constraint.+ append p1 p2 =+ Push (n1 + n2) $ \wf ->+ do p1 <: wf+ p2 <: \a i -> wf a (sizeConv n1 + i)+ where+ n1 = len p1+ n2 = len p2++ -- technicalities+ toDyn (Push n p) = Push (fromIntegral n) p+ fromDyn n (Push _ p) = Push n p+++instance Array (Push Block) where+ iota s = Push s $ \wf ->+ do+ forAll (sizeConv s) $ \ix -> wf ix ix+ replicate s e = Push s $ \wf ->+ do+ forAll (sizeConv s) $ \ix -> wf e ix+ aMap f (Push s p) = Push s $ \wf -> p (\e ix -> wf (f e) ix)+ ixMap f (Push s p) = Push s $ \wf -> p (\e ix -> wf e (f ix))++ -- unfortunately a Choice constraint.+ append p1 p2 =+ Push (n1 + n2) $ \wf ->+ do p1 <: wf+ p2 <: \a i -> wf a (sizeConv n1 + i)+ where+ n1 = len p1+ n2 = len p2++ -- technicalities+ toDyn (Push n p) = Push (fromIntegral n) p+ fromDyn n (Push _ p) = Push n p++instance Array (Push Grid) where+ iota s = error "iota: not supported as Grid"+ replicate s e = error "replicate: not supported as Grid"+ aMap f (Push s p) = Push s $ \wf -> p (\e ix -> wf (f e) ix)+ ixMap f (Push s p) = Push s $ \wf -> p (\e ix -> wf e (f ix))++ -- unfortunately a Choice constraint.+ append p1 p2 =+ mkPush (n1 + n2) $ \wf ->+ do p1 <: wf+ p2 <: \a i -> wf a (sizeConv n1 + i)+ where+ n1 = len p1+ n2 = len p2++ -- technicalities+ toDyn (Push n p) = Push (fromIntegral n) p+ fromDyn n (Push _ p) = Push n p++++ --------------------------------------------------------------------------- -- Functor instance Pull/Push arrays ----------------------------------------------------------------------------instance Array arr => Functor (arr w) where +instance Array (Push t) => Functor (Push t s) where fmap = aMap +instance Functor (Pull s) where+ fmap = aMap++ --------------------------------------------------------------------------- -- Pushable ----------------------------------------------------------------------------class Pushable t where- push :: ASize s => Pull s e -> Push t s e +-- | Convert a pull array to a push array.+push :: (t *<=* Block) => ASize s => Pull s e -> Push t s e+push (Pull n ixf) =+ mkPush n $ \wf ->+ forAll (sizeConv n) $ \i -> wf (ixf i) i -instance Pushable Thread where- push (Pull n ixf) =- Push n $ \wf -> seqFor (sizeConv n) $ \i -> wf (ixf i) i+-- Keep push, user may need to annotate with a type+-- Add specific+-- pushThread+-- pushWarp+-- pushBlock+pushThread :: ASize s => Pull s e -> Push Thread s e+pushThread = push -instance Pushable Warp where- push (Pull n ixf) =- Push n $ \wf ->- forAll (sizeConv n) $ \i -> wf (ixf i) i+pushWarp :: ASize s => Pull s e -> Push Warp s e+pushWarp = push -instance Pushable Block where- push (Pull n ixf) =- Push n $ \wf ->- forAll (sizeConv n) $ \i -> wf (ixf i) i+pushBlock :: ASize s => Pull s e -> Push Block s e+pushBlock = push -instance Pushable Grid where- push (Pull n ixf) =- Push n $ \wf ->- forAll (sizeConv n) $ \i -> wf (ixf i) i - ++------------------------------------------------------------+-- I think these are strange and should go away.+-- They are always id for Push arrays and push for pull arrays.+-- Most instances are nonsense. which is Weird.+--+-- class AsGrid a where+-- asGrid :: ASize s => a s e -> Push Grid s e++-- instance AsGrid Pull where+-- asGrid = error $ "asGrid: Trying to compute a pull array as a grid.\n" +++-- "This operation is not supported, there are too many choises\n" +++-- "involved that we just should not make automatically"+-- instance AsGrid (Push Grid) where+-- asGrid = id++-- instance AsGrid (Push Block) where+-- asGrid _ = error $ "asGrid: Trying to convert a Block to a Grid.\n" +++-- "This operation is not supported!"++-- instance AsGrid (Push Warp) where+-- asGrid _ = error $ "asGrid: Trying to convert a Warp to a Grid.\n" +++-- "This operation is not supported!"++-- instance AsGrid (Push Thread) where+-- asGrid _ = error $ "asGrid: Trying to convert a Thread to a Grid.\n" +++-- "This operation is not supported!"+++-- class AsBlock a where+-- asBlock :: a Word32 e -> SPush Block e++-- instance AsBlock Pull where+-- asBlock = push++-- instance AsBlock (Push Block) where+-- asBlock = id++-- instance AsBlock (Push Grid) where+-- asBlock _ = error $ "asBlock: Trying to convert a Grid to a Block.\n" +++-- "This operation is not supported!"++-- instance AsBlock (Push Warp) where+-- asBlock _ = error $ "asBlock: Trying to convert a Thread to a Block.\n" +++-- "This operation is not supported!"++-- instance AsBlock (Push Thread) where+-- asBlock _ = error $ "asBlock: Trying to convert a Thread to a Block.\n" +++-- "This operation is not supported!"+++-- class AsWarp a where+-- asWarp :: a Word32 e -> SPush Warp e++-- instance AsWarp Pull where+-- asWarp = push++-- instance AsWarp (Push Warp) where+-- asWarp = id++-- instance AsWarp (Push Grid) where+-- asWarp _ = error $ "asWarp: Trying to convert a Grid to a Warp.\n" +++-- "This operation is not supported!"++-- instance AsWarp (Push Block) where+-- asWarp _ = error $ "asWarp: Trying to convert a Block to a Warp.\n" +++-- "This operation is not supported!"++-- instance AsWarp (Push Thread) where+-- asWarp _ = error $ "asBlock: Trying to convert a Thread to a Warp.\n" +++-- "This operation is not supported!"+++-- class AsThread a where+-- asThread :: a Word32 e -> SPush Thread e++-- instance AsThread Pull where+-- asThread = push++-- instance AsThread (Push Thread) where+-- asThread = id++-- instance AsThread (Push Grid) where+-- asThread _ = error $ "asThread: Trying to convert a Grid to a Thread.\n" +++-- "This operation is not supported!"++-- instance AsThread (Push Block) where+-- asThread _ = error $ "asThread: Trying to convert a Block to a Thread.\n" +++-- "This operation is not supported!"++-- instance AsThread (Push Warp) where+-- asThread _ = error $ "asThread: Trying to convert a Warp to a Thread.\n" +++-- "This operation is not supported!"++++++-- class Pushable t where+-- push :: ASize s => Pull s e -> Push t s e++-- instance Pushable Thread where+-- push (Pull n ixf) =+-- Push n $ \wf -> seqFor (sizeConv n) $ \i -> wf (ixf i) i++-- instance Pushable Warp where+-- push (Pull n ixf) =+-- Push n $ \wf ->+-- forAll (sizeConv n) $ \i -> wf (ixf i) i++-- instance Pushable Block where+-- push (Pull n ixf) =+-- Push n $ \wf ->+-- forAll (sizeConv n) $ \i -> wf (ixf i) i++-- instance Pushable Grid where+-- push (Pull n ixf) =+-- Push n $ \wf ->+-- forAll (sizeConv n) $ \i -> wf (ixf i) i+ -- class PushableN t where -- pushN :: ASize s => Word32 -> Pull s e -> Push t s e -- instance PushableN Block where -- pushN n (Pull m ixf) = -- Push m $ \ wf -> forAll (sizeConv (m `div` fromIntegral n)) $ \tix ->--- warpForAll 1 $ \_ -> +-- warpForAll 1 $ \_ -> -- seqFor (fromIntegral n) $ \ix -> wf (ixf (tix * fromIntegral n + ix))--- (tix * fromIntegral n + ix) - +-- (tix * fromIntegral n + ix)+ -- instance PushableN Grid where -- pushN n (Pull m ixf) = -- Push m $ \ wf -> forAll (sizeConv (m `div` fromIntegral n)) $ \bix -> -- forAll (fromIntegral n) $ \tix -> wf (ixf (bix * fromIntegral n + tix))--- (bix * fromIntegral n + tix) +-- (bix * fromIntegral n + tix) -------------------------------------------------------------------------- -- Indexing, array creation. ----------------------------------------------------------------------------+pushApp :: Push t s a -> (a -> EWord32 -> TProgram ()) -> Program t () pushApp (Push _ p) a = p a infixl 9 <: (<:) :: Push t s a -> (a -> EWord32 -> Program Thread ()) -> Program t ()-(<:) = pushApp +(<:) = pushApp -infixl 9 ! -(!) :: Pull s e -> Exp Word32 -> e -(!) arr = pullFun arr +infixl 9 !+(!) :: Pull s e -> Exp Word32 -> e+(!) arr = pullFun arr +---------------------------------------------------------------------------+--+---------------------------------------------------------------------------
Obsidian/Atomic.hs view
@@ -1,3 +1,4 @@+ {-# LANGUAGE GADTs #-} {- Joel Svensson,
Obsidian/CodeGen/CUDA.hs view
@@ -1,12 +1,18 @@ +{- Joel Svensson 2013..2015 -module Obsidian.CodeGen.CUDA (genKernel) where + Notes:+ 25-Nov-2014: Making changes regarding memory allocation+ +-} +module Obsidian.CodeGen.CUDA (genKernel, genKernelParams, SharedMemConfig(..),ToProgram(..)) where + import Obsidian.CodeGen.Reify import Obsidian.CodeGen.CompileIM import Obsidian.CodeGen.Liveness-import Obsidian.CodeGen.Memory+import Obsidian.CodeGen.Memory2 import Text.PrettyPrint.Mainland import qualified Data.Map as M@@ -15,16 +21,44 @@ -- Generate CUDA kernels --------------------------------------------------------------------------- +-- | Generates kernel C code as a String+-- while assuming there is 48KB of shared mem in 32 banks genKernel :: ToProgram prg => Word32 -> String -> prg -> String-genKernel nt kn prg = prgStr+genKernel = genKernelParams sm_conf+ where + -- pretend we have 32 banks and 48kb shared mem+ sm_conf :: SharedMemConfig + sm_conf = SharedMemConfig 49152 32 True++++-- | Generates kernel C code as a String,+-- Programmer passes in the Shared memory configuration+genKernelParams :: ToProgram prg+ => SharedMemConfig+ -> Word32+ -> String+ -> prg+ -> String+genKernelParams sm_conf nt kn prg = prgStr where- prgStr = pretty 75 $ ppr $ compile PlatformCUDA (Config nt bytesShared) kn (a,rim) + prgStr = pretty 75+ $ ppr+ $ compileDeclsTop+ (Config nt bytesShared)+ name_loc+ kn (a,im) + -- $ compile PlatformCUDA+ -- (Config nt bytesShared)+ -- kn (a,rim) (a,im) = toProgram_ 0 prg iml = computeLiveness im- (m,mm) = mmIM iml sharedMem (M.empty)- bytesShared = size m - rim = renameIM mm iml+ (m,mm) = memMapIM sm_conf iml (M.empty)+ bytesShared = size m+ name_loc = M.assocs mm+ --rim = renameIM mm iml+
Obsidian/CodeGen/CompileIM.hs view
@@ -5,50 +5,41 @@ {- - Joel Svensson 2013-+ Joel Svensson 2013..2015 -} -module Obsidian.CodeGen.CompileIM where --import Language.C.Quote hiding (Block)-import Language.C.Quote.CUDA as CU+module Obsidian.CodeGen.CompileIM where -import qualified Language.C.Quote.OpenCL as CL +import Language.C.Quote.CUDA hiding (Block) -import qualified "language-c-quote" Language.C.Syntax as C +import qualified "language-c-quote" Language.C.Syntax as C import Obsidian.Exp (IExp(..),IBinOp(..),IUnOp(..))-import Obsidian.Types+import Obsidian.Types as T import Obsidian.DimSpec import Obsidian.CodeGen.Program -import Control.Monad.State- import Data.Word-import Data.Int -{- TODOs:- - * Pass a target "platform" to code generator.- - CUDA- - OpenCL- - Sequential C- - C with OpenMP ? - * rewrite some functions here to use a reader monad. +{- Notes: + * Generate only CUDA + * TODOs+ - Currently inserts some pointless code+ The values of tid,wid etc are updated before certain code blocks+ and reset to a standard value after, often to just be updated again to+ something else. It may result in prettier code if all code blocks+ that use a variable sets it to what it should be rather than expecting "standard".+ *** probably has no effect on performance *** - * TODO: Make sure tid always has correct Value + -} ------------------------------------------------------------------------------ Platform+-- Config ----------------------------------------------------------------------------data Platform = PlatformCUDA- | PlatformOpenCL- | PlatformC data Config = Config { configThreadsPerBlock :: Word32, configSharedMem :: Word32}@@ -59,26 +50,26 @@ --------------------------------------------------------------------------- -- compileExp (maybe a bad name) ----------------------------------------------------------------------------compileExp :: IExp -> Exp +compileExp :: IExp -> C.Exp compileExp (IVar name t) = [cexp| $id:name |] -- TODO: Fix all this! -compileExp (IBlockIdx X) = [cexp| $id:("bid")|] -- [cexp| $id:("blockIdx.x") |]-compileExp (IBlockIdx Y) = [cexp| $id:("blockIdx.y") |]-compileExp (IBlockIdx Z) = [cexp| $id:("blockIdx.z") |]+-- compileExp (IBlockIdx X) = [cexp| $id:("bid")|] -- [cexp| $id:("blockIdx.x") |]+-- compileExp (IBlockIdx Y) = [cexp| $id:("blockIdx.y") |]+-- compileExp (IBlockIdx Z) = [cexp| $id:("blockIdx.z") |] -compileExp (IThreadIdx X) = [cexp| $id:("threadIdx.x") |]-compileExp (IThreadIdx Y) = [cexp| $id:("threadIdx.y") |]-compileExp (IThreadIdx Z) = [cexp| $id:("threadIdx.z") |]+-- compileExp (IThreadIdx X) = [cexp| $id:("threadIdx.x") |]+-- compileExp (IThreadIdx Y) = [cexp| $id:("threadIdx.y") |]+-- compileExp (IThreadIdx Z) = [cexp| $id:("threadIdx.z") |] -compileExp (IBlockDim X) = [cexp| $id:("blockDim.x") |]-compileExp (IBlockDim Y) = [cexp| $id:("blockDim.y") |]-compileExp (IBlockDim Z) = [cexp| $id:("blockDim.z") |]+-- compileExp (IBlockDim X) = [cexp| $id:("blockDim.x") |]+-- compileExp (IBlockDim Y) = [cexp| $id:("blockDim.y") |]+-- compileExp (IBlockDim Z) = [cexp| $id:("blockDim.z") |] -compileExp (IGridDim X) = [cexp| $id:("GridDim.x") |]-compileExp (IGridDim Y) = [cexp| $id:("GridDim.y") |]-compileExp (IGridDim Z) = [cexp| $id:("GridDim.z") |]+-- compileExp (IGridDim X) = [cexp| $id:("GridDim.x") |]+-- compileExp (IGridDim Y) = [cexp| $id:("GridDim.y") |]+-- compileExp (IGridDim Z) = [cexp| $id:("GridDim.z") |] compileExp (IBool True) = [cexp|1|] compileExp (IBool False) = [cexp|0|]@@ -95,8 +86,41 @@ compileExp (IFloat n) = [cexp| $float:(toRational n) |] compileExp (IDouble n) = [cexp| $double:(toRational n) |] -compileExp (IIndex (i1,[e]) t) = [cexp| $(compileExp i1)[$(compileExp e)] |] +-- Implementing these may be a bit awkward+-- given there are no vector literals in cuda. +compileExp (IFloat2 n m) = error "IFloat2 unhandled"+compileExp (IFloat3 n m l) = error "IFloat3 unhandled"+compileExp (IFloat4 n m l k) = error "IFloat4 unhandled"+compileExp (IDouble2 n m) = error "IDouble2 unhandled" +compileExp (IInt8_2 n m) = error "FIXME"+compileExp (IInt8_3 n m k) = error "FIXME"+compileExp (IInt8_4 n m k l) = error "FIXME"+compileExp (IInt16_2 n m ) = error "FIXME"+compileExp (IInt16_3 n m k) = error "FIXME"+compileExp (IInt16_4 n m k l) = error "FIXME"+compileExp (IInt32_2 n m) = error "FIXME"+compileExp (IInt32_3 n m k) = error "FIXME"+compileExp (IInt32_4 n m k l) = error "FIXME"+compileExp (IInt64_2 n m) = error "FIXME"+compileExp (IInt64_3 n m k) = error "FIXME"+compileExp (IInt64_4 n m k l) = error "FIXME"+compileExp (IWord8_2 n m) = error "FIXME"+compileExp (IWord8_3 n m k) = error "FIXME"+compileExp (IWord8_4 n m k l) = error "FIXME"+compileExp (IWord16_2 n m ) = error "FIXME"+compileExp (IWord16_3 n m k) = error "FIXME"+compileExp (IWord16_4 n m k l) = error "FIXME"+compileExp (IWord32_2 n m) = error "FIXME"+compileExp (IWord32_3 n m k) = error "FIXME"+compileExp (IWord32_4 n m k l) = error "FIXME"+compileExp (IWord64_2 n m) = error "FIXME"+compileExp (IWord64_3 n m k) = error "FIXME"+compileExp (IWord64_4 n m k l) = error "FIXME" ++compileExp (IIndex (i1,[e]) t) = [cexp| $(compileExp i1)[$(compileExp e)] |]+compileExp a@(IIndex (_,_) _) = error $ "compileExp: Malformed index expression " ++ show a+ compileExp (ICond e1 e2 e3 t) = [cexp| $(compileExp e1) ? $(compileExp e2) : $(compileExp e3) |] compileExp (IBinOp op e1 e2 t) = go op @@ -118,7 +142,8 @@ go IOr = [cexp| $x || $y |] go IPow = case t of Float -> [cexp|powf($x,$y) |]- Double -> [cexp|pow($x,$y) |] + Double -> [cexp|pow($x,$y) |]+ _ -> error $ "IPow applied at wrong type" go IBitwiseAnd = [cexp| $x & $y |] go IBitwiseOr = [cexp| $x | $y |] go IBitwiseXor = [cexp| $x ^ $y |]@@ -129,6 +154,10 @@ x = compileExp e go IBitwiseNeg = [cexp| ~$x|] go INot = [cexp| !$x|]+ go IGetX = [cexp| $x.x|]+ go IGetY = [cexp| $x.y|]+ go IGetZ = [cexp| $x.z|]+ go IGetW = [cexp| $x.w|] compileExp (IFunCall name es t) = [cexp| $fc |] where@@ -138,7 +167,8 @@ compileExp (ICast e t) = [cexp| ($ty:(compileType t)) $e' |] where e' = compileExp e- ++compileType :: T.Type -> C.Type compileType (Int8) = [cty| typename int8_t |] compileType (Int16) = [cty| typename int16_t |] compileType (Int32) = [cty| typename int32_t |]@@ -149,14 +179,47 @@ compileType (Word64) = [cty| typename uint64_t |] compileType (Float) = [cty| float |] compileType (Double) = [cty| double |]++compileType (Vec2 Float) = [cty| float4|]+compileType (Vec3 Float) = [cty| float3|]+compileType (Vec4 Float) = [cty| float2|]++compileType (Vec2 Double) = [cty| double2|]++-- How does this interplay with my use of uint8_t etc. Here it is char!+compileType (Vec2 Int8) = [cty| char2|]+compileType (Vec3 Int8) = [cty| char3|]+compileType (Vec4 Int8) = [cty| char4|]+ +compileType (Vec2 Int16) = [cty| short2|]+compileType (Vec3 Int16) = [cty| short3|]+compileType (Vec4 Int16) = [cty| short4|]++compileType (Vec2 Int32) = [cty| int2|]+compileType (Vec3 Int32) = [cty| int3|]+compileType (Vec4 Int32) = [cty| int4|]++compileType (Vec2 Word8) = [cty| uchar2|]+compileType (Vec3 Word8) = [cty| uchar3|]+compileType (Vec4 Word8) = [cty| uchar4|]+ +compileType (Vec2 Word16) = [cty| ushort2|]+compileType (Vec3 Word16) = [cty| ushort3|]+compileType (Vec4 Word16) = [cty| ushort4|]++compileType (Vec2 Word32) = [cty| uint2|]+compileType (Vec3 Word32) = [cty| uint3|]+compileType (Vec4 Word32) = [cty| uint4|]+++compileType (Shared t) = [cty| __shared__ $ty:(compileType t) |] compileType (Pointer t) = [cty| $ty:(compileType t)* |]-compileType (Volatile t) = [cty| volatile $ty:(compileType t)|] +compileType (Volatile t) = [cty| volatile $ty:(compileType t)|]+compileType t = error $ "compileType: Not implemented " ++ show t ---compileType' (Volatile t) = [cty| volatile $ty:(compileType t)|] ---compileType' t = compileType t {--+ *** THIS IS SOLVED ! *** Solve the volatile issue. When operating in a warp without syncs Add Volatile to pointers and only there!@@ -180,74 +243,57 @@ -} ------------------------------------------------------------------------------ ** --- Compile IM--- ** ---------------------------------------------------------------------------- ---- newtype CInfo a = CInfo (State CInfoState a)--- deriving (Monad, MonadState CInfoState) ---- data CInfoState = CInfoState { cInfoTid :: (Bool,Exp),--- cInfoWarpID :: (Bool,Exp),--- cInfoWarpIx :: (Bool,Exp) } --- evalCInfo (CInfo s) = evalState s ----------------------------------------------------------------------------- -- Statement t to Stm --------------------------------------------------------------------------- --compileStm :: Platform -> Config -> Statement t -> [Stm]-compileStm p c (SAssign name [] e) =+compileStm :: Config -> Statement t -> [C.Stm]+compileStm c (SAssign name [] e) = [[cstm| $(compileExp name) = $(compileExp e);|]]-compileStm p c (SAssign name [ix] e) = +compileStm c (SAssign name [ix] e) = [[cstm| $(compileExp name)[$(compileExp ix)] = $(compileExp e); |]]-compileStm p c (SAtomicOp name ix atop) = +compileStm c (SAtomicOp name ix atop) = case atop of AtInc -> [[cstm| atomicInc(&$(compileExp name)[$(compileExp ix)],0xFFFFFFFF); |]] AtAdd e -> [[cstm| atomicAdd(&$(compileExp name)[$(compileExp ix)],$(compileExp e));|]] AtSub e -> [[cstm| atomicSub(&$(compileExp name)[$(compileExp ix)],$(compileExp e));|]] AtExch e -> [[cstm| atomicExch(&$(compileExp name)[$(compileExp ix)],$(compileExp e));|]] -compileStm p c (SCond be im) = [[cstm| if ($(compileExp be)) { $stms:body } |]]+compileStm c (SCond be im) = [[cstm| if ($(compileExp be)) { $stms:body } |]] where - body = compileIM p c im -- (compileIM p c im)-compileStm p c (SSeqFor loopVar n im) = + body = compileIM c im -- (compileIM p c im)+compileStm c (SSeqFor loopVar n im) = [[cstm| for (int $id:loopVar = 0; $id:loopVar < $(compileExp n); ++$id:loopVar) { $stms:body } |]] where- body = compileIM p c im -- (compileIM p c im)+ body = compileIM c im -- (compileIM p c im) -- Just relay to specific compileFunction-compileStm p c a@(SForAll lvl n im) = compileForAll p c a+compileStm c a@(SForAll lvl n im) = compileForAll c a -compileStm p c a@(SDistrPar lvl n im) = compileDistr p c a +compileStm c a@(SDistrPar lvl n im) = compileDistr c a -compileStm p c (SSeqWhile b im) =+compileStm c (SSeqWhile b im) = [[cstm| while ($(compileExp b)) { $stms:body}|]] where- body = compileIM p c im + body = compileIM c im -compileStm p c SSynchronize - = case p of- PlatformCUDA -> [[cstm| __syncthreads(); |]]- PlatformOpenCL -> [[cstm| barrier(CLK_LOCAL_MEM_FENCE); |]]+compileStm c SSynchronize = [[cstm| __syncthreads(); |]]+ -compileStm _ _ (SAllocate _ _ _) = []-compileStm _ _ (SDeclare name t) = []+compileStm _ (SAllocate _ _ _) = []+compileStm _ (SDeclare name t) = [] -compileStm _ _ a = error $ "compileStm: missing case "+compileStm _ a = error $ "compileStm: missing case " --------------------------------------------------------------------------- -- DistrPar ----------------------------------------------------------------------------compileDistr :: Platform -> Config -> Statement t -> [Stm] -compileDistr PlatformCUDA c (SDistrPar Block n im) = codeQ ++ codeR+compileDistr :: Config -> Statement t -> [C.Stm] +compileDistr c (SDistrPar Block n im) = codeQ ++ codeR -- New here is BLOCK virtualisation where- cim = compileIM PlatformCUDA c im+ cim = compileIM c im -- ++ [[cstm| __syncthreads();|]] numBlocks = [cexp| $id:("gridDim.x") |] @@ -269,13 +315,13 @@ -- I must look over the functions that can potentially create this IM. -- Can make a separate case for unknown 'n' but generate worse code. -- (That is true for all levels) -compileDistr PlatformCUDA c (SDistrPar Warp (IWord32 n) im) = codeQ ++ codeR +compileDistr c (SDistrPar Warp (IWord32 n) im) = codeQ ++ codeR -- Here the 'im' should be distributed over 'n'warps. -- 'im' uses a warpID variable to identify what warp it is. -- 'n' may be higher than the actual number of warps we have! -- So GPU warp virtualisation is needed. where- cim = compileIM PlatformCUDA c im+ cim = compileIM c im nWarps = fromIntegral $ configThreadsPerBlock c `div` 32 numWarps = [cexp| $int:nWarps|] @@ -297,17 +343,17 @@ [cstm| __syncthreads();|]] ------------------------------------------------------------------------------ ForAll is compiled differently for different platforms+-- ForAll ----------------------------------------------------------------------------compileForAll :: Platform -> Config -> Statement t -> [Stm]-compileForAll PlatformCUDA c (SForAll Warp (IWord32 n) im) = codeQ ++ codeR+compileForAll :: Config -> Statement t -> [C.Stm]+compileForAll c (SForAll Warp (IWord32 n) im) = codeQ ++ codeR where nt = 32 q = n `div` nt r = n `mod` nt - cim = compileIM PlatformCUDA c im + cim = compileIM c im codeQ = case q of@@ -315,7 +361,7 @@ 1 -> cim n -> [[cstm| for ( int vw = 0; vw < $int:q; ++vw) { $stms:body } |], [cstm| $id:("warpIx") = threadIdx.x % 32; |]]- -- [cstm| __syncthreads();|]]+ --[cstm| __syncthreads();|]] where body = [cstm|$id:("warpIx") = vw*$int:nt + (threadIdx.x % 32); |] : cim --body = [cstm|$id:("warpIx") = (threadIdx.x % 32) * q + vw; |] : cim@@ -329,9 +375,9 @@ -- [cstm| __syncthreads();|], [cstm| $id:("warpIx") = threadIdx.x % 32; |]] -compileForAll PlatformCUDA c (SForAll Block (IWord32 n) im) = goQ ++ goR +compileForAll c (SForAll Block (IWord32 n) im) = goQ ++ goR where- cim = compileIM PlatformCUDA c im+ cim = compileIM c im -- ++ [[cstm| __syncthreads();|]] nt = configThreadsPerBlock c @@ -348,7 +394,9 @@ -- stm <- updateTid [cexp| threadIdx.x |] -- return $ [cstm| $id:loopVar = threadIdx.x; |] : cim n -> [[cstm| for ( int i = 0; i < $int:q; ++i) { $stms:body } |], + -- __syncthreads(); } |], [cstm| $id:("tid") = threadIdx.x; |]]+ -- [cstm| __syncthreads();|]] where body = [cstm|$id:("tid") = i*$int:nt + threadIdx.x; |] : cim @@ -365,166 +413,67 @@ $stms:cim } |], [cstm| $id:("tid") = threadIdx.x; |]] -compileForAll PlatformCUDA c (SForAll Grid n im) = cim+compileForAll c (SForAll Grid n im) = error "compileForAll: Grid" -- cim -- The grid case is special. May need more thought- --error "compileForAll: Grid"- where- cim = compileIM PlatformCUDA c im--compileForAll PlatformC c (SForAll lvl (IWord32 n) im) = go+ -- + -- The problem with this case is that+ -- I need to come up with a blocksize (but without any guidance)+ -- from the programmer.+ -- Though! There is no way the programmer could provide any+ -- such info ... where- body = compileIM PlatformC c im - go = [ [cstm| for (int i = 0; i <$int:n; ++i) { $stms:body } |] ] - + cim = compileIM c im ------------------------------------------------------------------------------- compileWarp (needs fixing so that warpIx and warpID are always correct) ------------------------------------------------------------------------------- compileWarp :: Platform -> Config -> IExp -> IMList t -> [Stm]--- compileWarp PlatformCUDA c (IWord32 warps) im = --- concatMap (go . fst) im+-- compileForAll PlatformC c (SForAll lvl (IWord32 n) im) = go -- where--- go (SAllocate nom n t) = [] -- Skip allocations at this point. Been handled earlier. --- go (SForAll Warp (IWord32 n) im) =--- {---- warps : number of warps that should execute this code--- n : number of threads per each of these warps---- warps may be a larger number than there are actual warps--- n map be a larger number than 32 (the actual number of threads per warp)---- c : is the configuration. It knows the actual number of threads--- we are generating code for. --- -} --- if (wholeRealWarps <= 0)--- then error "compileWarp: Atleast one full warp of real threads needed!"--- else --- case (wholeRealWarps `compare` warps) of--- {---- 3 potential outcomes.--- #1 we have more real warps than we need--- #2 we have exactly the amount of warps that we need--- #3 we have fewer real warps than requested.--- -} --- GT -> [[cstm| if (threadIdx.x < $int:(warps*32)) {--- $stms:(goQ ++ goR) } |]] --- EQ -> goQR--- LT -> wQ ++ wR - --- where--- cim = compileIM PlatformCUDA c im- --- nt = configThreadsPerBlock c--- wholeRealWarps = nt `quot` 32--- threadsPartialWarp = nt `rem` 32 -- Do not use partial warps! - --- -- Maybe something else than s for goR--- -- It will depend on how --- goQR = goQ ++ goR - --- threadQ = n `quot` 32 -- Set up virtual threads within warp--- threadR = n `rem` 32 -- ----- ###################################################################### --- CLEAN THIS MESS UP ! !!! ! ! ! !! ! !!!! ! !! ! ! ! ! ! ! !!! ! ! ! !--- ######################################################################--- -- -- Compile for warps, potentially with virtual threads - -- goQ :: [Stm]- -- goQ = case threadQ of - -- 0 -> [] - -- 1 -> cim - -- n -> [[cstm| for (int i = 0; i < $int:threadQ; ++i) {- -- $stms:body } |],- -- [cstm| $id:("warpIx") = threadIdx.x % 32; |]]- -- where - -- body = [cstm| $id:("warpIx") = i*32 + threadIdx.x % 32; |] : cim- -- goR :: [Stm] - -- goR = case threadR of - -- 0 -> [] - -- _ -> [[cstm| if ( threadIdx.x % 32 < $int:threadR) { - -- $id:("warpIx") = $int:(threadQ*32) + (threadIdx.x % 32);- -- $stms:cim } |],- -- [cstm| $id:("warpIx") = threadIdx.x % 32;|]]-- -- -- Compile for virtual warps - -- warpQ = warps `quot` wholeRealWarps -- wholeRealWarps may be zero! - -- warpR = warps `rem` wholeRealWarps- -- -- each warp will pretend to be warpQ number of warps. - - -- -- warp jump size - -- -- warpQ is the number of virtual warps that EACH real warp - -- -- will pretend to be. - - -- -- warpR is a number of leftover warps (less than wholeRealWarps) - -- -- that need to be processed afterwards. - - -- wQ = case warpQ of - -- 0 -> [] - -- 1 -> goQR - -- n -> [[cstm| for (int vw = 0; vw < $int:warpQ; ++vw) { $stms:body } |],- -- [cstm| $id:("warpID") = threadIdx.x / 32; |]]- -- where - -- --- vw * wholeRealWarps is incorrect. - -- --- - -- body = [cstm| $id:("warpID") = vw + ((threadIdx.x / 32)*$int:warpQ); |] : goQR+-- body = compileIM PlatformC c im +-- go = [ [cstm| for (int i = 0; i <$int:n; ++i) { $stms:body } |] ] + - -- wR = case warpR of - -- 0 -> [] - -- -- Skip one division, - -- -- Premature optimisation? CUDA compiler will share the threadIdx.x / 32 result. - -- n -> [[cstm| if (threadIdx.x < $int:(n * 32)) {- -- $id:("warpID") = $int:(warpQ*wholeRealWarps) + threadIdx.x / 32; - -- $stms:(goQR) } |], -- cim- -- [cstm| $id:("warpID") = threadIdx.x / 32;|]] --------------------------------------------------------------------------- -- CompileIM to list of Stm --------------------------------------------------------------------------- -compileIM :: Platform -> Config -> IMList a -> [Stm]-compileIM pform conf im = concatMap ((compileStm pform conf) . fst) im-+compileIM :: Config -> IMList a -> [C.Stm]+compileIM conf im = concatMap ((compileStm conf) . fst) im --------------------------------------------------------------------------- -- Generate entire Kernel ----------------------------------------------------------------------------type Parameters = [(String,Obsidian.Types.Type)]+type Parameters = [(String,T.Type)] -compile :: Platform -> Config -> String -> (Parameters,IMList a) -> Definition-compile pform config kname (params,im)- = go pform +compile :: Config -> String -> (Parameters,IMList a) -> C.Definition+compile config kname (params,im)+ = go where- stms = compileIM pform config im+ stms = compileIM config im - ps = compileParams pform params- go PlatformCUDA- = [cedecl| extern "C" __global__ void $id:kname($params:ps) {$items:cudabody} |]- go PlatformOpenCL- = [CL.cedecl| __kernel void $id:kname($params:ps) {$stms:stms} |]- go PlatformC- = [cedecl| extern "C" void $id:kname($params:ps) {$items:cbody} |] + ps = compileParams params+ go = [cedecl| extern "C" __global__ void $id:kname($params:ps) {$items:cudabody} |]+ -- go PlatformOpenCL+ -- = [CL.cedecl| __kernel void $id:kname($params:ps) {$stms:stms} |]+ -- go PlatformC+ -- = [cedecl| extern "C" void $id:kname($params:ps) {$items:cbody} |] cudabody = (if (configSharedMem config > 0) -- then [BlockDecl [cdecl| extern volatile __shared__ typename uint8_t sbase[]; |]] - then [BlockDecl [cdecl| __shared__ typename uint8_t sbase[$uint:(configSharedMem config)]; |]] + then [C.BlockDecl [cdecl| __shared__ typename uint8_t sbase[$uint:(configSharedMem config)] ; |]] else []) ++ --[BlockDecl [cdecl| typename uint32_t tid = threadIdx.x; |]] ++ --[BlockDecl [cdecl| typename uint32_t warpID = threadIdx.x / 32; |], -- BlockDecl [cdecl| typename uint32_t warpIx = threadIdx.x % 32; |]] ++ -- [BlockDecl [cdecl| typename uint32_t bid = blockIdx.x; |]] ++ (if (usesGid im) - then [BlockDecl [cdecl| typename uint32_t gid = blockIdx.x * blockDim.x + threadIdx.x; |]]+ then [C.BlockDecl [cdecl| typename uint32_t gid = blockIdx.x * blockDim.x + threadIdx.x; |]] else []) ++ (if (usesBid im) - then [BlockDecl [cdecl| typename uint32_t bid = blockIdx.x; |]]+ then [C.BlockDecl [cdecl| typename uint32_t bid = blockIdx.x; |]] else []) ++ (if (usesTid im) - then [BlockDecl [cdecl| typename uint32_t tid = threadIdx.x; |]]+ then [C.BlockDecl [cdecl| typename uint32_t tid = threadIdx.x; |]] else []) ++ (if (usesWarps im) - then [BlockDecl [cdecl| typename uint32_t warpID = threadIdx.x / 32; |],- BlockDecl [cdecl| typename uint32_t warpIx = threadIdx.x % 32; |]] + then [C.BlockDecl [cdecl| typename uint32_t warpID = threadIdx.x / 32; |],+ C.BlockDecl [cdecl| typename uint32_t warpIx = threadIdx.x % 32; |]] else []) ++ -- All variables used will be unique and can be declared -- at the top level @@ -532,34 +481,90 @@ -- Not sure if I am using language.C correctly. -- Maybe compileSTM should create BlockStms ? -- TODO: look how Nikola does it. - map BlockStm stms+ map C.BlockStm stms cbody = -- add memory allocation - map BlockStm stms+ map C.BlockStm stms -- Declare variables. -declares (SDeclare name t,_) = [BlockDecl [cdecl| $ty:(compileType t) $id:name;|]]+declares :: (Statement t,t) -> [C.BlockItem]+declares (SDeclare name t,_) = [C.BlockDecl [cdecl| $ty:(compileType t) $id:name;|]] declares (SCond _ im,_) = concatMap declares im declares (SSeqWhile _ im,_) = concatMap declares im declares (SForAll _ _ im,_) = concatMap declares im-declares (SDistrPar _ _ im,_) = concatMap declares im --- declares (SForAllBlocks _ im,_) = concatMap declares im--- declares (SNWarps _ im,_) = concatMap declares im--- declares (SWarpForAll _ im,_) = concatMap declares im +declares (SDistrPar _ _ im,_) = concatMap declares im+declares (SSeqFor _ _ im,_) = concatMap declares im declares _ = [] --------------------------------------------------------------------------- -- Parameter lists for functions (kernel head) ----------------------------------------------------------------------------compileParams :: Platform -> Parameters -> [Param]-compileParams PlatformOpenCL = map go- where- go (name,Pointer t) = [CL.cparam| global $ty:(compileType t) $id:name |]- go (name, t) = [CL.cparam| $ty:(compileType t) $id:name |]+compileParams :: Parameters -> [C.Param]+-- compileParams PlatformOpenCL = map go+-- where+-- go (name,Pointer t) = [CL.cparam| global $ty:(compileType t) $id:name |]+-- go (name, t) = [CL.cparam| $ty:(compileType t) $id:name |] -- C or CUDA -compileParams _ = map go+compileParams = map go where go (name,t) = [cparam| $ty:(compileType t) $id:name |] ++---------------------------------------------------------------------------+-- Compile with shared memory arrays declared at top+---------------------------------------------------------------------------+-- CODE DUPLICATION FOR NOW++compileDeclsTop :: Config -> [(String,((Word32,Word32),T.Type))] -> String -> (Parameters,IMList a) -> C.Definition+compileDeclsTop config toplevelarrs kname (params,im)+ = go + where+ stms = compileIM config im+ + ps = compileParams params+ go = [cedecl| extern "C" __global__ void $id:kname($params:ps) {$items:cudabody} |]+ -- go PlatformOpenCL+ -- = [CL.cedecl| __kernel void $id:kname($params:ps) {$stms:stms} |]+ -- go PlatformC+ -- = [cedecl| extern "C" void $id:kname($params:ps) {$items:cbody} |] ++ cudabody = (if (configSharedMem config > 0)+ -- then [BlockDecl [cdecl| extern volatile __shared__ typename uint8_t sbase[]; |]] + then [C.BlockDecl [cdecl| __shared__ typename uint8_t sbase[$uint:(configSharedMem config)]; |]] + else []) +++ --[BlockDecl [cdecl| typename uint32_t tid = threadIdx.x; |]] +++ --[BlockDecl [cdecl| typename uint32_t warpID = threadIdx.x / 32; |],+ -- BlockDecl [cdecl| typename uint32_t warpIx = threadIdx.x % 32; |]] +++-- [BlockDecl [cdecl| typename uint32_t bid = blockIdx.x; |]] +++ (if (usesGid im) + then [C.BlockDecl [cdecl| typename uint32_t gid = blockIdx.x * blockDim.x + threadIdx.x; |]]+ else []) ++ + (if (usesBid im) + then [C.BlockDecl [cdecl| typename uint32_t bid = blockIdx.x; |]]+ else []) ++ + (if (usesTid im) + then [C.BlockDecl [cdecl| typename uint32_t tid = threadIdx.x; |]]+ else []) +++ (if (usesWarps im) + then [C.BlockDecl [cdecl| typename uint32_t warpID = threadIdx.x / 32; |],+ C.BlockDecl [cdecl| typename uint32_t warpIx = threadIdx.x % 32; |]] + else []) +++ -- declare all arrays used+ concatMap declareArr toplevelarrs +++ -- All variables used will be unique and can be declared + -- at the top level + concatMap declares im ++ + -- Not sure if I am using language.C correctly. + -- Maybe compileSTM should create BlockStms ?+ -- TODO: look how Nikola does it. + map C.BlockStm stms++ cbody = -- add memory allocation + map C.BlockStm stms+++declareArr :: (String, ((Word32,Word32),T.Type)) -> [C.BlockItem]+declareArr (arr,((_,addr),t)) =+ [C.BlockDecl [cdecl| $ty:(compileType t) $id:arr = ($ty:(compileType t))(sbase + $int:addr);|]]
Obsidian/CodeGen/Liveness.hs view
@@ -1,6 +1,6 @@ -{- Joel Svensson 2012, 2013+{- Joel Svensson 2012..2015 notes: added case for SeqFor Jan-21-2013@@ -40,6 +40,9 @@ TODO: Can ixs contain array names ? (Most likely yes! think about the counting sort example) + BUG: Arrays from the "outside" needs special+ treatment within a loop body. + -} @@ -114,11 +117,20 @@ -- Is this correct ? Same question, all below return (SCond bexp iml,ns) ++-- *** THIS IS FIXED *** +-- This needs to change. +-- arrays from the "outside" that +-- are used within the loop needs special treatment. process (SSeqFor nom n im,_) = do + -- get names alive after loop s <- get let iml = computeLiveness1 s im + -- l is liveness info "leaving" im l = safeHead iml+ -- alive at these points are those things in l + -- plus the things before (s) ns = s `Set.union` l put ns return (SSeqFor nom n iml,ns) @@ -134,7 +146,6 @@ do s <- get return (SBreak,s)- process (SForAll lvl n im,_) = do
− Obsidian/CodeGen/Memory.hs
@@ -1,249 +0,0 @@--{- Joel Svensson 2012, 2013- - notes:- Added a SeqFor case Jan-21-2013-- -} -module Obsidian.CodeGen.Memory - (MemMap,- Memory,- allocate,- free,- freeAll,- size, - sharedMem, - Address,- Bytes,- mmIM,- renameIM ) - where --import qualified Data.List as List-import qualified Data.Set as Set-import Data.Word--import Obsidian.Types-import Obsidian.Globs--import Obsidian.Exp -import Obsidian.CodeGen.Program-import Obsidian.CodeGen.Liveness--import qualified Data.Map as Map -------------------------------------------------------------------------------- Memory layout------------------------------------------------------------------------------type MemMap = Map.Map Name (Word32,Type)--type Address = Word32-type Bytes = Word32 --data Memory = Memory {freeList :: [(Address,Bytes)] ,- allocated :: [(Address,Bytes)] , - size :: Bytes} -- how much used- deriving Show - - --- 48 kilobytes of smem -sharedMem = Memory [(0,49152)] [] 0---updateMax :: Memory -> Memory -updateMax mem = let m = maximum [a+b|(a,b) <- allocated mem]- m' = max m (size mem)- in mem {size = m'}---- This one needs to check that shared memory is not full.-allocate :: Memory -> Bytes -> (Memory,Address)-allocate m b = - let adress = filter (\(x,y) -> y >= b) (freeList m) -- get a list of candidates- getTop mem = let (a,b) = case null (allocated m) of - False -> maximum $ List.sort (allocated m) - True -> (0,0)- in a+b- in case adress of - -- use the first candidate (try better approaches - -- such as searching for best match, so that not to waste memory)- ((a,bytes):_) -> let fl = filter (\(addr,_) -> a /= addr) (freeList m)- fl' = if b < bytes - then (a+b,bytes-b):fl- else fl- in (updateMax (m {freeList = fl', - allocated = (a,b):allocated m}) ,a)- [] -> error "out of shared memory"- - -free :: Memory -> Address -> Memory-free m a = mem - where - bytes = lookup a (allocated m)- al = filter (\(addr,_) -> a /= addr) (allocated m)-- -- TODO: Investigate this much closer.- -- Is it a bug or is freeing a non allocated memory area- -- OK?- - mem = case bytes of - Nothing -> m- {-- error $ "error: Address " ++ show a ++- " not found in allocated list" ++- "\n" ++ show m- -} - Just b -> m {freeList = compress ((a,b):(freeList m)),- allocated = al}--freeAll :: Memory -> [Address] -> Memory -freeAll m [] = m-freeAll m (a:as) = freeAll (free m a) as--compress = merge . List.sort --merge [] = [] -merge [x] = [x]-merge ((x,b):(y,b2):xs) = if (x+b == y)- then merge ((x,b+b2):xs) - else (x,b):merge((y,b2):xs)-------------------------------------------------------------------------------- Memory map the new IM-----------------------------------------------------------------------------mmIM :: IML -> Memory -> MemMap -> (Memory, MemMap)-mmIM im memory memmap = r im (memory,memmap)- where- r [] m = m- r (x:xs) (m,mm) =- let- (m',mm') = process x m mm- - freeable = getFreeableSet x xs- freeableAddrs = mapM (flip Map.lookup mm') (filter dontMap (Set.toList freeable))- dontMap name = not ((List.isPrefixOf "input" name) || - (List.isPrefixOf "output" name))- mNew =- case freeableAddrs of- (Just as) -> freeAll m' (map fst as)- Nothing -> m'- in r xs (mNew,mm')-mmIM' :: IML -> Memory -> MemMap -> (Memory, MemMap)-mmIM' im memory memmap = r im (memory,memmap)- where- r [] m = m- r (x:xs) (m,mm) =- let- (m',mm') = process x m mm- - freeable = getFreeableSet x xs- freeableAddrs = mapM (flip Map.lookup mm') (filter dontMap (Set.toList freeable))- dontMap name = not ((List.isPrefixOf "input" name) || - (List.isPrefixOf "output" name))- mNew =- case freeableAddrs of- (Just as) -> m' -- freeAll m' (map fst as)- Nothing -> m'- in r xs (mNew,mm')-- -process (SAllocate name size t,_) m mm = (m',mm') - where (m',addr) = allocate m size- mm' = case Map.lookup name mm of- Nothing -> Map.insert name (addr,t) mm- (Just (a, t)) -> error $ "mmIm: " ++ name ++ " is already mapped to " ++ show a-- -- A tricky case. --- process (SForAllBlocks n im,_) m mm = mmIM im m mm- -- Another tricky case. -process (SSeqFor _ n im,_) m mm = mmIM im m mm-process (SSeqWhile b im,_) m mm = mmIM im m mm - -- Yet another tricky case.-process (SForAll _ n im,_) m mm = mmIM im m mm-process (SDistrPar Warp n im,_) m mm = mmIM' im m mm -- mmIM im m mm-process (SDistrPar Block n im,_) m mm = mmIM im m mm - -- The worst of them all.--- process (SForAllThreads n im,_) m mm = mmIM im m mm--- process (SNWarps _ im,_) m mm = mmIM im m mm--- process (SWarpForAll _ im,_) m mm = mmIM im m mm -- -- process im m mm = error $ printStm im -- "process: WHat!"-process (_,_) m mm = (m,mm) ---- Friday (2013 Mars 29, discovered bug) -getFreeableSet :: (Statement Liveness,Liveness) -> IML -> Liveness -getFreeableSet (_,l) [] = Set.empty -- not l ! -getFreeableSet (_,l) ((_,l1):_) = l Set.\\ l1-------------------------------------------------------------------------------- Rename arrays in IM---------------------------------------------------------------------------- --renameIM :: MemMap -> IML -> IMList ()-renameIM mm im = zip (map (go . fst) im) (repeat ())- where- go (SAssign name ix e) = SAssign (renameIVar mm name)- (map (renameIExp mm) ix)- (renameIExp mm e)- go (SAtomicOp name ix atop) = SAtomicOp (renameIVar mm name)- (renameIExp mm ix)- (renameAtOp mm atop) - go (SCond be im) = SCond (renameIExp mm be)- (renameIM mm im)- go (SSeqFor str n im) = SSeqFor str (renameIExp mm n)- (renameIM mm im)- go SBreak = SBreak- go (SSeqWhile n im) = SSeqWhile (renameIExp mm n)- (renameIM mm im)- go (SForAll lvl n im) = SForAll lvl (renameIExp mm n)- (renameIM mm im)- go (SDistrPar lvl n im) = SDistrPar lvl (renameIExp mm n)- (renameIM mm im) ---- go (SForAllBlocks n im) = SForAllBlocks (renameIExp mm n)--- (renameIM mm im)--- go (SNWarps n im) = SNWarps (renameIExp mm n)--- (renameIM mm im)--- go (SWarpForAll n im) = SWarpForAll (renameIExp mm n)--- (renameIM mm im) - -- Strip this out earlier. - go (SAllocate name n t) = SAllocate name n t - go (SDeclare name t) = SDeclare name t- go SSynchronize = SSynchronize -------------------------------------------------------------------------------- Memory map the arrays in an CExpr-----------------------------------------------------------------------------renameIExp mm e@(IVar nom t) = renameIVar mm e -renameIExp mm (IIndex (e1,es) t) = IIndex (renameIExp mm e1, map (renameIExp mm) es) t-renameIExp mm (IBinOp op e1 e2 t) = IBinOp op (renameIExp mm e1) (renameIExp mm e2) t-renameIExp mm (IUnOp op e t) = IUnOp op (renameIExp mm e) t -renameIExp mm (IFunCall nom exprs t) = IFunCall nom (map (renameIExp mm) exprs) t-renameIExp mm (ICast e t) = ICast (renameIExp mm e) t-renameIExp mm (ICond e1 e2 e3 t) = ICond (renameIExp mm e1)- (renameIExp mm e2)- (renameIExp mm e3)- t-renameIExp mm a = a---renameIVar mm (IVar name t) =- case Map.lookup name mm of - Just (addr,t) -> - let core = sbaseIExp addr - cast c = ICast c t- in cast core- - Nothing -> IVar name t- where- sbaseIExp 0 = IVar "sbase" (Pointer Word8) - sbaseIExp addr = IBinOp IAdd (IVar "sbase" (Pointer Word8)) - (IWord32 addr) - (Pointer Word8) --renameAtOp mm AtInc = AtInc-renameAtOp mm (AtAdd e) = AtAdd (renameIExp mm e)-renameAtOp mm (AtSub e) = AtSub (renameIExp mm e)-renameAtOp mm (AtExch e) = AtExch (renameIExp mm e) --
+ Obsidian/CodeGen/Memory2.hs view
@@ -0,0 +1,457 @@++{- Joel Svensson 2012..2015++ Notes:+ Jan-27-2015: Bug fix related to arrays alive when+ entering into loops.+ Fix seems to solve the problem.+ Need to investigate that it still frees arrays+ as soon as possible. + Nov-25-2014: Changes to memory management+ + Jan-21-2013: Added a SeqFor case ++ -} +module Obsidian.CodeGen.Memory2 + (MemMap,+ Memory,+ allocate,+ free,+ freeAll,+ size, + Address,+ Bytes,+ memMapIM,+ renameIM,+ SharedMemConfig(..),+ createSharedMem+ ) + where ++import qualified Data.List as List+import qualified Data.Set as Set+import Data.Word+import Data.Maybe ++import Obsidian.Types+import Obsidian.Globs++import Obsidian.Exp +import Obsidian.CodeGen.Program+import Obsidian.CodeGen.Liveness++import Debug.Trace +++import qualified Data.Map as Map +---------------------------------------------------------------------------+-- Planned improvements+---------------------------------------------------------------------------+-- # Always start a shared memory array at a "bank-aligned" address+-- + So that programmer can really effect access patterns. +-- # Do not rename with pointers+-- instead output a list of (type ,name,address) quadruples+-- that can be used to create an alias (at top scope of program)+-- + Prettier code -> easier debugging +-- + Potential efficiency issues, from less casting etc+-- # (LONG-TERM) Clever memory allocation+-- + The future is known! So encode the optimal memory+-- allocation schema++---------------------------------------------------------------------------+-- Memory layout+---------------------------------------------------------------------------++type MemMap = Map.Map Name (AlignedAddress,Type)++type AlignedAddress = (Address,Address) ++type Address = Word32+type Bytes = Word32 ++data Memory = Memory {freeList :: [(Address,Bytes)] ,+ allocated :: [(Address,Bytes)] , + size :: Bytes} -- how much used+ deriving Show ++updateMax :: Memory -> Memory +updateMax mem = let m = maximum [a+b|(a,b) <- allocated mem]+ m' = max m (size mem)+ in mem {size = m'}+ ++---------------------------------------------------------------------------+-- Shared memory configurations+---------------------------------------------------------------------------+data SharedMemConfig =+ SharedMemConfig { smSize :: Bytes -- amount of shared mem+ , smBanks :: Word32 -- Number of banks 16/32+ , smBankAlign :: Bool + }+++createSharedMem :: SharedMemConfig -> Memory+createSharedMem conf = Memory [(0,smSize conf)] [] 0 ++-- bank allign an address, returning a new aligned address+-- and the number of EXTRA bytes that needs to be present+-- from the old provided address in order to store aligned data+-- at the new location.+bank_align :: SharedMemConfig -> Address -> (AlignedAddress, Bytes)+bank_align conf address =+ case (how_far_off == 0) of+ True -> ((address,address),0)+ False -> ((address,address + bump), bump) + + where banks = smBanks conf+ -- if address % bank_alignment == 0+ -- the address is aligned+ bank_alignment = banks * 4 -- number of banks * 4 bytes+ + how_far_off = address `mod` bank_alignment+ bump = bank_alignment - how_far_off ++---------------------------------------------------------------------------+-- Allocate memory+--------------------------------------------------------------------------- +allocate :: SharedMemConfig -> Memory -> Bytes -> (Memory,AlignedAddress)+allocate conf m b =+ case smBankAlign conf of+ True ->+ -- Does any memory location exist that+ -- allows for the allocation of this array + case catMaybes new_candidates of+ [] -> error $ "allocate: out of shared memory:" +++ "\n Allocating: " ++ show b ++ " bytes" +++ "\n Free List: " ++ show (freeList m) ++ + "\n Potentials: " ++ show address_candidates ++ + "\n Fit with align: " ++ show new_candidates + + ((aligned_address,free_space,alloc_size):_) ->+ -- update free list+ -- Clear the allocated address from the free list + let fl = filter (\(addr,_) -> (fst aligned_address /= addr)) (freeList m)+ -- if the chosen space is larger than what we need+ -- add the unused chunk to the free list + fl' = if alloc_size < free_space+ then (fst aligned_address + alloc_size,+ free_space - alloc_size):fl+ else fl+ -- Update memory and return a result address+ in (updateMax $ m { freeList = fl'+ , allocated =+ (fst aligned_address,alloc_size):allocated m}+ , aligned_address)+ + False ->+ case map (pretend_align b) address_candidates of+ [] -> error "out of shared memory"+ + ((aligned_address,free_space,alloc_size):_) ->+ -- update free list+ -- Clear the allocated address from the free list + let fl = filter (\(addr,_) -> (fst aligned_address /= addr)) (freeList m)+ -- if the chosen space is larger than what we need+ -- add the unused chunk to the free list + fl' = if alloc_size < free_space+ then (fst aligned_address + alloc_size,+ free_space - alloc_size):fl+ else fl+ -- Update memory and return a result address+ in (updateMax $ m { freeList = fl'+ , allocated =+ (fst aligned_address,alloc_size):allocated m}+ , aligned_address)++ + ++ where+ -- Candidates after aligning+ new_candidates = map (tryCandidate b) address_candidates+ -- Original address canditades + address_candidates = filter (\(_,y) -> y >= b) $ freeList m + -- Create silly AlignedAddress (that are not really aligned at all)+ pretend_align bytes (addr, free_space) = ((addr,addr),free_space,bytes) ++ -- try to align an address+ -- results in an AlignedAdress+ tryCandidate bytes (addr, free_space) =+ let (aligned_addr, extra_bytes) = bank_align conf addr+ alloc_size = bytes + extra_bytes+ in+ case free_space >= alloc_size of + True -> Just (aligned_addr,free_space,alloc_size)+ False -> Nothing ++---------------------------------------------------------------------------+-- Free memory +---------------------------------------------------------------------------+free :: Memory -> AlignedAddress -> Memory+free m (alloc_addr,_) = mem + where + bytes = lookup (alloc_addr) (allocated m)+ al = filter (\(addr,_) -> alloc_addr /= addr) (allocated m)++ -- TODO: Investigate this much closer.+ -- Is it a bug or is freeing a non allocated memory area+ -- OK?+ -- 2014-Nov-25: I dont remember what this refers to+ -- But, if a problem resurfaces, look here. + + mem = case bytes of + Nothing -> m+ {-+ error $ "error: Address " ++ show a +++ " not found in allocated list" +++ "\n" ++ show m+ -} + Just b -> m {freeList = compress ((alloc_addr,b):(freeList m)),+ allocated = al}++freeAll :: Memory -> [AlignedAddress] -> Memory +freeAll m [] = m+freeAll m (a:as) = freeAll (free m a) as+++compress :: [(Address,Bytes)] -> [(Address,Bytes)]+compress = merge . List.sort+ where+ merge :: [(Address,Bytes)] -> [(Address,Bytes)]+ merge [] = [] + merge [x] = [x]+ merge ((x,b):(y,b2):xs) = if (x+b == y)+ then merge ((x,b+b2):xs) + else (x,b):merge((y,b2):xs)+++---------------------------------------------------------------------------+-- Memory map the new IM+---------------------------------------------------------------------------++memMapIM :: SharedMemConfig -> IML -> MemMap -> (Memory, MemMap)+memMapIM conf im memmap = mmIM conf im memory memmap+ where+ memory = createSharedMem conf+ +mmIM :: SharedMemConfig -> IML -> Memory -> MemMap -> (Memory, MemMap)+mmIM conf im memory memmap = r im (memory,memmap)+ where + r [] m = m+ r (x:xs) (m,mm) =+ let+ (m',mm') = process conf x m mm+ + freeable = getFreeableSet x xs+ freeableAddrs = mapM (flip Map.lookup mm') (filter dontMap (Set.toList freeable))+ dontMap name = not ((List.isPrefixOf "input" name) || + (List.isPrefixOf "output" name))+ mNew =+ case freeableAddrs of+ (Just as) -> freeAll m' (map fst as)+ Nothing -> m'+ in -- trace ("freeable: " ++ show freeable ++ "\n") $ + r xs (mNew,mm')+ + process :: SharedMemConfig -> (Statement Liveness,Liveness) -> Memory -> MemMap -> (Memory,MemMap)+ process conf (SAllocate name size t,_) m mm = (m',mm') + where (m',addr) = allocate conf m size+ mm' =+ case Map.lookup name mm of+ Nothing -> Map.insert name (addr,t) mm+ (Just (a, t)) -> error $ "mmIm: " ++ name ++ " is already mapped to " ++ show a++ -- Boilerplate+ -- BUG: Bug in memory management related to seqloops+ -- It may be better to try to fix this bug here.+ -- A special mmIM for the loop case may be needed. + process conf (SSeqFor _ n im,alive) m mm = mmIMLoop conf alive im m mm+ process conf (SSeqWhile b im,_) m mm = mmIM conf im m mm + process conf (SForAll _ n im,_) m mm = mmIM conf im m mm+ -- 2014-Nov-25:+ -- This one used mmIM' which was identical to mmIM.+ -- This must have been a leftover from when I thought+ -- warp memory needed some special attention here. + process conf (SDistrPar Warp n im,_) m mm = mmIMDistrWarp conf im m mm + process conf (SDistrPar Block n im,_) m mm = mmIM conf im m mm + process conf (_,_) m mm = (m,mm) ++-- Friday (2013 Mars 29, discovered bug)+-- 2014-Nov-25: was the "l" the bug ? (some details help)+getFreeableSet :: (Statement Liveness,Liveness) -> IML -> Liveness +getFreeableSet (_,l) [] = Set.empty -- not l ! +getFreeableSet (_,l) ((_,l1):_) = l Set.\\ l1+++---------------------------------------------------------------------------+-- +---------------------------------------------------------------------------+mmIMLoop conf nonfreeable im memory memmap = r im (memory,memmap)+ where + r [] m = m+ r (x:xs) (m,mm) =+ let+ (m',mm') = process conf x m mm+ + freeable' = getFreeableSet x xs+ freeable = freeable' Set.\\ nonfreeable+ freeableAddrs = mapM (flip Map.lookup mm') (filter dontMap (Set.toList freeable))+ dontMap name = not ((List.isPrefixOf "input" name) || + (List.isPrefixOf "output" name))+ mNew =+ case freeableAddrs of+ (Just as) -> freeAll m' (map fst as)+ Nothing -> m'+ in --trace ("freeable': " ++ show freeable' ++ "\n" +++ -- "freeable: " ++ show freeable ++ "\n" ++ + -- "nonfreeable: " ++ show nonfreeable) $+ r xs (mNew,mm')+ + process :: SharedMemConfig -> (Statement Liveness,Liveness) -> Memory -> MemMap -> (Memory,MemMap)+ process conf (SAllocate name size t,_) m mm = (m',mm') + where (m',addr) = allocate conf m size+ mm' =+ case Map.lookup name mm of+ Nothing -> Map.insert name (addr,t) mm+ (Just (a, t)) -> error $ "mmIm: " ++ name ++ " is already mapped to " ++ show a++ -- Boilerplate+ process conf (SSeqFor _ n im,alive) m mm = mmIMLoop conf (nonfreeable `Set.union` alive) im m mm+ process conf (SSeqWhile b im,_) m mm = mmIMLoop conf nonfreeable im m mm + process conf (SForAll _ n im,_) m mm = mmIMLoop conf nonfreeable im m mm+ -- 2014-Nov-25:+ -- This one used mmIM' which was identical to mmIM.+ -- This must have been a leftover from when I thought+ -- warp memory needed some special attention here. + process conf (SDistrPar Warp n im,_) m mm = mmIMLoop conf nonfreeable im m mm + process conf (SDistrPar Block n im,_) m mm = mmIMLoop conf nonfreeable im m mm + process conf (_,_) m mm = (m,mm) +++-- NOTE: This is a hack to make programs distributed+-- over warps not "free" its arrays.+-- Distributing over warps introduces entirely new+-- shared memory behaviour..+-- This needs a review and some real thought!++mmIMDistrWarp conf im memory memmap = r im (memory,memmap)+ where + r [] m = m+ r (x:xs) (m,mm) =+ let+ (m',mm') = process conf x m mm+ + freeable = getFreeableSet x xs+ --freeable = freeable' Set.\\ nonfreeable+ freeableAddrs = mapM (flip Map.lookup mm') (filter dontMap (Set.toList freeable))+ dontMap name = not ((List.isPrefixOf "input" name) || + (List.isPrefixOf "output" name))+ mNew =+ case freeableAddrs of+ (Just as) -> m' -- freeAll m' (map fst as)+ Nothing -> m'+ in --trace ("freeable': " ++ show freeable' ++ "\n" +++ -- "freeable: " ++ show freeable ++ "\n" ++ + -- "nonfreeable: " ++ show nonfreeable) $+ r xs (mNew,mm')+ + process :: SharedMemConfig -> (Statement Liveness,Liveness) -> Memory -> MemMap -> (Memory,MemMap)+ process conf (SAllocate name size t,_) m mm = (m',mm') + where (m',addr) = allocate conf m size+ mm' =+ case Map.lookup name mm of+ Nothing -> Map.insert name (addr,t) mm+ (Just (a, t)) -> error $ "mmIm: " ++ name ++ " is already mapped to " ++ show a++ -- Boilerplate+ process conf (SSeqFor _ n im,alive) m mm = mmIMDistrWarp conf im m mm+ process conf (SSeqWhile b im,_) m mm = mmIMDistrWarp conf im m mm + process conf (SForAll _ n im,_) m mm = mmIMDistrWarp conf im m mm+ -- 2014-Nov-25:+ -- This one used mmIM' which was identical to mmIM.+ -- This must have been a leftover from when I thought+ -- warp memory needed some special attention here. + process conf (SDistrPar Warp n im,_) m mm = mmIMDistrWarp conf im m mm + process conf (SDistrPar Block n im,_) m mm = mmIMDistrWarp conf im m mm + process conf (_,_) m mm = (m,mm) ++++++---------------------------------------------------------------------------+-- Rename arrays in IM+--------------------------------------------------------------------------- ++renameIM :: MemMap -> IML -> IMList ()+renameIM mm im = zip (map (go . fst) im) (repeat ())+ where+ go (SAssign name ix e) = SAssign (renameIVar mm name)+ (map (renameIExp mm) ix)+ (renameIExp mm e)+ go (SAtomicOp name ix atop) = SAtomicOp (renameIVar mm name)+ (renameIExp mm ix)+ (renameAtOp mm atop) + go (SCond be im) = SCond (renameIExp mm be)+ (renameIM mm im)+ go (SSeqFor str n im) = SSeqFor str (renameIExp mm n)+ (renameIM mm im)+ go SBreak = SBreak+ go (SSeqWhile n im) = SSeqWhile (renameIExp mm n)+ (renameIM mm im)+ go (SForAll lvl n im) = SForAll lvl (renameIExp mm n)+ (renameIM mm im)+ go (SDistrPar lvl n im) = SDistrPar lvl (renameIExp mm n)+ (renameIM mm im) ++-- go (SForAllBlocks n im) = SForAllBlocks (renameIExp mm n)+-- (renameIM mm im)+-- go (SNWarps n im) = SNWarps (renameIExp mm n)+-- (renameIM mm im)+-- go (SWarpForAll n im) = SWarpForAll (renameIExp mm n)+-- (renameIM mm im) + -- Strip this out earlier. + go (SAllocate name n t) = SAllocate name n t + go (SDeclare name t) = SDeclare name t+ go SSynchronize = SSynchronize ++---------------------------------------------------------------------------+-- Memory map the arrays in an CExpr+---------------------------------------------------------------------------+renameIExp :: MemMap -> IExp -> IExp +renameIExp mm e@(IVar _ _) = renameIVar mm e +renameIExp mm (IIndex (e1,es) t) = IIndex (renameIExp mm e1, map (renameIExp mm) es) t+renameIExp mm (IBinOp op e1 e2 t) = IBinOp op (renameIExp mm e1) (renameIExp mm e2) t+renameIExp mm (IUnOp op e t) = IUnOp op (renameIExp mm e) t +renameIExp mm (IFunCall nom exprs t) = IFunCall nom (map (renameIExp mm) exprs) t+renameIExp mm (ICast e t) = ICast (renameIExp mm e) t+renameIExp mm (ICond e1 e2 e3 t) = ICond (renameIExp mm e1)+ (renameIExp mm e2)+ (renameIExp mm e3)+ t+renameIExp _ a = a++renameIVar :: MemMap -> IExp -> IExp +renameIVar mm (IVar name t) =+ -- t == t1 should be true + case Map.lookup name mm of + Just ((_,real_addr),t1) -> + let core = sbaseIExp (real_addr)+ cast c = ICast c t1+ in cast core+ + Nothing -> IVar name t+ where+ sbaseIExp 0 = IVar "sbase" (Pointer Word8) + sbaseIExp addr = IBinOp IAdd (IVar "sbase" (Pointer Word8)) + (IWord32 addr) + (Pointer Word8)+renameIVar _ _ = error "renameIVar: incorrect expression" ++renameAtOp :: MemMap -> AtOp -> AtOp +renameAtOp _ AtInc = AtInc+renameAtOp mm (AtAdd e) = AtAdd (renameIExp mm e)+renameAtOp mm (AtSub e) = AtSub (renameIExp mm e)+renameAtOp mm (AtExch e) = AtExch (renameIExp mm e) ++
Obsidian/CodeGen/Program.hs view
@@ -3,10 +3,12 @@ FlexibleInstances #-} {-# LANGUAGE TypeSynonymInstances #-} {-# LANGUAGE GeneralizedNewtypeDeriving #-}+{-# LANGUAGE ConstraintKinds #-}+{-# LANGUAGE DataKinds #-} {- CodeGen.Program. - Joel Svensson 2012, 2013+ Joel Svensson 2012..2015 Notes: 2013-03-17: Codegeneration is changing@@ -21,7 +23,6 @@ import Obsidian.Atomic import qualified Obsidian.Program as P-import Obsidian.Program (Step,Zero) import Data.Word import Data.Supply@@ -30,6 +31,7 @@ import System.IO.Unsafe import Control.Monad.State+import Control.Applicative ---------------------------------------------------------------------------@@ -40,7 +42,7 @@ type IM = IMList () --- out :: +out :: a -> [(a,())] out a = [(a,())] -- Atomic operations@@ -68,11 +70,6 @@ | SForAll HLevel IExp (IMList t) | SDistrPar HLevel IExp (IMList t) --- | SForAllBlocks IExp (IMList t)--- | SNWarps IExp (IMList t)--- | SWarpForAll IExp (IMList t) --- | SWarpForAll String String IExp (IMList t) - -- Memory Allocation.. | SAllocate Name Word32 Type | SDeclare Name Type@@ -84,11 +81,12 @@ --------------------------------------------------------------------------- -- Collect and pass around data during first step compilation data Context = Context { ctxNWarps :: Maybe Word32,+ ctxNThreads :: Maybe Word32, ctxGLBUsesTid :: Bool, ctxGLBUsesWid :: Bool} newtype CM a = CM (State Context a)- deriving (Monad, MonadState Context)+ deriving (Monad, MonadState Context, Functor, Applicative) runCM :: CM a -> Context -> a --runCM (CM cm) ctx = evalState cm ctx@@ -106,15 +104,24 @@ enterWarp :: Word32 -> CM () enterWarp n = modify $ \ctx -> ctx { ctxNWarps = Just n } +enterThread :: Word32 -> CM ()+enterThread n = modify $ \ctx -> ctx {ctxNThreads = Just n} + clearWarp :: CM () clearWarp = modify $ \ctx -> ctx {ctxNWarps = Nothing} getNWarps :: CM (Maybe Word32) getNWarps = do ctx <- get- return $ ctxNWarps ctx + return $ ctxNWarps ctx -emptyCtx = Context Nothing False False+getNThreads :: CM (Maybe Word32)+getNThreads = do+ ctx <- get+ return $ ctxNThreads ctx ++emptyCtx :: Context+emptyCtx = Context Nothing Nothing False False --------------------------------------------------------------------------- @@ -131,11 +138,12 @@ where go (SDistrPar _ _ im) = usesTid im go (SForAll Block _ _) = True+ go (SSeqFor _ _ im) = usesTid im go _ = False usesBid :: IMList t -> Bool usesBid = any (go . fst) where- go (SDistrPar Block _ im) = True -- usesBid im+ go (SDistrPar Block _ _) = True -- usesBid im -- go (SForAll Block _ _) = True go _ = False usesGid :: IMList t -> Bool@@ -173,14 +181,30 @@ (a,im) <- compile i2 p return ((),out $ SSeqFor nom (expToIExp n) im)- + compile s (P.Allocate nom n t) = do+ (Just nt) <- getNThreads -- must be a Just at this point+ nw' <- getNWarps++ let nw = case nw' of+ Nothing -> 1+ Just i -> i+ + return ((),out $ SAllocate nom (nt*nw*n) t)+ compile _ (P.Sync) =+ return ((),[])+ compile s p = cs s p -- Compile Warp program instance Compile P.Warp where compile s (P.DistrPar n f) = error "Currently not supported to distribute over the threads, use ForAll instead!"- compile s (P.ForAll n f) = do+ compile s (P.ForAll n@(Literal n') f) = do++ -- setup context to know number of threads+ -- executing + enterThread n'+ let p = f (variable "warpIx") (a,im) <- compile s p return (a, out $ SForAll Warp (expToIExp n) im)@@ -195,47 +219,57 @@ return (b,(im1 ++ im2)) compile s (P.Return a) = return (a,[]) compile s (P.Identifier) = return (supplyValue s, [])-+ compile s (P.Sync) = return ((),[])+ -- Why no fallthrough here ?+ -- Adding (must have been a horrible mistake!)+ compile s p = cs s p+ -- Compile Block program instance Compile P.Block where- compile s (P.ForAll n f) = do+ compile s (P.ForAll n@(Literal n') f) = do++ -- Set up the context to know the number of+ -- concurrent thread programs that are executing. + enterThread n'+ setUsesTid+ let nom = "tid" v = variable nom p = f v- setUsesTid + (a,im) <- compile s p -- in this case a could be () (since it is guaranteed to be anyway). a return (a,out (SForAll Block (expToIExp n) im)) compile s (P.DistrPar n'@(Literal n) f) = do - {- Distribute work over warps! -} + {- Distribute work over warps! -}+ -- Set up the context for the compilation+ -- of the Warp code.+ -- BUG: Something like this is needed for distribution+ -- over threads too!+ -- FIXED: Bug mentioned above should be (at least) partially fixed. enterWarp n- -- Need to generate some IM here that the backend can read the- -- Number of desired warps from. + -- Number of active warps are stored in the context. (a,im) <- compile s (f (variable "warpID")) return (a, out (SDistrPar Warp (expToIExp n') im))+ compile s (P.Allocate id n t) = return ((),out (SAllocate id n t))+ compile s (P.Sync) = return ((),out (SSynchronize)) compile s p = cs s p -- Compile a Grid Program instance Compile P.Grid where- compile s (P.ForAll n f) = do- - -- Incorrect, need to compute global thread ids and apply - let p = f gid -- (BlockIdx X)- gid = variable "gid" - (a,im) <- compile s p - return (a, out (SForAll Grid (expToIExp n) im))- {- Distribute over blocks -} compile s (P.DistrPar n f) = do -- Need to generate IM here that the backend can read desired number of blocks from- let p = f (BlockIdx X) + let p = f (variable "bid") -- "blockIdx.x") -- (BlockIdx X) (a, im) <- compile s p -- (f (BlockIdx X)) return (a, out (SDistrPar Block (expToIExp n) im))+ compile s (P.Allocate _ _ _) = error "Allocate at level Grid" compile s p = cs s p + {- ForAll cannot happen here! -} @@ -250,27 +284,14 @@ cs i (P.AtomicOp name ix atom) = case atom of AtomicInc -> return $ ((),out (SAtomicOp (IVar name Word32) (expToIExp ix) AtInc))- AtomicAdd e -> undefined- AtomicSub e -> undefined- AtomicExch e -> undefined - -- (vres, out im)- -- where- -- res = "a" ++ show (supplyValue i)- -- vres = IVar res (typeOf (undefined :: a))- -- vname = IVar name (typeOf (undefined :: a))- -- im = SAtomicOp vres vname ix atom----cs i (P.AtomicOp name ix at) = (v,out im)--- where --- nom = "a" ++ show (supplyValue i)--- v = variable nom--- im = SAtomicOp nom name ix at+ AtomicAdd e -> error $ "CodeGen.Program: AtomicAdd is not implemented"+ AtomicSub e -> error $ "CodeGen.Program: AtomicSub is not implemented" + AtomicExch e -> error $ "CodeGen.Program: AtomicExch is not implemented" cs i (P.Cond bexp p) = do ((),im) <- compile i p return ((),out (SCond (expToIExp bexp) im)) - cs i (P.SeqFor n f) = do let (i1,i2) = split2 i nom = "i" ++ show (supplyValue i1)@@ -289,11 +310,9 @@ cs i (P.Break) = return ((), out SBreak) -cs i (P.Allocate id n t) = return ((),out (SAllocate id n t))-cs i (P.Declare id t) = return ((),out (SDeclare id t))--cs i (P.Sync) = return ((),out (SSynchronize))+cs i (P.Allocate id n t) = error $ "CodeGen.Program: Allocate without a hierarchy designation" +cs i (P.Declare id t) = return ((),out (SDeclare id t)) cs i (P.Bind p f) = do let (s1,s2) = split2 i@@ -307,10 +326,10 @@ -- Unhandled cases -cs i p = error $ "#Program.hs# unhandled in cs: " ++ P.printPrg p -- compile i p +cs i p = error $ "CodeGen.Program: unhandled in cs: " ++ P.printPrg p -- compile i p ------------------------------------------------------------------------------ Turning IM to strings+-- Turning IM to strings (outdated and broken) --------------------------------------------------------------------------- printIM :: Show a => IMList a -> String @@ -357,20 +376,6 @@ printStm (SDistrPar lvl n im,m) = "forAll gid in [0.." ++ show n ++"] do" ++ meta m ++ concatMap printStm im ++ "\ndone;\n"---- printStm (SWarpForAll n im ,m) =--- "forAll(InWarp) tid" ++ " in [0.." ++ show n ++"] do" ++ meta m ++--- concatMap printStm im ++ "\ndone;\n"---- printStm (SNWarps n im,m) = "Run " ++ show n++ " Warps {\n" ++--- printIM im ++ "\n }"---printStm (SForAllThreads n im,m) =--- "forAllThreads i in [0.." ++ show n ++"] do" ++ meta m ++ --- concatMap printStm im ++ "\ndone;\n"--- --- printStm (a,m) = error $ show m meta :: Show a => a -> String meta m = "\t//" ++ show m ++ "\n"
Obsidian/CodeGen/Reify.hs view
@@ -1,3 +1,4 @@+ {-# LANGUAGE FlexibleInstances, OverlappingInstances, UndecidableInstances,@@ -22,27 +23,28 @@ import Obsidian.Mutable import Obsidian.Types-import Obsidian.Globs import Obsidian.Program-import Obsidian.Force-import Obsidian.Memory--import Obsidian.Names import Obsidian.Library+import Obsidian.Globs+import Obsidian.Names import qualified Obsidian.CodeGen.Program as CG import Obsidian.CodeGen.CompileIM import Data.Word-import Data.Int-import qualified Data.Map as M ------------------------------------------------------------------------------ New approach (hopefully)---------------------------------------------------------------------------- -- "reify" Haskell functions into CG.Programs+--------------------------------------------------------------------------- +{- TODO:+ needs to be revamped to support more than one output array+ from a single program.+ * needs a class for "outputs" generation.+-} ++ --------------------------------------------------------------------------- -- --------------------------------------------------------------------------- @@ -163,6 +165,12 @@ lengthVar = variable n input = namedMutable nom lengthVar t = typeOf_ (undefined :: t)++namedMutable :: Name -> s -> Mutable mloc s a+namedMutable s v = Mutable v (Single s)+undefinedMutable :: s -> Mutable mloc s a +undefinedMutable v = Mutable v undefined + instance (ToProgram b, Scalar t) => ToProgram ((Exp t) -> b) where
+ Obsidian/Data.hs view
@@ -0,0 +1,18 @@++module Obsidian.Data where+++import Obsidian.Exp+import Obsidian.Memory+++--------------------------------------------------------------------------- +-- Data (should match up with storable instances for completeness)+-- Use this to ensure nonnestednes where required +---------------------------------------------------------------------------+class (Storable a, Choice a) => Data a+instance Scalar a => Data (Exp a)+instance (Data a, Data b) => Data (a,b)+instance (Data a, Data b, Data c) => Data (a,b,c)+instance (Data a, Data b, Data c, Data d) => Data (a,b,c,d) +-- Storable has up to triples
Obsidian/Exp.hs view
@@ -5,8 +5,9 @@ UndecidableInstances, OverlappingInstances, RankNTypes #-} +{-# LANGUAGE CPP #-} -{- Joel Svensson 2012 -} +{- Joel Svensson 2012, 2013, 2014 -} module Obsidian.Exp (module Obsidian.Exp,@@ -33,9 +34,11 @@ --type Data a = Exp a +-- Sizes of these are platform dependent type EInt = Exp Int type EWord = Exp Word +-- Types with platform independent size type EInt8 = Exp Int8 type EInt16 = Exp Int16 type EInt32 = Exp Int32@@ -56,7 +59,6 @@ type EW32 = Exp Word32 type EW64 = Exp Word64 - type EFloat = Exp Float type EDouble = Exp Double type EBool = Exp Bool @@ -70,61 +72,119 @@ sizeOf :: Exp a -> Int -- typeOf :: Exp a -> Type -- Good enough for me ... +#define SCALAR(t,s) instance Scalar t where { \+ sizeOf _ = s; \+ typeOf _ = t;} -instance Scalar Bool where - sizeOf _ = Storable.sizeOf (undefined :: Int)- typeOf _ = Bool +SCALAR(Bool,Storable.sizeOf (1 :: Int))+SCALAR(Int, Storable.sizeOf (1 :: Int)) +SCALAR(Int8, 1)+SCALAR(Int16,2)+SCALAR(Int32,4)+SCALAR(Int64,8) -instance Scalar Int where - sizeOf _ = Storable.sizeOf (undefined :: Int)- typeOf _ = Int+SCALAR(Float,Storable.sizeOf (1.0 :: Float))+SCALAR(Double,Storable.sizeOf (1.0 :: Double)) -instance Scalar Int8 where - sizeOf _ = 1- typeOf _ = Int8+SCALAR(Word,Storable.sizeOf (1 :: Word))+SCALAR(Word8, 1)+SCALAR(Word16,2)+SCALAR(Word32,4)+SCALAR(Word64,8) -instance Scalar Int16 where - sizeOf _ = 2- typeOf _ = Int16+#define SCALARVEC2(t,vt) instance Scalar (Vector2 t) where \+ {sizeOf _ = 2 * sizeOf (0 :: Exp t); \+ typeOf _ = vt} -instance Scalar Int32 where - sizeOf _ = 4- typeOf _ = Int32+#define SCALARVEC3(t,vt) instance Scalar (Vector3 t) where \+ {sizeOf _ = 3 * sizeOf (0 :: Exp t); \+ typeOf _ = vt} -instance Scalar Int64 where - sizeOf _ = 8 - typeOf _ = Int64- -instance Scalar Float where- sizeOf _ = Storable.sizeOf (undefined :: Float)- typeOf _ = Float- -instance Scalar Double where - sizeOf _ = 8 -- Storable.sizeOf (undefined :: Double) - typeOf _ = Double+#define SCALARVEC4(t,vt) instance Scalar (Vector4 t) where \+ {sizeOf _ = 4 * sizeOf (0 :: Exp t); \+ typeOf _ = vt} -instance Scalar Word where- sizeOf _ = Storable.sizeOf (undefined :: Word) - typeOf _ = Word- -instance Scalar Word8 where- sizeOf _ = 1- typeOf _ = Word8 -instance Scalar Word16 where - sizeOf _ = 2- typeOf _ = Word16- -instance Scalar Word32 where - sizeOf _ = 4 - typeOf _ = Word32+SCALARVEC2(Double,Vec2 Double)+SCALARVEC2(Float,Vec2 Float)+SCALARVEC3(Float,Vec3 Float)+SCALARVEC4(Float,Vec4 Float) ++SCALARVEC2(Int8,Vec2 Int8)+SCALARVEC3(Int8,Vec3 Int8)+SCALARVEC4(Int8,Vec4 Int8) ++SCALARVEC2(Int16,Vec2 Int16)+SCALARVEC3(Int16,Vec3 Int16)+SCALARVEC4(Int16,Vec4 Int16) ++SCALARVEC2(Int32,Vec2 Int32)+SCALARVEC3(Int32,Vec3 Int32)+SCALARVEC4(Int32,Vec4 Int32) +++---------------------------------------------------------------------------+-- Support CUDA Vector types+---------------------------------------------------------------------------++data Vector2 a = Vector2 !a !a+data Vector3 a = Vector3 !a !a !a+data Vector4 a = Vector4 !a !a !a !a++instance Show a => Show (Vector2 a) where+ show (Vector2 a b) = "Vector2 " ++ show a ++ " " ++ show b ++instance Show a => Show (Vector3 a) where+ show (Vector3 a b c ) = "Vector3 " ++ show a ++ " " ++ show b ++ " " ++ show c -instance Scalar Word64 where - sizeOf _ = 8 - typeOf _ = Word64+instance Show a => Show (Vector4 a) where+ show (Vector4 a b c d) = "Vector4 " ++ show a ++ " " ++ show b ++ " " ++ show c ++ " " ++ show d +instance Eq a => Eq (Vector2 a) where+ (Vector2 a b) == (Vector2 c d) = a == c && b == d +instance Eq a => Eq (Vector3 a) where+ (Vector3 a b c) == (Vector3 d e f) = a == d && b == e && c == f +instance Eq a => Eq (Vector4 a) where+ (Vector4 a b c d) == (Vector4 e f g h) = a == e && b == f && c == g && d == h +++class Vector v where+ getX :: (Scalar (v a), Scalar a) => Exp (v a) -> Exp a+ getY :: (Scalar (v a), Scalar a) => Exp (v a) -> Exp a+ getZ :: (Scalar (v a), Scalar a) => Exp (v a) -> Exp a+ getW :: (Scalar (v a), Scalar a) => Exp (v a) -> Exp a+++instance Vector Vector2 where+ getX (Literal (Vector2 x _)) = Literal x + getX v = UnOp GetX v+ getY (Literal (Vector2 _ y)) = Literal y + getY v = UnOp GetY v+ getZ v = error "getZ not allowed on Vector2"+ getW v = error "getW not allowed on Vector2"++instance Vector Vector3 where+ getX (Literal (Vector3 x _ _)) = Literal x+ getX v = UnOp GetX v+ getY (Literal (Vector3 _ y _)) = Literal y+ getY v = UnOp GetY v+ getZ (Literal (Vector3 _ _ z)) = Literal z + getZ v = UnOp GetZ v + getW v = error "getW not allowed on Vector3"++instance Vector Vector4 where+ getX (Literal (Vector4 x _ _ _)) = Literal x + getX v = UnOp GetX v+ getY (Literal (Vector4 _ y _ _)) = Literal y + getY v = UnOp GetY v+ getZ (Literal (Vector4 _ _ z _)) = Literal z + getZ v = UnOp GetZ v+ getW (Literal (Vector4 _ _ _ w)) = Literal w+ getW v = UnOp GetW v ++ --------------------------------------------------------------------------- -- Expressions data Exp a where@@ -140,14 +200,14 @@ can be translated into the CUDA/OpenCL specific concept later in the codegeneration -}- WarpSize :: Exp Word32+ -- WarpSize :: Exp Word32 - BlockDim :: DimSpec -> Exp Word32+ -- BlockDim :: DimSpec -> Exp Word32 - BlockIdx :: DimSpec - -> Exp Word32- ThreadIdx :: DimSpec- -> Exp Word32+ -- BlockIdx :: DimSpec + -- -> Exp Word32+ -- ThreadIdx :: DimSpec+ -- -> Exp Word32 Index :: Scalar a => (Name,[Exp Word32]) @@ -243,17 +303,22 @@ Int32ToWord32 :: Op (Int32 -> Word32) Word32ToInt32 :: Op (Word32 -> Int32) Word32ToFloat :: Op (Word32 -> Float)- Word32ToWord8 :: Op (Word32 -> Word8) - + Word32ToWord8 :: Op (Word32 -> Word8) ++ -- Vector Access+ GetX :: Vector v => Op (v a -> a) + GetY :: Vector v => Op (v a -> a)+ GetZ :: Vector v => Op (v a -> a)+ GetW :: Vector v => Op (v a -> a) --------------------------------------------------------------------------- -- helpers variable name = Index (name,[]) index name ix = Index (name,[ix]) -warpSize :: Exp Word32-warpSize = WarpSize+--warpSize :: Exp Word32+--warpSize = WarpSize --------------------------------------------------------------------------- -- Typecasts@@ -672,15 +737,23 @@ ifThenElse b e1' e2', ifThenElse b e1'' e2'') +instance (Choice a, Choice b, Choice c, Choice d)+ => Choice (a,b,c,d) where+ ifThenElse b (e1,e1',e1'',e1''') (e2,e2',e2'',e2''')+ = (ifThenElse b e1 e2,+ ifThenElse b e1' e2',+ ifThenElse b e1'' e2'',+ ifThenElse b e1''' e2''')+ --------------------------------------------------------------------------- -- Print Expressions --------------------------------------------------------------------------- printExp :: Scalar a => Exp a -> String-printExp (BlockIdx X) = "blockIdx.x"-printExp (ThreadIdx X) = "threadIdx.x"-printExp (BlockDim X) = "blockDim.x"+--printExp (BlockIdx X) = "blockIdx.x"+--printExp (ThreadIdx X) = "threadIdx.x"+--printExp (BlockDim X) = "blockDim.x" printExp (Literal a) = show a printExp (Index (name,[])) = name printExp (Index (name,es)) = @@ -714,28 +787,69 @@ printOp Sin = " Sin " printOp Cos = " Cos "+printOp Sqrt = " Sqrt " printOp BitwiseAnd = " & " printOp BitwiseOr = " | " printOp BitwiseXor = " ^ " printOp BitwiseNeg = " ~ " +printOp GetX = "getX"+printOp GetY = "getY"+printOp GetZ = "getZ"+printOp GetW = "getW" + --------------------------------------------------------------------------- -- Internal exp (not a GADT) --------------------------------------------------------------------------- data IExp = IVar Name Type- | IBlockIdx DimSpec- | IThreadIdx DimSpec- | IBlockDim DimSpec- | IGridDim DimSpec+ -- | IBlockIdx DimSpec+ -- | IThreadIdx DimSpec+ -- | IBlockDim DimSpec+ -- | IGridDim DimSpec +-- Break out: Values and Vectors this is too messy. | IBool Bool | IInt8 Int8 | IInt16 Int16 | IInt32 Int32 | IInt64 Int64 | IWord8 Word8 | IWord16 Word16 | IWord32 Word32 | IWord64 Word64 | IFloat Float | IDouble Double +-- Vector Types (Clean this up, somehow)+ | IFloat2 Float Float+ | IFloat3 Float Float Float+ | IFloat4 Float Float Float Float+ | IDouble2 Double Double+ | IInt8_2 Int8 Int8+ | IInt8_3 Int8 Int8 Int8+ | IInt8_4 Int8 Int8 Int8 Int8 + | IInt16_2 Int16 Int16+ | IInt16_3 Int16 Int16 Int16+ | IInt16_4 Int16 Int16 Int16 Int16 + | IInt32_2 Int32 Int32+ | IInt32_3 Int32 Int32 Int32+ | IInt32_4 Int32 Int32 Int32 Int32 + | IInt64_2 Int64 Int64+ | IInt64_3 Int64 Int64 Int64+ | IInt64_4 Int64 Int64 Int64 Int64 + | IWord8_2 Word8 Word8+ | IWord8_3 Word8 Word8 Word8+ | IWord8_4 Word8 Word8 Word8 Word8 + | IWord16_2 Word16 Word16+ | IWord16_3 Word16 Word16 Word16+ | IWord16_4 Word16 Word16 Word16 Word16 + | IWord32_2 Word32 Word32+ | IWord32_3 Word32 Word32 Word32+ | IWord32_4 Word32 Word32 Word32 Word32 + | IWord64_2 Word64 Word64+ | IWord64_3 Word64 Word64 Word64+ | IWord64_4 Word64 Word64 Word64 Word64 + + + -- ... much more to add. ++-- Operations | IIndex (IExp,[IExp]) Type | ICond IExp IExp IExp Type | IBinOp IBinOp IExp IExp Type@@ -753,6 +867,7 @@ deriving (Eq, Ord, Show) data IUnOp = IBitwiseNeg | INot+ | IGetX | IGetY | IGetZ | IGetW deriving (Eq, Ord, Show) @@ -788,10 +903,15 @@ binOpToIBinOp ShiftR = IShiftR unOpToIUnOp :: Op t -> IUnOp-unOpToIUnOp BitwiseNeg = IBitwiseNeg-unOpToIUnOp Not = INot +unOpToIUnOp BitwiseNeg = IBitwiseNeg+unOpToIUnOp Not = INot+unOpToIUnOp GetX = IGetX+unOpToIUnOp GetY = IGetY+unOpToIUnOp GetZ = IGetZ+unOpToIUnOp GetW = IGetW + --------------------------------------------------------------------------- -- Turn Exp a to IExp with type information. ---------------------------------------------------------------------------@@ -835,7 +955,7 @@ expToIExp (Literal a) = IDouble a expToIExp a = expToIExpGeneral a --- This is strange. +-- This is strange. (... WHY??? (2014)) instance ExpToIExp Word where expToIExp (Literal a) = IWord32 (fromIntegral a) expToIExp a = expToIExpGeneral a @@ -856,13 +976,51 @@ expToIExp (Literal a) = IWord64 a expToIExp a = expToIExpGeneral a +---------------------------------------------------------------------------+-- Vector Exp to IExp+---------------------------------------------------------------------------+#define ETOIEVEC2(t,ct) instance ExpToIExp (Vector2 t) where \+ {expToIExp (Literal (Vector2 a b)) = ct a b; \+ expToIExp a = expToIExpGeneral a} +#define ETOIEVEC3(t,ct) instance ExpToIExp (Vector3 t) where \+ {expToIExp (Literal (Vector3 a b c)) = ct a b c; \+ expToIExp a = expToIExpGeneral a}++#define ETOIEVEC4(t,ct) instance ExpToIExp (Vector4 t) where \+ {expToIExp (Literal (Vector4 a b c d)) = ct a b c d; \+ expToIExp a = expToIExpGeneral a}++-- CPP string concatenation seems to not work with GHC CPP. +-- So this gets a bit more wordy. ++ETOIEVEC2(Float,IFloat2) +ETOIEVEC3(Float,IFloat3)+ETOIEVEC4(Float,IFloat4) ++ETOIEVEC2(Int8,IInt8_2) +ETOIEVEC3(Int8,IInt8_3)+ETOIEVEC4(Int8,IInt8_4) ++ETOIEVEC2(Int16,IInt16_2) +ETOIEVEC3(Int16,IInt16_3)+ETOIEVEC4(Int16,IInt16_4) ++ETOIEVEC2(Int32,IInt32_2) +ETOIEVEC3(Int32,IInt32_3)+ETOIEVEC4(Int32,IInt32_4) ++instance ExpToIExp (Vector2 Double) where+ expToIExp (Literal (Vector2 a b)) = IDouble2 a b+ expToIExp a = expToIExpGeneral a +---------------------------------------------------------------------------+-- translation from Exp to IExp in the general case. expToIExpGeneral :: ExpToIExp a => Exp a -> IExp-expToIExpGeneral WarpSize = IVar "warpsize" Word32 -expToIExpGeneral (BlockIdx d) = IBlockIdx d-expToIExpGeneral (BlockDim d) = IBlockDim d -expToIExpGeneral (ThreadIdx d) = IThreadIdx d+--expToIExpGeneral WarpSize = IVar "warpsize" Word32 +--expToIExpGeneral (BlockIdx d) = IBlockIdx d+--expToIExpGeneral (BlockDim d) = IBlockDim d +--expToIExpGeneral (ThreadIdx d) = IThreadIdx d expToIExpGeneral e@(Index (name,[])) = IVar name (typeOf e) expToIExpGeneral e@(Index (name,xs))
Obsidian/Force.hs view
@@ -1,13 +1,20 @@+ {-# LANGUAGE ScopedTypeVariables #-} {-# LANGUAGE FlexibleInstances #-} {-# LANGUAGE FlexibleContexts #-} {-# LANGUAGE UndecidableInstances #-}-{-# LANGUAGE TypeFamilies #-} {-# LANGUAGE MultiParamTypeClasses #-}+{-# LANGUAGE TypeOperators #-}+{-# LANGUAGE GADTs #-}+----------------------------------------+{- LANGUAGE KindSignatures -}+{-# LANGUAGE TypeFamilies #-} +{-# LANGUAGE DataKinds #-}+{-# LANGUAGE ConstraintKinds #-} -{- Joel Svensson 2012, 2013 +{- Joel Svensson 2012..2015 Notes: 2014-03-28: Changed API.@@ -23,105 +30,135 @@ -} -module Obsidian.Force (Forceable, force, forcePull, unsafeForce) where--- Write, force, forcePull, unsafeForce, unsafeWritePush) where +{-+ Warning: This module is full of magic! +-} ++module Obsidian.Force ( unsafeWritePush+ , unsafeWritePull+ , compute_+ , computePull_+ , ComputeAs(..)+ , Compute) where++ import Obsidian.Program import Obsidian.Exp import Obsidian.Array-import Obsidian.Types-import Obsidian.Globs import Obsidian.Memory import Obsidian.Names+import Obsidian.Data import Data.Word -import Control.Monad+ ---------------------------------------------------------------------------+--+-------------------------------------------------------------------------++-- | Compute constraint. +type Compute t = (Write t, t *<=* Block)++-- | Arrays can be computed at level t if level t allows compute.+class Compute t => ComputeAs t a where+ compute :: Data e => a Word32 e -> Program t (Pull Word32 e)+ +instance Compute t => ComputeAs t Pull where+ compute = computePull_ ++{- + The key to this instance is that when matching up instances the+ -whatever matcher that does that-+ matches only against the head, ignoring the constraint.+ meaning that all variations of t, t1 is caught by this+ instance. Though, those where t and t1 are not equal+ a type error is the result (rather than a missing instance).++ This means that the constraint "Compute Block (Push Thread)"+ matches this instance, but is a type error.+-} +instance (t ~ t1, Compute t) => ComputeAs t (Push t1) where+ compute = compute_++compute_ :: (Data a, Compute t)+ => Push t Word32 a -> Program t (Pull Word32 a) +compute_ arr = do+ rval <- unsafeWritePush False arr+ sync+ return rval++computePull_ :: (t *<=* Block, Data a, Compute t)+ => Pull Word32 a -> Program t (Pull Word32 a) +computePull_ arr = + if (len arr <= 32)+ then do+ rval <- unsafeWritePush True parr+ return rval+ else do+ rval <- unsafeWritePush False parr + sync+ return rval+ where parr = push arr++ +--------------------------------------------------------------------------- -- Force local (requires static lengths!) --------------------------------------------------------------------------- class Write t where unsafeWritePush :: Storable a => Bool -> Push t Word32 a -> Program t (Pull Word32 a)- -- unsafeWritePull :: MemoryOps a => Pull Word32 a -> Program t (Pull Word32 a) - -instance Write Warp where- unsafeWritePush volatile p =++-- What to do about volatile here?+-- Ignoring that parameter for now.+-- Thought: It does not matter.+-- Thought: Is this function correct at all?+-- What happens if a thread program allocates memory+-- DONE: The above problem has been fixed! +instance Write Thread where+ unsafeWritePush _ p = do+ (snames :: Names a) <- names "arr" ++ -- Here I know that this pattern match will succeed let n = len p- names <- names "arr"- allocateVolatileArray names n+ + allocateArray snames n+ p <: threadAssignArray snames (variable "tid") n + + return $ threadPullFrom snames (variable "tid") n++instance Write Warp where+ unsafeWritePush _ p =+ do+ let n = len p+ noms <- names "arr"+ allocateVolatileArray noms n - p <: warpAssignArray names (variable "warpID") n - return $ warpPullFrom names (variable "warpID") n+ p <: warpAssignArray noms (variable "warpID") n + return $ warpPullFrom noms (variable "warpID") n instance Write Block where unsafeWritePush volatile p = do let n = len p- names <- names "arr"+ noms <- names "arr" if (volatile)- then allocateVolatileArray names n- else allocateArray names n+ then allocateVolatileArray noms n+ else allocateArray noms n - p <: assignArray names - return $ pullFrom names n+ p <: assignArray noms + return $ pullFrom noms n -instance Write Thread where- unsafeWritePush volatile p =- do- (snames :: Names a) <- names "arr" - -- Here I know that this pattern match will succeed- let n = len p- - allocateArray snames n- p <: assignArray snames - - return $ pullFrom snames n ------------------------------------------------------------------------------ Force functions +-- unsafe! ------------------------------------------------------------------------------- | It is possible to force at level T if we can Write and Sync at that level. -class (Sync t, Write t) => Forceable t--instance Forceable Thread-instance Forceable Warp-instance Forceable Block----- | force turns a @Push@ array to a @Program@ generating a @Pull@ array.--- The returned array represents reading from an array manifest in memory.-force :: (Storable a, Forceable t)- => Push t Word32 a -> Program t (Pull Word32 a) -force arr = do- rval <- unsafeWritePush False arr- sync- return rval---- | Make a @Pull@ array manifest in memory. -forcePull :: (Storable a, Forceable t, Pushable t) - => Pull Word32 a -> Program t (Pull Word32 a) -forcePull = unsafeForce . push ---- | unsafeForce is dangerous on @Push@ arrays as it does not--- insert synchronization primitives. -unsafeForce :: (Storable a, Forceable t) =>- Push t Word32 a -> Program t (Pull Word32 a) -unsafeForce arr = - if (len arr <= 32)- then do- rval <- unsafeWritePush True arr- return rval- else do- rval <- unsafeWritePush False arr - sync- return rval -+unsafeWritePull :: (t *<=* Block, Write t, Storable a) => Bool -> Pull Word32 a -> Program t (Pull Word32 a)+unsafeWritePull t = unsafeWritePush t . push
Obsidian/Library.hs view
@@ -1,4 +1,4 @@-{- Joel Svensson 2012, 2013, 2014 +{- Joel Svensson 2012..2015 Mary Sheeran 2012 Notes:@@ -17,15 +17,86 @@ ScopedTypeVariables, TypeFamilies, GADTs #-}+{-# LANGUAGE TypeOperators #-}+{-# LANGUAGE FlexibleContexts #-}+{-# LANGUAGE MultiParamTypeClasses #-}+{-# LANGUAGE UndecidableInstances #-} -module Obsidian.Library where +module Obsidian.Library+ (logBaseI+ , reverse+ , splitAt+ , splitUp+ , coalesce+ , halve+ , evenOdds+ , evens+ , odds+ , everyNth+ , singleton+ , generate+ , last+ , first+ , take+ , drop+ , head+ , tail + , fold1+ , shiftLeft+ , unzip+ , zip+ , unzip3+ , zip3+ , zipWith+ , zipWith3+ , pair+ , unpair+ , unsafeBinSplit+ , binSplit + , concP+ , unpairP + , load -- RENAME THIS + , store -- RENAME THIS+ -- Hierarchy programming + , asThread+ , asThreadMap+ , asGrid+ , asGridMap+ , AsWarp(..)+ , AsBlock(..)+ , liftPar -- generic hierarchy programming+ , liftSeq -- generic hierarchy programming+ , liftIn -- generic hierarchy programming + -- Repeat a program+ , rep + + -- Executing programs+ , ExecProgram(..) +-- , ExecBlock(..)+-- , ExecWarp(..)+-- , ExecThread(..)+ , execBlock+ , execBlock'+ , execThread+ , execThread'+ , execWarp+ , execWarp' +-- , execBlock+-- , execThread+-- , execWarp+++ -- Leftovers from past days + , singletonPush+ , runPush + )where + import Obsidian.Array import Obsidian.Exp import Obsidian.Program-import Obsidian.Types--import Obsidian.Force+import Obsidian.Data+import Obsidian.Mutable -- needed for threadsPerBlock analysis -- import qualified Obsidian.CodeGen.Program as P @@ -35,18 +106,33 @@ import Data.Bits import Data.Word -import Prelude hiding (splitAt,zipWith,replicate,reverse)+import Prelude hiding (splitAt,zipWith,replicate,reverse,unzip,zip,zip3,unzip3,zipWith3, last, take, drop, head, tail) + ---------------------------------------------------------------------------+-- Helper +---------------------------------------------------------------------------+logBaseI :: Integral a => a -> a -> a+logBaseI b x+ = if x < b+ then 0+ else+ let+ l = 2 * logBaseI (b*b) x+ doDiv x l = if x < b then l else doDiv (x`div`b) (l+1)+ in+ doDiv (x`div`(b^l)) l+++--------------------------------------------------------------------------- -- Reverse an array by indexing in it backwards --------------------------------------------------------------------------- -- | Reverses a Pull array.-reverse :: ASize l => Pull l a -> Pull l a -reverse arr = mkPull n (\ix -> arr ! ((sizeConv m) - ix)) - where m = n-1- n = len arr- +reverse :: (Array array, ArrayLength array, ASize l) => array l a -> array l a+reverse arr = ixMap (\ix -> (sizeConv m) - ix) arr+ where m = n-1+ n = len arr --------------------------------------------------------------------------- -- splitAt (name clashes with Prelude.splitAt) ---------------------------------------------------------------------------@@ -65,8 +151,9 @@ n = len arr n2 = n `div` 2 + -- | Splits a Pull array into chunks of size n. Result is a Pull of Pull arrays.-splitUp :: ASize l => Word32 -> Pull l a -> Pull l (SPull a)+splitUp :: (ASize l, ASize s, Integral s) => s -> Pull l a -> Pull l (Pull s a) splitUp n arr {-(Pull m ixf)-} = mkPull (len arr `div` fromIntegral n) $ \i -> mkPull n $ \j -> arr ! (i * (sizeConv n) + j) @@ -78,7 +165,8 @@ mkPull s $ \i -> mkPull n $ \j -> arr ! (i + (sizeConv s) * j) where s = len arr `div` fromIntegral n- ++ --------------------------------------------------------------------------- -- elements at even indices to fst output, odd to snd. ---------------------------------------------------------------------------@@ -114,7 +202,7 @@ singleton :: (Array a, ASize l) => e -> a l e singleton a = replicate 1 a --- | Generate a pull or push array usign a function from Index to element.+-- | Generate a pull or push array using a function from Index to element. generate :: (Functor (a s), Array a, ASize s) => s -> (EWord32 -> b) -> a s b generate n f = fmap f (iota n) @@ -141,6 +229,17 @@ drop n arr = setSize (len arr - n) $ ixMap (\ix -> ix + sizeConv n) arr ---------------------------------------------------------------------------+-- Head and Tail on pull arrays+---------------------------------------------------------------------------++head :: ASize l => Pull l a -> a+head arr = arr ! 0++tail :: ASize l => Pull l a -> Pull l a+tail = drop 1 +++--------------------------------------------------------------------------- -- fold (sequential , unrolled) --------------------------------------------------------------------------- -- | Fold a nonempty pull array using a given operator. The result a singleton array (push or pull). @@ -153,40 +252,9 @@ --------------------------------------------------------------------------- -- Shift arrays ------------------------------------------------------------------------------ shiftRight :: (ASize l, Choice a) => Word32 -> a -> Pull l a -> Pull l a--- shiftRight dist elt arr = setSize (len arr)--- $ replicate (fromIntegral dist) elt `conc` arr---- -- TODO: incorrect! --- shiftLeft :: (ASize l, Choice a) => Word32 -> a -> Pull l a -> Pull l a--- shiftLeft dist elt arr = mkPull (len arr)--- $ \ix -> (arr `conc` replicate (fromIntegral dist) elt)--- ! (ix + fromIntegral dist) - ---- opposite of evenOdds ---shuffle :: ASize l => PT t -> Pull l a -> Pull l a -> Push t l a---shuffle Block a1 a2 =--- Push (len a1 + len a2) $--- \ wf -> ForAll (sizeConv (len a1)) $--- \ tid -> do--- wf (a1 ! tid) (tid * 2) --- wf (a2 ! tid) (tid * 2 + 1) --------------------------------------------------------------------------------- Concatenate the arrays------------------------------------------------------------------------------- -- | Concatenate two Pull arrays (Potentially inefficient).--- conc :: (ASize l, Choice a) => Pull l a -> Pull l a -> Pull l a --- conc a1 a2 = mkPull (n1+n2)--- $ \ix -> ifThenElse (ix <* (sizeConv n1)) --- (a1 ! ix) --- (a2 ! (ix - (sizeConv n1)))--- where --- n1 = len a1--- n2 = len a2 -+shiftLeft :: ASize l => Word32 -> Pull l a -> Pull l a+shiftLeft dist arr = setSize (len arr - (fromIntegral dist))+ $ ixMap (\ix -> ix + (fromIntegral dist)) arr --------------------------------------------------------------------------- -- zipp unzipp@@ -251,19 +319,54 @@ (fst (arr ! (ix `shiftR` 1))) (snd (arr ! (ix `shiftR` 1)))) +-- | Triple up consecutive elements in a Pull array.+triple :: ASize l => Pull l a -> Pull l (a,a,a)+triple arr =+ mkPull (len arr `div` 3) $ \ix ->+ (arr ! (ix*3), arr ! (ix*3+1), arr ! (ix*3+2)) +-- | Flatten a Pull array of triples. +untriple :: ASize l => Choice a => Pull l (a,a,a) -> Pull l a+untriple arr =+ mkPull (3*len arr) $ \ix ->+ let (k,j) = divMod ix 3+ (a0,a1,a2) = arr ! k+ in ifThenElse (j ==* 0) a0 $+ ifThenElse (j ==* 1) a1 a2+++-- | Quadruple up consecutive elements in a Pull array.+quadruple :: ASize l => Pull l a -> Pull l (a,a,a,a)+quadruple arr =+ mkPull (len arr `div` 4) $ \ix ->+ (arr ! (ix*4), arr ! (ix*4+1), arr ! (ix*4+2), arr ! (ix*4+3))++-- | Flatten a Pull array of triples. +unquadruple :: ASize l => Choice a => Pull l (a,a,a,a) -> Pull l a+unquadruple = unpair . unpair . fmap (\(a0,a1,a2,a3) -> ((a0,a1), (a2,a3)))+++ --------------------------------------------------------------------------- -- twoK (untested for proper functionality) --------------------------------------------------------------------------- -- | Recursively split an array in the middle. Apply an array to array computation--- on each part. @binSplit 3@ divides the array into 8 pieces. -binSplit ::Int+-- on each part. @binSplit 3@ divides the array into 8 pieces.+-- UNSAFE+unsafeBinSplit :: Int -> (Pull Word32 a -> Pull Word32 b) -> Pull Word32 a -> Pull Word32 b -binSplit = twoK+unsafeBinSplit = twoK +binSplit :: Data a+ => Int+ -> (Pull Word32 a -> Pull Word32 b)+ -> Mutable Shared Word32 a+ -> Pull Word32 b+binSplit n f = unsafeBinSplit n f . mutableToPull + -- See if this should be specifically for Static size pull arrays twoK :: Int -> (Pull Word32 a -> Pull Word32 b) -> Pull Word32 a -> Pull Word32 b twoK 0 f = f -- divide 0 times and apply f@@ -276,7 +379,9 @@ nl2 = len (f (mkPull m (\j -> arr ! variable "X"))) lt = nl2 `shiftL` n in arr'+ + --------------------------------------------------------------------------- -- *** PUSHY LIBRARY *** --- ---------------------------------------------------------------------------@@ -287,7 +392,7 @@ -- | Concatenate two push arrays. concP :: ASize l- => Push t l a -> Push t l a -> Push t l a + => Push t l a -> Push t l a -> Push t l a concP p1 p2 = mkPush (n1 + n2) $ \wf -> do@@ -295,9 +400,20 @@ p2 <: \a i -> wf a (sizeConv n1 + i) where n1 = len p1- n2 = len p2 + n2 = len p2 +-- | Flatten a Pull array of pairs. Result is a push array+unpairP :: ASize l => Choice a => Push t l (a,a) -> Push t l a+unpairP arr =+ mkPush (2 * len arr) $ \ wf ->+ do + -- even iterations+ arr <: \ (a,_) i -> wf a ((sizeConv i) `shiftL` 1)+ -- Odd iterations + arr <: \ (_,b) i -> wf b ((sizeConv i) `shiftL` 1 + 1) ++ -- Implement unpair on pusharrays, -- Impement for more tuples than pairs. --unpair :: ASize l => Choice a => Pull l (a,a) -> Pull l a@@ -340,13 +456,16 @@ --------------------------------------------------------------------------- -- Parallel concatMap -----------------------------------------------------------------------------pConcatMap f = pConcat . fmap f-pUnCoalesceMap f = pUnCoalesce . fmap f-pConcatMapJoin f = pConcat . fmap (runPush.f)-pUnCoalesceMapJoin f = pUnCoalesce . fmap (runPush.f)-pCoalesceMap n f = pUnCoalesce . fmap f . coalesce n-pSplitMap n f = pConcat . fmap f . splitUp n+-- pConcatMap :: ASize s+-- => (a -> SPush (Below t) b)+-- -> Pull s a+-- -> Push t s b +-- pConcatMap f = pConcat . fmap f+-- pUnCoalesceMap f = pUnCoalesce . fmap f+-- pConcatMapJoin f = pConcat . fmap (runPush.f)+-- pUnCoalesceMapJoin f = pUnCoalesce . fmap (runPush.f)+-- pCoalesceMap n f = pUnCoalesce . fmap f . coalesce n+-- pSplitMap n f = pConcat . fmap f . splitUp n --------------------------------------------------------------------------- -- Step into the Hierarchy by distributing a@@ -354,10 +473,251 @@ -- at a specific level in the Hierarchy. --------------------------------------------------------------------------- +-- "Compute as" family of functions ++liftPar :: ASize l => Pull l (SPush t a) -> Push (Step t) l a+liftPar = pConcat++liftSeq :: ASize l => Pull l (SPush t a) -> Push t l a+liftSeq = sConcat ++liftIn :: (t *<=* Block, ASize l)+ => Pull l (SPush Thread b)+ -> Push t l b +liftIn = tConcat++---------------------------------------------------------------------------+-- AsBlock+---------------------------------------------------------------------------+-- Rename this to asBlock, asGrid, asThread+class (t *<=* Block) => AsBlock t where+ asBlock :: SPull (SPush t a) ->+ SPush Block a+ asBlockMap :: (a -> SPush t b) + -> SPull a+ -> SPush Block b ++instance AsBlock Thread where + asBlock = tConcat+ asBlockMap f = tConcat . fmap f+ +instance AsBlock Warp where+ asBlock = pConcat+ asBlockMap f = pConcat . fmap f+ +instance AsBlock Block where+ asBlock = sConcat+ asBlockMap f = sConcat . fmap f ++---------------------------------------------------------------------------+-- AsWarp+--------------------------------------------------------------------------- +class (t *<=* Warp) => AsWarp t where+ asWarp :: SPull (SPush t a) ->+ SPush Warp a + asWarpMap :: (a -> SPush t b) + -> SPull a+ -> SPush Warp b ++ +instance AsWarp Thread where+ asWarp = tConcat+ asWarpMap f = tConcat . fmap f+ +instance AsWarp Warp where+ asWarp = sConcat+ asWarpMap f = sConcat . fmap f++---------------------------------------------------------------------------+-- LiftThread+---------------------------------------------------------------------------+ +asThread :: ASize l+ => Pull l (SPush Thread b)+ -> Push Thread l b+asThread = tConcat++asThreadMap :: (a -> SPush Thread b) + -> SPull a+ -> SPush Thread b+asThreadMap f = tConcat . fmap f+++---------------------------------------------------------------------------+-- AsGrid+---------------------------------------------------------------------------++asGrid :: ASize l => Pull l (SPush Block a)+ -> Push Grid l a+asGrid = pConcat++asGridMap :: ASize l => (a -> SPush Block b) + -> Pull l a+ -> Push Grid l b +asGridMap f = pConcat . fmap f+++---------------------------------------------------------------------------+-- Repeat a program+---------------------------------------------------------------------------++-- | Repeat a program (iterate it) +rep :: Word32 -> (a -> Program t a) -> a -> Program t a+rep 0 _ a = return a +rep n prg a = do+ b <- rep (n-1) prg a+ prg b + ++---------------------------------------------------------------------------+-- RunPush +---------------------------------------------------------------------------++class ExecProgram t a where+ exec :: Data e+ => Program t (a Word32 e)+ -> Push t Word32 e ++instance (t *<=* Block) => ExecProgram t Pull where+ exec = runPush . liftM push ++--instance ExecProgram t (Push t) where+-- exec = runPush++-- Here we also want the type error behaviour.+-- it is a type error to try to "execute" a push t at any level different from t +instance (t ~ t1) => ExecProgram t (Push t1) where+ exec = runPush +++execThread :: (ExecProgram Thread a, Data e)+ => Program Thread (a Word32 e)+ -> Push Thread Word32 e +execThread = exec++execThread' :: Data a => Program Thread a -> SPush Thread a +execThread' = singletonPush++execBlock :: (ExecProgram Block a, Data e)+ => Program Block (a Word32 e)+ -> Push Block Word32 e +execBlock = exec ++execBlock' :: Data a => Program Block a -> SPush Block a+execBlock' = singletonPush ++execWarp :: (ExecProgram Warp a, Data e)+ => Program Warp (a Word32 e)+ -> Push Warp Word32 e+execWarp = exec ++execWarp' :: Data a => Program Warp a -> SPush Warp a+execWarp' = singletonPush +++-- class ExecThread a where+-- execThread :: Data e+-- => Program Thread (a Word32 e)+-- -> Push Thread Word32 e+ +-- instance ExecThread (Push Thread) where+-- execThread = runPush++-- instance ExecThread Pull where+-- execThread = execThread . liftM pushThread+++-- class ExecBlock a where+-- execBlock :: Data e+-- => Program Block (a Word32 e)+-- -> Push Block Word32 e+++-- instance ExecBlock (Push Block) where+-- execBlock = runPush++-- instance ExecBlock Pull where+-- execBlock = execBlock . liftM pushBlock++-- class ExecWarp a where+-- execWarp :: Data e+-- => Program Warp (a Word32 e)+-- -> Push Warp Word32 e++-- instance ExecWarp (Push Warp) where+-- execWarp = runPush++-- instance ExecWarp Pull where+-- execWarp = execWarp . liftM pushWarp+++-- | Fuses the program that computes a Push array into the Push array. +runPush :: Program t (Push t s a) -> Push t s a+runPush prg =+ mkPush n $ \wf -> do+ parr <- prg+ parr <: wf+ -- It is a bit scary that I need to "evaluate" programs here. + where n = len $ fst $ runPrg 0 prg++-- | Lifts @runPush@ to one input functions.+runPush1 :: (a -> Program t (Push t s b)) -> a -> Push t s b+runPush1 f a = runPush (f a)++-- | Lifts @runPush@ to two input functions.+runPush2 :: (a -> b -> Program t (Push t s c)) -> a -> b -> Push t s c+runPush2 f a b = runPush (f a b) ++-- | Converts a program computing a pull Array to a Push array+runPull :: (t *<=* Block, ASize s) => Program t (Pull s a) -> Push t s a+runPull = runPush . liftM push ++-- | Lifts @runPull@ to one input functions.+runPull1 :: (t *<=* Block, ASize s) => (a -> Program t (Pull s b)) -> a -> Push t s b+runPull1 f a = runPull (f a)++-- | Lifts @runPull@ to two input functions.+runPull2 :: (t *<=* Block, ASize s) => (a -> b -> Program t (Pull s c)) -> a -> b -> Push t s c+runPull2 f a b = runPull (f a b)++---------------------------------------------------------------------------+-- +---------------------------------------------------------------------------+pushPrg :: (t *<=* Block) => Program t a -> SPush t a+pushPrg = singletonPush+++---------------------------------------------------------------------------+-- Singleton push arrays +---------------------------------------------------------------------------++-- Danger! use only with Scalar a's +-- -- | Create a singleton Push array.+--singletonPush :: a -> SPush t a+--singletonPush = singletonPushP . return ++-- | Monadic version of @singleton@.+singletonPush :: (t *<=* Block) => Program t a -> SPush t a+singletonPush prg =+ mkPush 1 $ \wf -> do+ a <- prg+ forAll 1 $ \_ -> + wf a 0+++++++---------------------------------------------------------------------------+-- Old stuff that should nolonger be exported!+-- * It is still used internally +---------------------------------------------------------------------------+ -- | A way to enter into the hierarchy -- A bunch of Thread computations, spread across the threads of either -- a Warp, block or grid. (or performed sequentially in a single thread) -tConcat :: ASize l+tConcat :: (t *<=* Block, ASize l) => Pull l (SPush Thread b) -> Push t l b tConcat arr =@@ -371,7 +731,7 @@ s = len (arr ! 0) --(f (variable "tid")) -- arr -- | Variant of @tConcat@. -tDistribute :: ASize l+tDistribute :: (t *<=* Block, ASize l) => l -> (EWord32 -> SPush Thread b) -> Push t l b@@ -392,7 +752,10 @@ rn = len (arr ! 0) -- All arrays are same length -- | Distribute work across the parallel resources at a given level of the GPU hierarchy-pDistribute :: ASize l => l -> (EWord32 -> SPush t a) -> Push (Step t) l a+pDistribute :: ASize l+ => l+ -> (EWord32 -> SPush t a)+ -> Push (Step t) l a pDistribute n f = pConcat (mkPull n f) -- | Sequential concatenation of a Pull of Push.@@ -401,13 +764,14 @@ mkPush (n * fromIntegral rn) $ \wf -> do seqFor (sizeConv n) $ \bix ->- let p = arr ! bix -- (Push _ p) = arr ! bix+ let p = arr ! bix wf' a ix = wf a (bix * sizeConv rn + ix) in p <: wf' where n = len arr rn = len $ arr ! 0 + -- | Variant of sConcat. sDistribute :: ASize l => l -> (EWord32 -> SPush t a) -> Push t l a sDistribute n f = sConcat (mkPull n f) @@ -415,7 +779,9 @@ -- pUnCoalesce adapted from Niklas branch. -- | Combines work that was distributed in a Coalesced way. -- | Applies a permutation on stores.-pUnCoalesce :: ASize l => Pull l (SPush t a) -> Push (Step t) l a+pUnCoalesce :: ASize l+ => Pull l (SPush t a)+ -> Push (Step t) l a pUnCoalesce arr = mkPush (n * fromIntegral rn) $ \wf -> distrPar (sizeConv n) $ \bix ->@@ -427,61 +793,3 @@ rn = len $ arr ! 0 s = sizeConv rn g wf a i = wf a (i `div` s + (i`mod`s)*(sizeConv n))-------------------------------------------------------------------------------- RunPush -------------------------------------------------------------------------------- | Fuses the program that computes a Push array into the Push array. -runPush :: Program t (Push t s a) -> Push t s a-runPush prg =- mkPush n $ \wf -> do- parr <- prg- parr <: wf- -- It is a bit scary that I need to "evaluate" programs here. - where n = len $ fst $ runPrg 0 prg---- | Lifts @runPush@ to one input functions.-runPush1 :: (a -> Program t (Push t s b)) -> a -> Push t s b-runPush1 f a = runPush (f a)---- | Lifts @runPush@ to two input functions.-runPush2 :: (a -> b -> Program t (Push t s c)) -> a -> b -> Push t s c-runPush2 f a b = runPush (f a b) ---- | Converts a program computing a pull Array to a Push array-runPull :: (Pushable t, ASize s) => Program t (Pull s a) -> Push t s a-runPull = runPush . liftM push ---- | Lifts @runPull@ to one input functions.-runPull1 :: (Pushable t, ASize s) => (a -> Program t (Pull s b)) -> a -> Push t s b-runPull1 f a = runPull (f a)---- | Lifts @runPull@ to two input functions.-runPull2 :: (Pushable t, ASize s) => (a -> b -> Program t (Pull s c)) -> a -> b -> Push t s c-runPull2 f a b = runPull (f a b)-------------------------------------------------------------------------------- -----------------------------------------------------------------------------pushPrg :: Program t a -> SPush t a-pushPrg = singletonPush--------------------------------------------------------------------------------- Singleton push arrays -------------------------------------------------------------------------------- Danger! use only with Scalar a's --- -- | Create a singleton Push array.---singletonPush :: a -> SPush t a---singletonPush = singletonPushP . return ---- | Monadic version of @singleton@.-singletonPush :: Program t a -> SPush t a-singletonPush prg =- mkPush 1 $ \wf -> do- a <- prg- forAll 1 $ \ix -> - wf a 0-
Obsidian/Memory.hs view
@@ -16,7 +16,6 @@ import Obsidian.Program import Obsidian.Exp import Obsidian.Types-import Obsidian.Globs import Obsidian.Array -- Importing this feels a bit strange. import Obsidian.Names @@ -34,16 +33,15 @@ -- Array operations assignArray :: Names a -> a -> Exp Word32 -> Program Thread () allocateArray :: Names a -> Word32 -> Program t ()- pullFrom :: Names a -> Word32 -> Pull Word32 a+ pullFrom :: ASize s => Names a -> s -> Pull s a - -- Scalar operations assignScalar :: Names a -> a -> Program Thread ()- allocateScalar :: Names a -> Program t () + allocateScalar :: Names a -> Program t ()+ allocateSharedScalar :: Names a -> Program t () readFrom :: Names a -> a - -- Warp level operations warpAssignArray :: Names a -> EWord32@@ -52,6 +50,14 @@ -> EWord32 -> Program Thread () warpPullFrom :: Names a -> EWord32 -> Word32 -> Pull Word32 a++ threadAssignArray :: Names a+ -> EWord32+ -> Word32+ -> a+ -> EWord32+ -> Program Thread ()+ threadPullFrom :: Names a -> EWord32 -> Word32 -> Pull Word32 a -- Extra allocateVolatileArray :: Names a -> Word32 -> Program t ()@@ -76,17 +82,27 @@ -- Scalar ops allocateScalar (Single name) =- Declare name (typeOf (undefined :: Exp a)) + Declare name (typeOf (undefined :: Exp a))+ allocateSharedScalar (Single name) =+ Declare name (Shared $ typeOf (undefined :: Exp a))+ assignScalar (Single name) a = Assign name [] a readFrom (Single name) = variable name -- Warp ops - warpAssignArray (Single name) warpID step a ix =- Assign name [warpID * fromIntegral step + ix] a + warpAssignArray (Single name) warpId step a ix =+ Assign name [warpId * fromIntegral step + ix] a - warpPullFrom (Single name) warpID n- = mkPull n (\i -> index name (warpID * fromIntegral n + i))+ warpPullFrom (Single name) warpId n+ = mkPull n (\i -> index name (warpId * fromIntegral n + i)) + -- Thread ops+ threadAssignArray (Single name) threadId step a ix =+ Assign name [threadId * fromIntegral step + ix] a++ threadPullFrom (Single name) threadId n+ = mkPull n (\i -> index name (threadId * fromIntegral n + i)) + -- Extra allocateVolatileArray (Single name) n = Allocate name (n * fromIntegral (sizeOf (undefined :: Exp a)))@@ -112,6 +128,11 @@ allocateScalar (Tuple ns1 ns2) = allocateScalar ns1 >> allocateScalar ns2++ allocateSharedScalar (Tuple ns1 ns2) =+ allocateSharedScalar ns1 >> + allocateSharedScalar ns2+ assignArray (Tuple ns1 ns2) (a,b) ix = assignArray ns1 a ix >>@@ -120,6 +141,10 @@ warpAssignArray (Tuple ns1 ns2) warpID step (a,b) ix = warpAssignArray ns1 warpID step a ix >> warpAssignArray ns2 warpID step b ix++ threadAssignArray (Tuple ns1 ns2) threadId step (a,b) ix =+ threadAssignArray ns1 threadId step a ix >> + threadAssignArray ns2 threadId step b ix assignScalar (Tuple ns1 ns2) (a,b) =@@ -134,8 +159,13 @@ warpPullFrom (Tuple ns1 ns2) warpID n = let p1 = warpPullFrom ns1 warpID n p2 = warpPullFrom ns2 warpID n- in mkPull n (\ix -> (p1 ! ix, p2 ! ix)) + in mkPull n (\ix -> (p1 ! ix, p2 ! ix)) + threadPullFrom (Tuple ns1 ns2) threadId n+ = let p1 = threadPullFrom ns1 threadId n+ p2 = threadPullFrom ns2 threadId n+ in mkPull n (\ix -> (p1 ! ix, p2 ! ix))+ readFrom (Tuple ns1 ns2) = let p1 = readFrom ns1 p2 = readFrom ns2@@ -166,7 +196,13 @@ allocateScalar ns1 >> allocateScalar ns2 >> allocateScalar ns3- ++ allocateSharedScalar (Triple ns1 ns2 ns3) =+ allocateSharedScalar ns1 >>+ allocateSharedScalar ns2 >>+ allocateSharedScalar ns3++ assignArray (Triple ns1 ns2 ns3) (a,b,c) ix = assignArray ns1 a ix >> assignArray ns2 b ix >>@@ -175,9 +211,14 @@ warpAssignArray (Triple ns1 ns2 ns3) warpID step (a,b,c) ix = warpAssignArray ns1 warpID step a ix >> warpAssignArray ns2 warpID step b ix >>- warpAssignArray ns3 warpID step c ix - + warpAssignArray ns3 warpID step c ix + threadAssignArray (Triple ns1 ns2 ns3) threadID step (a,b,c) ix =+ threadAssignArray ns1 threadID step a ix >>+ threadAssignArray ns2 threadID step b ix >>+ threadAssignArray ns3 threadID step c ix++ assignScalar (Triple ns1 ns2 ns3) (a,b,c) = assignScalar ns1 a >> assignScalar ns2 b >>@@ -195,9 +236,102 @@ p3 = warpPullFrom ns3 warpID n in mkPull n (\ix -> (p1 ! ix, p2 ! ix, p3 ! ix)) + threadPullFrom (Triple ns1 ns2 ns3) threadId n+ = let p1 = threadPullFrom ns1 threadId n+ p2 = threadPullFrom ns2 threadId n+ p3 = threadPullFrom ns3 threadId n + in mkPull n (\ix -> (p1 ! ix, p2 ! ix,p3 ! ix))++ readFrom (Triple ns1 ns2 ns3) = let p1 = readFrom ns1 p2 = readFrom ns2 p3 = readFrom ns3 in (p1,p2,p3) +instance (Storable a, Storable b, Storable c, Storable d) => Storable (a, b, c, d) where+ names pre =+ do+ (a :: Names a) <- names pre+ (b :: Names b) <- names pre+ (c :: Names c) <- names pre+ (d :: Names d) <- names pre+ return $ Quadruple a b c d++ allocateArray (Quadruple ns1 ns2 ns3 ns4) n =+ allocateArray ns1 n >>+ allocateArray ns2 n >>+ allocateArray ns3 n >>+ allocateArray ns4 n++ allocateVolatileArray (Quadruple ns1 ns2 ns3 ns4) n =+ allocateVolatileArray ns1 n >>+ allocateVolatileArray ns2 n >>+ allocateVolatileArray ns3 n >>+ allocateVolatileArray ns4 n++ allocateScalar (Quadruple ns1 ns2 ns3 ns4) =+ allocateScalar ns1 >>+ allocateScalar ns2 >>+ allocateScalar ns3 >>+ allocateScalar ns4++ allocateSharedScalar (Quadruple ns1 ns2 ns3 ns4) =+ allocateSharedScalar ns1 >>+ allocateSharedScalar ns2 >>+ allocateSharedScalar ns3 >>+ allocateSharedScalar ns4+++ assignArray (Quadruple ns1 ns2 ns3 ns4) (a,b,c,d) ix =+ assignArray ns1 a ix >>+ assignArray ns2 b ix >>+ assignArray ns3 c ix >>+ assignArray ns4 d ix++ warpAssignArray (Quadruple ns1 ns2 ns3 ns4) warpID step (a,b,c,d) ix =+ warpAssignArray ns1 warpID step a ix >>+ warpAssignArray ns2 warpID step b ix >>+ warpAssignArray ns3 warpID step c ix >>+ warpAssignArray ns4 warpID step d ix++ threadAssignArray (Quadruple ns1 ns2 ns3 ns4) threadID step (a,b,c,d) ix =+ threadAssignArray ns1 threadID step a ix >>+ threadAssignArray ns2 threadID step b ix >>+ threadAssignArray ns3 threadID step c ix >>+ threadAssignArray ns4 threadID step d ix+ ++ assignScalar (Quadruple ns1 ns2 ns3 ns4) (a,b,c,d) =+ assignScalar ns1 a >>+ assignScalar ns2 b >>+ assignScalar ns3 c >>+ assignScalar ns4 d++ pullFrom (Quadruple ns1 ns2 ns3 ns4) n =+ let p1 = pullFrom ns1 n+ p2 = pullFrom ns2 n+ p3 = pullFrom ns3 n+ p4 = pullFrom ns4 n+ in mkPull n (\ix -> (p1 ! ix, p2 ! ix, p3 ! ix, p4 ! ix))++ warpPullFrom (Quadruple ns1 ns2 ns3 ns4) warpID n+ = let p1 = warpPullFrom ns1 warpID n+ p2 = warpPullFrom ns2 warpID n+ p3 = warpPullFrom ns3 warpID n+ p4 = warpPullFrom ns4 warpID n+ in mkPull n (\ix -> (p1 ! ix, p2 ! ix, p3 ! ix, p4 ! ix))++ threadPullFrom (Quadruple ns1 ns2 ns3 ns4) threadID n+ = let p1 = threadPullFrom ns1 threadID n+ p2 = threadPullFrom ns2 threadID n+ p3 = threadPullFrom ns3 threadID n+ p4 = threadPullFrom ns4 threadID n+ in mkPull n (\ix -> (p1 ! ix, p2 ! ix, p3 ! ix, p4 ! ix))++ readFrom (Quadruple ns1 ns2 ns3 ns4) =+ let p1 = readFrom ns1+ p2 = readFrom ns2+ p3 = readFrom ns3+ p4 = readFrom ns4+ in (p1,p2,p3,p4)
Obsidian/Mutable.hs view
@@ -3,31 +3,30 @@ EmptyDataDecls, FlexibleInstances #-} -{- Joel Svensson 2013 -}+{- Joel Svensson 2013, 2014 -} module Obsidian.Mutable ( Mutable(Mutable) , Shared- , Global - , newS- , forceTo+ , Global+ , MShared+ , MGlobal+ , newSharedMutable+ , writeToSync , writeTo- , pullFrom+ , assignMutable+ , indexMutable + , mutableToPull , atomicInc--- , mutlen -- hack- , namedMutable- , undefinedMutable ) where import Obsidian.Exp-import Obsidian.Types-import Obsidian.Globs import Obsidian.Program import Obsidian.Memory import Obsidian.Names import Obsidian.Array-import Obsidian.Atomic +import Obsidian.Atomic import Data.Word @@ -65,17 +64,14 @@ len (Mutable n _) = n ---namedMutable s v = Mutable v (Single s)-undefinedMutable v = Mutable v undefined --------------------------------------------------------------------------- -- Create Mutable Shared memory arrays -- # allocates shared memory --------------------------------------------------------------------------- -newS :: Storable a => SPush Block a -> Program Block (Mutable Shared Word32 a)-newS arr = do+-- | Create a new Mutable array in shared memory +newSharedMutable :: Storable a => SPush Block a -> Program Block (Mutable Shared Word32 a)+newSharedMutable arr = do (snames :: Names a) <- names "arr" allocateArray snames n let mut = Mutable n snames@@ -87,7 +83,10 @@ --------------------------------------------------------------------------- -- forceTo & writeTo ------------------------------------------------------------------------------ Much Hacking here +-- Much Hacking here++-- | Write a Push array into a mutable array.+-- There is no synchronisation inserted after the write writeTo :: Storable a => Mutable Shared Word32 a -> Push Block Word32 a@@ -100,20 +99,35 @@ -- Add forceTo with offsets (why? just thought it might be useful)-forceTo :: Storable a- => Mutable Shared Word32 a- -> Push Block Word32 a- -> Program Block ()-forceTo m arr =+-- | Write a Push array into a mutable array and sync+writeToSync :: Storable a+ => Mutable Shared Word32 a+ -> Push Block Word32 a+ -> Program Block ()+writeToSync m arr = do writeTo m arr- Sync + Sync+ ------------------------------------------------------------------------------ pullFrom +-- Low level operations --------------------------------------------------------------------------- -toPull :: Storable a => Mutable Shared Word32 a -> SPull a-toPull (Mutable n snames) = pullFrom snames n +-- | Write a value into a storable array. +assignMutable :: Storable a => Mutable loc l a -> EWord32 -> a -> Program Thread ()+assignMutable (Mutable _ snames) ix a = assignArray snames a ix + +indexMutable :: (ASize s, Storable a) => Mutable loc s a -> EWord32 -> a +indexMutable mut ix = mutableToPull mut ! ix + +---------------------------------------------------------------------------+-- mutable to pull conversion+---------------------------------------------------------------------------++-- | Convert a Mutable array to a Pull array +mutableToPull :: (ASize s, Storable a) => Mutable l s a -> Pull s a+mutableToPull (Mutable n snames) = pullFrom snames n + ---------------------------------------------------------------------------
Obsidian/Names.hs view
@@ -9,6 +9,7 @@ Single :: Name -> Names a Tuple :: Names a -> Names b -> Names (a,b) Triple :: Names a -> Names b -> Names c -> Names (a,b,c) + Quadruple :: Names a -> Names b -> Names c -> Names d -> Names (a,b,c,d) --------------------------------------------------------------------------- -- helpers@@ -20,4 +21,9 @@ mapNamesM_ f (Triple n1 n2 n3) = mapNamesM_ f n1 >> mapNamesM_ f n2 >> mapNamesM_ f n3+mapNamesM_ f (Quadruple n1 n2 n3 n4) = mapNamesM_ f n1 >>+ mapNamesM_ f n2 >>+ mapNamesM_ f n3 >>+ mapNamesM_ f n4+
Obsidian/Program.hs view
@@ -1,6 +1,7 @@-{- Joel Svensson 2012,2013+{- Joel Svensson 2012..2015 Notes:+ 2014 : starting a big overhauling 2013-04-02: Added a Break statement to the language. Use it to break out of sequential loops. 2013-01-08: removed number-of-blocks field from ForAllBlocks@@ -8,51 +9,53 @@ -} {-# LANGUAGE GADTs #-}-{-# LANGUAGE TypeFamilies #-} {-# LANGUAGE EmptyDataDecls #-} {-# LANGUAGE FlexibleInstances #-} {-# LANGUAGE MultiParamTypeClasses #-} {-# LANGUAGE ScopedTypeVariables #-}--+{-# LANGUAGE TypeOperators #-} +{-# LANGUAGE FlexibleContexts #-}+{-# LANGUAGE UndecidableInstances #-} +{-# LANGUAGE ExplicitNamespaces #-} +{-# LANGUAGE TypeFamilies #-}+{-# LANGUAGE DataKinds #-}+{-# LANGUAGE ConstraintKinds #-}+------------------------------------------+{- LANGUAGE FunctionalDependencies -} module Obsidian.Program ( -- Hierarchy - Thread, Block, Grid, Step, Zero, Warp, - -- Program type- -- CoreProgram(..),- Program(..), -- all exported.. for now- TProgram, BProgram, GProgram, WProgram(..), + Thread, Block, Grid, Warp, Step, + Program(..), + TProgram, BProgram, GProgram, WProgram, + -- Class- Sync, + type (*<=*), -- helpers printPrg, runPrg, uniqueNamed, uniqueNamed_, - assign, allocate, declare,+ allocate, declare, atomicOp, + -- Programming interface- seqFor, forAll, forAll2, seqWhile, sync, distrPar,+ seqFor, forAll, seqWhile, sync, distrPar, forAll2,+ singleThread++ ) where import Data.Word-import Data.Monoid import Obsidian.Exp import Obsidian.Types import Obsidian.Globs import Obsidian.Atomic-import Obsidian.Names --- Package value-supply-import Data.Supply-import System.IO.Unsafe--import Control.Monad import Control.Applicative @@ -60,25 +63,31 @@ -- Thread/Block/Grid --------------------------------------------------------------------------- -+data Thread+data Step t --- A hierarchy! -data Step a -- A step in the hierarchy-data Zero - -type Thread = Zero-type Warp = Step Thread-type Block = Step Warp-type Grid = Step Block+-- I would like these to be "somehow" closed. +type Warp = Step Thread+type Block = Step Warp+type Grid = Step Block +-- | Type level less-than-or-equal test.+type family LessThanOrEqual a b where+ LessThanOrEqual Thread Thread = True+ LessThanOrEqual Thread (Step m) = True+ LessThanOrEqual (Step n) (Step m) = LessThanOrEqual n m+ LessThanOrEqual x y = False+ +-- | This constraint is a more succinct way of requiring that @a@ be less than or equal to @b@.+type a *<=* b = (LessThanOrEqual a b ~ True) --------------------------------------------------------------------------- type Identifier = Int- + --------------------------------------------------------------------------- -- Program datatype --------------------------------------------------------------------------@@ -110,48 +119,48 @@ Break :: Program Thread () - -- use threads along one level- -- Warp, Block, Grid. - ForAll :: EWord32 + -- Thread, Warp, Block.+ -- Make sure Code generation works when t ~ Thread+ ForAll :: (t *<=* Block) => EWord32 -> (EWord32 -> Program Thread ())- -> Program t () -- (really atleast Step t) ! + -> Program t () - -- Distribute over Warps yielding a Block -- Distribute over Blocks yielding a Grid DistrPar :: EWord32 -> (EWord32 -> Program t ()) -> Program (Step t) ()- -- (really Step t -> Step (Step t) ) - - SeqFor :: EWord32 -> (EWord32 -> Program t ())- -> Program t () - - - -- #w warpId - --NWarps :: EWord32 -> (EWord32 -> Program Warp ()) -> Program Block () + -- BUG: I Need to recognize sequential distribution of+ -- work too in order to set up storage correctly for+ -- arrays allocated in within sequentially distributed work.+ -- This has been "fixed" (a bit of a hack) in CodeGen/Memory2.hs+ - --WarpForAll :: EWord32 - -- -> (EWord32 -> Program Thread ()) - -- -> Program Warp ()- -- WarpAllocate :: Name -> Word32 -> Type -> Program Warp () -- For now. + -- Sequential for loop+ SeqFor :: EWord32 -> (EWord32 -> Program t ())+ -> Program t () -- Allocate shared memory in each MP+ -- Can be done from any program level.+ -- Since the allocation happens block-wise though+ -- it is important to figure out how many instances of+ -- that t level program that needs memory! (messy) Allocate :: Name -> Word32 -> Type -> Program t () -- Automatic Variables Declare :: Name -> Type -> Program t () - Sync :: Program Block ()+ Sync :: (t *<=* Block) => Program t () - -- Parallel composition of Programs- -- TODO: Will I use this ? - --Par :: Program p () ->- -- Program p () ->- -- Program p () + -- WarpShuffle :: ... -> Program Warp () + -- Think about how to to allow multikernel programs.++ -- HardSync :: Program Grid () + -- How to decide arguments to the different kernels ? + -- Monad Return :: a -> Program t a Bind :: Program t a -> (a -> Program t b) -> Program t b@@ -164,22 +173,6 @@ type BProgram = Program Block type GProgram = Program Grid --- -- Programs are a reader monad --- newtype Program t a = Program (EWord32 -> CoreProgram t a)---- instance Monad (Program t )where--- return x = Program $ \ _ -> return x--- -- :: WProgram a -> (a -> WProgram b) -> WProgram b--- (Program h) >>= f = Program--- $ \w ->--- do--- a <- h w--- let (Program g) = f a--- g w----- core :: Program t a -> EWord32 -> CoreProgram t a--- core (Program f) id = f id --------------------------------------------------------------------------- -- Helpers --------------------------------------------------------------------------- @@ -196,7 +189,7 @@ --------------------------------------------------------------------------- -- Memory ----------------------------------------------------------------------------assign :: Scalar a => Name -> [Exp Word32] -> (Exp a) -> Program Thread ()+assign :: Scalar a => Name -> [Exp Word32] -> Exp a -> Program Thread () assign nom ix e = Assign nom ix e allocate :: Name -> Word32 -> Type -> Program t () @@ -218,23 +211,18 @@ --------------------------------------------------------------------------- -- forAll ----------------------------------------------------------------------------forAll :: EWord32 -> (EWord32 -> Program Thread ()) -> Program t ()-forAll n f = ForAll n $ \ix -> f ix+forAll :: (t *<=* Block) => EWord32+ -> (EWord32 -> Program Thread ())+ -> Program t ()+forAll n f = ForAll n f --- forAll :: EWord32 -> (EWord32 -> Program t ()) -> Program (Step t) ()--- forAll n f = Program $ \id -> ForAll n $ \ix -> core (f ix) id--forAll2 :: EWord32+forAll2 :: (t *<=* Block) => EWord32 -> EWord32 -> (EWord32 -> EWord32 -> Program Thread ())- -> Program (Step (Step t)) ()-forAll2 b n f = DistrPar b $ \bs -> ForAll n $ \ix -> f bs ix ---- forAll2 :: EWord32--- -> EWord32--- -> (EWord32 -> EWord32 -> Program t ())--- -> Program (Step (Step t)) ()--- forAll2 b n f = forAll b $ \bs -> forAll n (f bs)+ -> Program (Step t) ()+forAll2 b n f =+ DistrPar b $ \bs ->+ ForAll n $ \ix -> f bs ix distrPar :: EWord32 -> (EWord32 -> Program t ())@@ -242,17 +230,18 @@ distrPar b f = DistrPar b $ \bs -> f bs ------------------------------------------------------------------------------ warpForAll -------------------------------------------------------------------------------warpForAll :: EWord32 -> (EWord32 -> Program Thread ()) -> Program Warp ()--- warpForAll n f = Program $ \id -> WarpForAll n $ \ix -> core (f ix) id - +-- Let a single thread of a block/Warp perform a given+-- Thread program ---------------------------------------------------------------------------+singleThread :: (t *<=* Block) => Program Thread () -> Program t ()+singleThread p =+ forAll 1 (\_ -> p) + -- seqFor --------------------------------------------------------------------------- seqFor :: EWord32 -> (EWord32 -> Program t ()) -> Program t () seqFor (Literal 1) f = f 0-seqFor n f = SeqFor n $ \ix -> f ix+seqFor n f = SeqFor n f --------------------------------------------------------------------------- -- seqWhile@@ -284,23 +273,10 @@ fmap f fa ------------------------------------------------------------------------------ Class Sync+-- sync function ----------------------------------------------------------------------------class Sync t where- sync :: Program t () --instance Sync Warp where- sync = return ()--instance Sync Thread where- sync = return ()--instance Sync Block where- sync = Sync--instance Sync Grid where- sync = error "sync: not implemented on grid computations" - -- (implement this using counters and locks)+sync :: (t *<=* Block) => Program t ()+sync = Sync --------------------------------------------------------------------------- -- runPrg (RETHINK!) (Works for Block programs, but all?)@@ -314,7 +290,9 @@ runPrg i (Bind m f) = let (a,i') = runPrg i m in runPrg i' (f a)- ++-- All other constructors have () result + runPrg i (Sync) = ((),i) runPrg i (ForAll n ixf) = let (p,i') = runPrg i (ixf (variable "tid")) @@ -323,11 +301,11 @@ let (p,i') = runPrg i (f (variable "DUMMY")) in (p,i') -- What can this boolean depend upon ? its quite general!--- (we know p returns a ()... ) +-- p here is a Program Thread () runPrg i (Cond b p) = ((),i) runPrg i (Declare _ _) = ((),i) runPrg i (Allocate _ _ _ ) = ((),i)-runPrg i (Assign _ _ a) = ((),i) -- Probaby wrong.. +runPrg i (Assign _ _ a) = ((),i) runPrg i (AtomicOp _ _ _) = ((),i) -- variable ("new"++show i),i+1) {- What do I want from runPrg ?
Obsidian/Run/CUDA/Exec.hs view
@@ -2,10 +2,54 @@ ScopedTypeVariables, TypeFamilies, TypeSynonymInstances,- FlexibleInstances #-} + FlexibleInstances #-}+{-# LANGUAGE BangPatterns #-} -module Obsidian.Run.CUDA.Exec where+{-+ Gah!+ Most of this needs to be improved. + # Add functionality to pass arrays of tuples in and out+ from kernels. +-} +++module Obsidian.Run.CUDA.Exec ( mkRandomVec+ , CUDAVector+ , getDevices+ , propsSummary+ , sharedMemConfig+ , CUDA+ , KernelT+ , (<:>)+ , (<>)+ , (<==)+ , syncAll+ , exec+ , withCUDA+ , withCUDA'+ , destroyCtx+ , Context(..)+ , initialise+ , capture+ , captureIO+ , useVector+ , allocaVector+ , allocaFillVector+ , withVector+ , withFillVector + , allocaVector_+ , mallocVector+ , mallocVectorIO+ , freeVector+ , freeVectorIO+ , fill+ , peekCUDAVector+ , copyOut+ , copyInIO+ , copyOutIO+ ) where+ --------------------------------------------------------------------------- -- -- Low level interface to CUDA functionality from Obsidian@@ -14,58 +58,98 @@ import qualified Foreign.CUDA.Driver as CUDA-import qualified Foreign.CUDA.Driver.Device as CUDA-import qualified Foreign.CUDA.Analysis.Device as CUDA-import qualified Foreign.CUDA.Driver.Stream as CUDAStream import Obsidian.CodeGen.Reify import Obsidian.CodeGen.CUDA -import Obsidian.Types -- experimental-import Obsidian.Exp+-- import Obsidian.Types -- experimental+import Obsidian.Exp hiding (sizeOf) import Obsidian.Array import Obsidian.Program (Grid, GProgram) import Obsidian.Mutable -import Foreign.Marshal.Array+-- import Foreign.Marshal.Array import Foreign.ForeignPtr.Unsafe -- (req GHC 7.6 ?)-import Foreign.ForeignPtr hiding (unsafeForeignPtrToPtr)+import Foreign.ForeignPtr -import qualified Data.Vector.Storable as V +import qualified Data.Vector.Storable as V+import Foreign.Storable+import Foreign.Ptr import Data.Word import Data.Int-import Data.Supply-import Data.List-import qualified Data.Map as M-import Data.Maybe--import System.IO.Unsafe import System.Process import System.Random.MWC-import System.CPUTime.Rdtsc import Control.Monad.State +debug :: Bool+debug = False +---------------------------------------------------------------------------+--+---------------------------------------------------------------------------+instance V.Storable a => V.Storable (Vector4 a) where+ sizeOf _ = sizeOf (undefined :: a) * 4+ alignment _ = alignment (undefined :: a)+ + {-# INLINE peek #-}+ peek p = do+ a <- peekElemOff q 0+ b <- peekElemOff q 1+ c <- peekElemOff q 2+ d <- peekElemOff q 3+ return (Vector4 a b c d)+ where+ q = castPtr p+ {-# INLINE poke #-}+ poke p (Vector4 a b c d) = do+ pokeElemOff q 0 a+ pokeElemOff q 1 b+ pokeElemOff q 2 c+ pokeElemOff q 3 d+ where+ q = castPtr p -debug = False+instance V.Storable a => V.Storable (Vector2 a) where+ sizeOf _ = sizeOf (undefined :: a) * 2+ alignment _ = alignment (undefined :: a)+ + {-# INLINE peek #-}+ peek p = do+ a <- peekElemOff q 0+ b <- peekElemOff q 1+ return (Vector2 a b )+ where+ q = castPtr p+ {-# INLINE poke #-}+ poke p (Vector2 a b) = do+ pokeElemOff q 0 a+ pokeElemOff q 1 b+ where+ q = castPtr p + ++ --------------------------------------------------------------------------- -- Tools ---------------------------------------------------------------------------+-- | Generate a random vector mkRandomVec :: forall a.(V.Storable a, Variate a) => Int -> IO (V.Vector a) mkRandomVec k = withSystemRandom $ \g -> uniformVector g k :: IO (V.Vector a) --------------------------------------------------------------------------- -- An array located in GPU memory ----------------------------------------------------------------------------data CUDAVector a = CUDAVector {cvPtr :: CUDA.DevicePtr a,- cvLen :: Word32} +-- | Represents vectors (arrays) in GPU memory.+data CUDAVector a = CUDAVector {cvPtr :: !(CUDA.DevicePtr a),+ cvLen :: !Word32} --------------------------------------------------------------------------- -- Get a list of devices from the CUDA driver ---------------------------------------------------------------------------+-- | Get a list of CUDA devices available in the system getDevices :: IO [(CUDA.Device,CUDA.DeviceProperties)] getDevices = do num <- CUDA.count @@ -76,6 +160,7 @@ --------------------------------------------------------------------------- -- Print a Summary of a device's properties. ---------------------------------------------------------------------------+-- | A string showing some of the properties os a CUDA device. propsSummary :: CUDA.DeviceProperties -> String propsSummary props = unlines ["Device Name: " ++ CUDA.deviceName props,@@ -90,29 +175,44 @@ "Num MP: " ++ show (CUDA.multiProcessorCount props), "Mem bus width: " ++ show (CUDA.memBusWidth props)] +---------------------------------------------------------------------------+-- Extract shared mem config+---------------------------------------------------------------------------+-- | Extract shared memory configuration. This information is used+-- by the memory layout manager.+sharedMemConfig :: CUDA.DeviceProperties -> SharedMemConfig +sharedMemConfig props = SharedMemConfig sm_bytes num_banks True+ where+ sm_bytes = fromIntegral $ CUDA.sharedMemPerBlock props+ num_banks = if x < 2+ then 16+ else 32 + (CUDA.Compute x _) = CUDA.computeCapability props+ -------------------------------------------------------------------------- -- Environment to run CUDA computations in. -- # Needs to keep track of generated and loaded functions etc. ----------------------------------------------------------------------------+-- | An environment to perform CUDA computations within. data CUDAState = CUDAState { csIdent :: Int,+ csCtx :: CUDA.Context, csProps :: CUDA.DeviceProperties} -type CUDA a = StateT CUDAState IO a+type CUDA a = StateT CUDAState IO a -data Kernel = Kernel {kFun :: CUDA.Fun,- kThreadsPerBlock :: Word32,- kSharedBytes :: Word32} -- Change so that the type parameter to KernelT--- represents the "captured" type, with CUDAVectors instead of Pull, Push vectors. -data KernelT a = KernelT {ktFun :: CUDA.Fun,- ktThreadsPerBlock :: Word32,- ktSharedBytes :: Word32,- ktInputs :: [CUDA.FunParam],- ktOutput :: [CUDA.FunParam] }+-- represents the "captured" type, with CUDAVectors instead of Pull, Push vectors. +-- | A representation of a CUDA kernel. This represents the+-- kernel after it has been loaded from an Object file. +data KernelT a = KernelT {ktFun :: !CUDA.Fun,+ ktThreadsPerBlock :: !Word32,+ ktSharedBytes :: !Word32,+ ktInputs :: ![CUDA.FunParam],+ ktOutput :: ![CUDA.FunParam] }+ --------------------------------------------------------------------------- -- Kernel Input and Output classes ---------------------------------------------------------------------------@@ -131,7 +231,7 @@ instance KernelI Int32 where type KInput Int32 = Exp Int32 addInParam (KernelT f t s i o) a =- KernelT f t s (i ++ [CUDA.IArg $ fromIntegral a]) o + KernelT f t s (i ++ [CUDA.IArg $ fromIntegral a]) o instance KernelI Word32 where type KInput Word32 = Exp Word32@@ -158,13 +258,15 @@ --------------------------------------------------------------------------- -- (<>) apply a kernel to an input ---------------------------------------------------------------------------+-- | Apply a kernel to an input (<>) :: KernelI a => (Word32,KernelT (KInput a -> b)) -> a -> (Word32,KernelT b) (<>) (blocks,kern) a = (blocks,addInParam kern a) --------------------------------------------------------------------------- -- Assign a mutable input/output to a kernel---------------------------------------------------------------------------- +---------------------------------------------------------------------------+-- | apply a kernel to a mutable array (<:>) :: KernelM a => (Word32, KernelT (KMutable a -> b)) -> a -> (Word32, KernelT b) (<:>) (blocks,kern) a = (blocks, addMutable kern a)@@ -173,6 +275,8 @@ --------------------------------------------------------------------------- -- Execute a kernel and store output to an array ---------------------------------------------------------------------------+-- | Infix operator taking an allocated output array on LHS and+-- a (Number_Of_Blocks,Kernel) pair on the right. (<==) :: KernelO b => b -> (Word32, KernelT (KOutput b)) -> CUDA () (<==) o (nb,kern) = do@@ -202,7 +306,9 @@ -- (ktInputs k ++ ktOutput k) -- params -- lift $ CUDA.sync -- t2 <- lift rdtsc--- return (t2 - t1) +-- return (t2 - t1)++-- | Wait for all launched work on the device to finish syncAll :: CUDA () syncAll = lift $ CUDA.sync @@ -214,6 +320,8 @@ --------------------------------------------------------------------------- -- Execute a kernel that has no Output ( a -> GProgram ()) KernelT (GProgram ()) ---------------------------------------------------------------------------+-- | Execute a kernel that has no specified output.+-- It may write data into a Global Mutable array. exec :: (Word32, KernelT (GProgram ())) -> CUDA () exec (nb, k) = lift $ CUDA.launchKernel@@ -233,19 +341,43 @@ modify (\s -> s {csIdent = i+1 }) return i +data Context = Context {context :: !CUDA.Context,+ props :: !CUDA.DeviceProperties} +initialise :: IO (Context) +initialise =+ do CUDA.initialise []+ devs <- getDevices+ + case devs of+ [] -> error "No CUDA device found!"+ (x:_) ->+ do ctx <- CUDA.create (fst x) [CUDA.SchedAuto]+ return (Context ctx (snd x)) ++destroyCtx :: Context -> IO ()+destroyCtx (Context ctx props) = CUDA.destroy ctx++ +withCUDA' :: Context -> CUDA a -> IO a+withCUDA' (Context ctx props) p =+ do+ (a,_) <- runStateT p (CUDAState 0 ctx props)+ return a+ --------------------------------------------------------------------------- -- Run a CUDA computation ---------------------------------------------------------------------------+withCUDA :: CUDA a -> IO a withCUDA p = do CUDA.initialise [] devs <- getDevices case devs of [] -> error "No CUDA device found!" - (x:xs) ->+ (x:_) -> do - ctx <- CUDA.create (fst x) [CUDA.SchedAuto] + !ctx <- CUDA.create (fst x) [CUDA.SchedAuto] (a,_) <- runStateT p (CUDAState 0 ctx (snd x)) CUDA.destroy ctx return a@@ -267,27 +399,66 @@ let kn = "gen" ++ show i fn = kn ++ ".cu"- cub = fn ++ ".cubin"+ --cub = fn ++ ".cubin" - prgstr = genKernel threadsPerBlock kn f+ sm_conf = sharedMemConfig props+ + prgstr = genKernelParams sm_conf threadsPerBlock kn f header = "#include <stdint.h>\n" -- more includes ? -+ + when debug $ do lift $ putStrLn $ prgstr let arch = archStr props - lift $ storeAndCompile arch (fn) (header ++ prgstr)+ cub_file <- lift $ storeAndCompile arch fn (header ++ prgstr) - mod <- liftIO $ CUDA.loadFile cub- fun <- liftIO $ CUDA.getFun mod kn + cubin <- liftIO $ CUDA.loadFile cub_file+ fun <- liftIO $ CUDA.getFun cubin kn {- After loading the binary into the running process can I delete the .cu and the .cu.cubin ? -} return $! KernelT fun threadsPerBlock 0 {-bytesShared-} [] [] ++-- | Compile a program to a CUDA kernel and then load it into memory.+captureIO :: ToProgram prg+ => String -- ^ name of Kernel (and files stored to disk)+ -> CUDA.DeviceProperties -- ^ deviceproperties to compile for + -> Word32 -- ^ Threads per block+ -> prg -- ^ Program to capture+ -> IO (KernelT prg) +captureIO nom props threadsPerBlock f =+ do+ let kn = nom+ fn = kn ++ ".cu"++ sm_conf = sharedMemConfig props+ + prgstr = genKernelParams sm_conf threadsPerBlock kn f+ header = "#include <stdint.h>\n" -- more includes ? + + + when debug $ + do + putStrLn $ prgstr++ let arch = archStr props+ + cub_file <- storeAndCompile arch fn (header ++ prgstr)+ + cubin <- liftIO $ CUDA.loadFile cub_file+ fun <- liftIO $ CUDA.getFun cubin kn ++ {- After loading the binary into the running process+ can I delete the .cu and the .cu.cubin ? -} + + return $! KernelT fun threadsPerBlock 0 {-bytesShared-} [] []++ --------------------------------------------------------------------------- -- useVector: Copies a Data.Vector from "Haskell" onto the GPU Global mem --------------------------------------------------------------------------- @@ -305,9 +476,12 @@ lift $ CUDA.free dptr return b +-- THESE SHOULD BE RENAMED -- --------------------------------------------------------------------------- -- allocaVector: allocates room for a vector in the GPU Global mem ---------------------------------------------------------------------------+{-# DEPRECATED allocaVector "Use withVector instead" #-} +-- | allocate a vector in GPU Device memory allocaVector :: V.Storable a => Int -> (CUDAVector a -> CUDA b) -> CUDA b allocaVector n f =@@ -316,13 +490,72 @@ let cvector = CUDAVector dptr (fromIntegral n) b <- f cvector -- dptr lift $ CUDA.free dptr- return b + return b +withVector :: V.Storable a =>+ Int -> (CUDAVector a -> CUDA b) -> CUDA b+withVector n f = allocaVector n f + + ---------------------------------------------------------------------------+-- Low level memory allocation/deallocation+---------------------------------------------------------------------------+{-# DEPRECATED allocaVector_ "Use mallocVector instead" #-} +-- | allocate a vector in GPU Device memory +allocaVector_ :: V.Storable a =>+ Int -> CUDA (CUDAVector a) +allocaVector_ n =+ do+ dptr <- lift $ CUDA.mallocArray n+ return $ CUDAVector dptr (fromIntegral n)++mallocVector :: V.Storable a =>+ Int -> CUDA (CUDAVector a)+mallocVector = allocaVector_++mallocVectorIO :: V.Storable a =>+ Int -> IO (CUDAVector a)+mallocVectorIO n =+ do+ dptr <- CUDA.mallocArray n+ return $ CUDAVector dptr (fromIntegral n) ++freeVector :: CUDAVector a -> CUDA ()+freeVector (CUDAVector dptr _) = lift $ CUDA.free dptr++freeVectorIO :: CUDAVector a -> IO ()+freeVectorIO (CUDAVector dptr _) = CUDA.free dptr ++copyInIO :: V.Storable a =>+ V.Vector a -> IO (CUDAVector a) +copyInIO v =+ do+ let (hfptr,n) = V.unsafeToForeignPtr0 v+ + dptr <- CUDA.mallocArray n+ let hptr = unsafeForeignPtrToPtr hfptr+ CUDA.pokeArray n hptr dptr+ let cvector = CUDAVector dptr (fromIntegral (V.length v)) + -- b <- f cvector -- dptr+ -- lift $ CUDA.free dptr+ return cvector++copyOutIO :: V.Storable a => CUDAVector a -> IO (V.Vector a)+copyOutIO (CUDAVector dptr n) =+ do+ (fptr :: ForeignPtr a) <- mallocForeignPtrArray (fromIntegral n)+ let ptr = unsafeForeignPtrToPtr fptr+ CUDA.peekArray (fromIntegral n) dptr ptr+ return $ V.unsafeFromForeignPtr fptr 0 (fromIntegral n)++++--------------------------------------------------------------------------- -- Allocate and fill with default value ---------------------------------------------------------------------------+-- | Allocate and Fill a vector in GPU Device memory allocaFillVector :: V.Storable a => - Int -> a -> (CUDAVector a -> CUDA b) -> CUDA b + Int -> a -> (CUDAVector a -> CUDA b) -> CUDA b allocaFillVector n a f = do dptr <- lift $ CUDA.mallocArray n@@ -332,9 +565,15 @@ lift $ CUDA.free dptr return b +withFillVector :: V.Storable a => + Int -> a -> (CUDAVector a -> CUDA b) -> CUDA b+withFillVector = allocaFillVector++ --------------------------------------------------------------------------- -- Fill a Vector ---------------------------------------------------------------------------+-- | Fill a vector fill :: V.Storable a => CUDAVector a -> a -> CUDA ()@@ -344,10 +583,13 @@ --------------------------------------------------------------------------- -- Peek in a CUDAVector (Simple "copy back") ---------------------------------------------------------------------------+-- | Copy a vector from the device to the host+-- as a list. (SLOW!) peekCUDAVector :: V.Storable a => CUDAVector a -> CUDA [a] peekCUDAVector (CUDAVector dptr n) = lift $ CUDA.peekListArray (fromIntegral n) dptr- ++-- | Copy a vector from the device to the host. (FAST!) copyOut :: V.Storable a => CUDAVector a -> CUDA (V.Vector a) copyOut (CUDAVector dptr n) = do@@ -363,13 +605,7 @@ archStr :: CUDA.DeviceProperties -> String archStr props = "-arch=sm_" ++ archStr' (CUDA.computeCapability props) where- -- Updated for Cuda bindings version 0.5.0.0 archStr' (CUDA.Compute h l) = show h ++ show l- --archStr' (CUDA.Compute 1 0) = "10"- --archStr' (CUDA.Compute 1 2) = "12"- --archStr' (CUDA.Compute 2 0) = "20" - --archStr' (CUDA.Compute 3 0) = "30"- --archStr' x = error $ "Unknown architecture: " ++ show x --------------------------------------------------------------------------- -- Compile to Cubin (interface with nvcc)@@ -383,5 +619,6 @@ (_,_,_,pid) <- createProcess (shell ("nvcc " ++ arch ++ " -cubin -o " ++ nfp ++ " " ++ fp))- exitCode <- waitForProcess pid+ -- This could fail. Should be error check here (waitForProcess).+ _ <- waitForProcess pid return nfp
Obsidian/SeqLoop.hs view
@@ -1,4 +1,9 @@-{-# LANGUAGE ScopedTypeVariables #-} +{-# LANGUAGE ScopedTypeVariables #-}+{-# LANGUAGE TypeOperators #-}+{-# LANGUAGE FlexibleContexts #-}+{-# LANGUAGE ConstraintKinds #-}+{-# LANGUAGE GADTs #-}+ {- sequential loops with state@@ -14,11 +19,12 @@ import Obsidian.Array import Obsidian.Memory import Obsidian.Names--import Data.Word+import Obsidian.Data+import Obsidian.Force+import qualified Obsidian.Library as Lib -- TODO: Rename module to something better-+-- Or make part of Library.hs --------------------------------------------------------------------------- -- seqReduce ---------------------------------------------------------------------------@@ -69,8 +75,9 @@ -- --------------------------------------------------------------------------- -- | iterate a function until a condition holds. Results in a while loop--- with a break in the generated code. +-- with a break in the generated code. seqUntil :: Storable a+-- The comment above seems to not desribe reality ! => (a -> a) -> (a -> EBool) -> a@@ -111,16 +118,16 @@ n = len arr -- | Sequential scan that takes a carry-in. -seqScanCin :: Storable a- => (a -> a -> a)- -> a -- cin +seqScanCin :: (Storable a, Storable b) + => (b -> a -> b)+ -> b -- cin -> SPull a- -> SPush Thread a-seqScanCin op a arr {-(Pull n ixf)-} =+ -> SPush Thread b+seqScanCin op a arr = mkPush n $ \wf -> do- (ns :: Names a) <- names "v" -- (ixf 0) - allocateScalar ns -- (ixf 0)- assignScalar ns a -- (ixf 0)+ (ns :: Names a) <- names "v" + allocateScalar ns + assignScalar ns a -- wf (readFrom ns) 0 SeqFor (sizeConv n) $ \ix -> do assignScalar ns $ readFrom ns `op` (arr ! ix)@@ -128,3 +135,65 @@ where n = len arr +-- | Sequential scan with separate types for input, output and accumulator.+mapAccumL :: (ASize s, Storable a, Storable b, Storable acc)+ => (acc -> a -> (acc,b))+ -> acc -- cin+ -> Pull s a+ -> Push Thread s b+mapAccumL op acc arr {-(Pull n ixf)-} =+ mkPush n $ \wf -> do+ (ns :: Names a) <- names "v" -- (ixf 0)+ allocateScalar ns -- (ixf 0)+ assignScalar ns acc -- (ixf 0)+ -- wf (readFrom ns) 0+ SeqFor (sizeConv n) $ \ix -> do+ let (newAcc, b) = op (readFrom ns) (arr ! ix)+ -- order of writing matters, because readFrom is evaluated twice+ wf b ix+ assignScalar ns newAcc+ where+ n = len arr++mapAccumR :: (ASize s, Storable a, Storable b, Storable acc)+ => (acc -> a -> (acc,b))+ -> acc -- cin+ -> Pull s a+ -> Push Thread s b+mapAccumR op acc =+ Lib.reverse . mapAccumL op acc . Lib.reverse+++--------------------------------------------------------------------------- +-- sMapAccum+-- Generalisation of the old sConcat functionality.++sMapAccum :: (Compute t, Data acc, ASize l)+ => (acc -> Pull l a -> Program t (acc,Push t l b))+ -> acc+ -> Pull l (Pull l a)+ -> Push t l b+sMapAccum f acc arr =+ + mkPush (n * fromIntegral rn) $ \wf ->+ do+ (noms :: Names acc) <- names "v"+ --(noms2 :: Names acc) <- names "APA"+ + allocateSharedScalar noms+ -- allocateScalar noms2+ -- a single thread in the group, performs an assignment+ -- May need synchronization! + singleThread $ assignScalar noms acc+ sync+ seqFor (sizeConv n) $ \bix -> do+ --singleThread $ assignScalar noms2 acc + (newAcc, b) <- f (readFrom noms) (arr ! bix)+ singleThread $ assignScalar noms newAcc+ sync+ let wf' a ix = wf a (bix * sizeConv rn + ix) + b <: wf'+ + where + n = len arr+ rn = len $ arr ! 0
Obsidian/Types.hs view
@@ -1,7 +1,7 @@ {-| Module : Types Description : Type information, used internally by Obsídian.-Copyright : (c) Joel Svensson, 2014+Copyright : (c) Joel Svensson, 2014, 2015 License : BSD Maintainer : bo.joel.svensson@gmail.com Stability : experimental@@ -19,17 +19,37 @@ = Bool | Int | Word -- A bit problematic since the size of -- of these are platform dependent+ +-- Vector types supported by CUDA (Add more) + | FloatV2 | FloatV3 | FloatV4+ | DoubleV2+ + -- | Int8V2 | Int8V3 | Int8V4+ -- | Int16V2 | Int16V3 | Int16V4+ -- | Int32V2 | Int32V3 | Int32V4+ -- | Int64V2+ + -- | Word8V2 | Word8V3 | Word8V4+ -- | Word16V2 | Word16V3 | Word16V4+ -- | Word32V2 | Word32V3 | Word32V4+ -- | Word64V2+ | Int8 | Int16 | Int32 | Int64 | Word8 | Word16 | Word32 | Word64 - | Float | Double + | Float | Double+-- Vector Types+ | Vec2 Type | Vec3 Type | Vec4 Type + -- Used by CUDA, C And OpenCL generators+ | Shared Type | Volatile Type -- For warp local computations. | Pointer Type -- Pointer to a @type@ | Global Type -- OpenCL thing | Local Type -- OpenCL thing deriving (Eq, Ord, Show)- ++typeSize :: Num a => Type -> a typeSize Int8 = 1 typeSize Int16 = 2 typeSize Int32 = 4@@ -40,4 +60,6 @@ typeSize Word64 = 8 typeSize Bool = 4 typeSize Float = 4-typeSize Double = 8 +typeSize Double = 8+typeSize (Shared t) = typeSize t +typeSize t = error $ "typeSize: this is bad!: " ++ show t