diff --git a/Obsidian.cabal b/Obsidian.cabal --- a/Obsidian.cabal +++ b/Obsidian.cabal @@ -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 Author: Joel Svensson -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 diff --git a/Obsidian.hs b/Obsidian.hs --- a/Obsidian.hs +++ b/Obsidian.hs @@ -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 diff --git a/Obsidian/Array.hs b/Obsidian/Array.hs --- a/Obsidian/Array.hs +++ b/Obsidian/Array.hs @@ -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 +--------------------------------------------------------------------------- +-- +--------------------------------------------------------------------------- diff --git a/Obsidian/Atomic.hs b/Obsidian/Atomic.hs --- a/Obsidian/Atomic.hs +++ b/Obsidian/Atomic.hs @@ -1,3 +1,4 @@ + {-# LANGUAGE GADTs #-} {- Joel Svensson, diff --git a/Obsidian/CodeGen/CUDA.hs b/Obsidian/CodeGen/CUDA.hs --- a/Obsidian/CodeGen/CUDA.hs +++ b/Obsidian/CodeGen/CUDA.hs @@ -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 + diff --git a/Obsidian/CodeGen/CompileIM.hs b/Obsidian/CodeGen/CompileIM.hs --- a/Obsidian/CodeGen/CompileIM.hs +++ b/Obsidian/CodeGen/CompileIM.hs @@ -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);|]] diff --git a/Obsidian/CodeGen/Liveness.hs b/Obsidian/CodeGen/Liveness.hs --- a/Obsidian/CodeGen/Liveness.hs +++ b/Obsidian/CodeGen/Liveness.hs @@ -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 diff --git a/Obsidian/CodeGen/Memory.hs b/Obsidian/CodeGen/Memory.hs deleted file mode 100644 --- a/Obsidian/CodeGen/Memory.hs +++ /dev/null @@ -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) - - diff --git a/Obsidian/CodeGen/Memory2.hs b/Obsidian/CodeGen/Memory2.hs new file mode 100644 --- /dev/null +++ b/Obsidian/CodeGen/Memory2.hs @@ -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) + + diff --git a/Obsidian/CodeGen/Program.hs b/Obsidian/CodeGen/Program.hs --- a/Obsidian/CodeGen/Program.hs +++ b/Obsidian/CodeGen/Program.hs @@ -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" diff --git a/Obsidian/CodeGen/Reify.hs b/Obsidian/CodeGen/Reify.hs --- a/Obsidian/CodeGen/Reify.hs +++ b/Obsidian/CodeGen/Reify.hs @@ -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 diff --git a/Obsidian/Data.hs b/Obsidian/Data.hs new file mode 100644 --- /dev/null +++ b/Obsidian/Data.hs @@ -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 diff --git a/Obsidian/Exp.hs b/Obsidian/Exp.hs --- a/Obsidian/Exp.hs +++ b/Obsidian/Exp.hs @@ -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)) diff --git a/Obsidian/Force.hs b/Obsidian/Force.hs --- a/Obsidian/Force.hs +++ b/Obsidian/Force.hs @@ -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 diff --git a/Obsidian/Library.hs b/Obsidian/Library.hs --- a/Obsidian/Library.hs +++ b/Obsidian/Library.hs @@ -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 - diff --git a/Obsidian/Memory.hs b/Obsidian/Memory.hs --- a/Obsidian/Memory.hs +++ b/Obsidian/Memory.hs @@ -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) diff --git a/Obsidian/Mutable.hs b/Obsidian/Mutable.hs --- a/Obsidian/Mutable.hs +++ b/Obsidian/Mutable.hs @@ -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 + --------------------------------------------------------------------------- diff --git a/Obsidian/Names.hs b/Obsidian/Names.hs --- a/Obsidian/Names.hs +++ b/Obsidian/Names.hs @@ -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 + diff --git a/Obsidian/Program.hs b/Obsidian/Program.hs --- a/Obsidian/Program.hs +++ b/Obsidian/Program.hs @@ -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 ? diff --git a/Obsidian/Run/CUDA/Exec.hs b/Obsidian/Run/CUDA/Exec.hs --- a/Obsidian/Run/CUDA/Exec.hs +++ b/Obsidian/Run/CUDA/Exec.hs @@ -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 \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 \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 diff --git a/Obsidian/SeqLoop.hs b/Obsidian/SeqLoop.hs --- a/Obsidian/SeqLoop.hs +++ b/Obsidian/SeqLoop.hs @@ -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 diff --git a/Obsidian/Types.hs b/Obsidian/Types.hs --- a/Obsidian/Types.hs +++ b/Obsidian/Types.hs @@ -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