Obsidian (empty) → 0.0.0.1
raw patch · 25 files changed
+4279/−0 lines, 25 filesdep +basedep +containersdep +mtlsetup-changed
Dependencies added: base, containers, mtl, value-supply
Files
- LICENSE +28/−0
- Obsidian.cabal +79/−0
- Obsidian.hs +29/−0
- Obsidian/Array.hs +179/−0
- Obsidian/Atomic.hs +21/−0
- Obsidian/CodeGen/CUDA.hs +209/−0
- Obsidian/CodeGen/Common.hs +176/−0
- Obsidian/CodeGen/InOut.hs +155/−0
- Obsidian/CodeGen/Liveness.hs +144/−0
- Obsidian/CodeGen/Memory.hs +152/−0
- Obsidian/CodeGen/PP.hs +55/−0
- Obsidian/CodeGen/Program.hs +216/−0
- Obsidian/CodeGen/SPMDC.hs +590/−0
- Obsidian/DimSpec.hs +7/−0
- Obsidian/Exp.hs +816/−0
- Obsidian/Force.hs +68/−0
- Obsidian/Globs.hs +16/−0
- Obsidian/Library.hs +606/−0
- Obsidian/LibraryG.hs +129/−0
- Obsidian/Memory.hs +184/−0
- Obsidian/Names.hs +20/−0
- Obsidian/Program.hs +263/−0
- Obsidian/SeqLoop.hs +102/−0
- Obsidian/Types.hs +33/−0
- Setup.hs +2/−0
+ LICENSE view
@@ -0,0 +1,28 @@++BSD3 Full Text:+--------------------------------------------------------------------------------++Copyright (c) 2011-2013, Joel Svensson +All rights reserved.++Redistribution and use in source and binary forms, with or without+modification, are permitted provided that the following conditions are met:+ * Redistributions of source code must retain the above copyright+ notice, this list of conditions and the following disclaimer.+ * Redistributions in binary form must reproduce the above copyright+ notice, this list of conditions and the following disclaimer in the+ documentation and/or other materials provided with the distribution.+ * Neither the name of the <organization> nor the+ names of its contributors may be used to endorse or promote products+ derived from this software without specific prior written permission.++THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND+ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED+WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE+DISCLAIMED. IN NO EVENT SHALL <COPYRIGHT HOLDER> BE LIABLE FOR ANY+DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES+(INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;+LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND+ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT+(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS+SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ Obsidian.cabal view
@@ -0,0 +1,79 @@+Name: Obsidian+Version: 0.0.0.1++License: BSD3+License-file: LICENSE+Stability: Beta+Maintainer: Joel Svensson<svenssonjoel@yahoo.se>+Author: Joel Svensson<svenssonjoel@yahoo.se>++Copyright: Copyright (c) 2011-2013 Joel Svensson + +Synopsis: Embedded language for GPU Programming +HomePage: https://github.com/svenssonjoel/Obsidian+Description: + Obsidian is an embedded language for general purpose programming targeting + GPU's. ++Category: Language+Cabal-Version: >=1.8+Tested-With: GHC == 7.6.1++build-type: Simple++--extra-source-files:+-- examples/tests/Test_DotProd.hs+++source-repository head+ type: git+ location: git://github.com/svenssonjoel/Obsidian.git++----------------------------------------------------------------------------------------------------+Library+ build-depends: base >= 4 && < 5+ , mtl >= 2.0 + , value-supply >= 0.6+ , containers >= 0.4.2.1+ + exposed-modules: Obsidian+ + + other-modules: Obsidian.Array+ , Obsidian.Atomic+ , Obsidian.DimSpec+ , Obsidian.Exp+ , Obsidian.Force+ , Obsidian.Globs+ , Obsidian.Library+ , Obsidian.LibraryG +-- , Obsidian.Lift+ , Obsidian.Memory+ , Obsidian.Names+ , Obsidian.Program + , Obsidian.SeqLoop + , Obsidian.Types +-- , Obsidian.CodeGen.C+ , Obsidian.CodeGen.CUDA+ , Obsidian.CodeGen.Common+ , Obsidian.CodeGen.InOut+ , Obsidian.CodeGen.Liveness+ , Obsidian.CodeGen.Memory+-- , Obsidian.CodeGen.OpenCL+ , Obsidian.CodeGen.PP + , Obsidian.CodeGen.Program+ , Obsidian.CodeGen.SPMDC++ + GHC-Options: +-- -O2 ++-- include-dirs:+++-- extra-lib-dirs: ++-- Includes: +-- Extra-libraries: + +
+ Obsidian.hs view
@@ -0,0 +1,29 @@+module Obsidian (module Obsidian.Array,+ module Obsidian.Program,+ module Obsidian.Exp, + module Obsidian.Types, + module Obsidian.Force, + module Obsidian.Library,+ module Obsidian.LibraryG,+ module Obsidian.CodeGen.InOut,+ module Obsidian.CodeGen.CUDA,+ module Obsidian.Atomic, + module Obsidian.SeqLoop, + module Obsidian.Memory, + module Obsidian.Names ) where++++import Obsidian.Program+import Obsidian.Exp+import Obsidian.Types+import Obsidian.Array+import Obsidian.Library+import Obsidian.LibraryG+import Obsidian.Force+import Obsidian.CodeGen.InOut+import Obsidian.CodeGen.CUDA +import Obsidian.Atomic+import Obsidian.SeqLoop+import Obsidian.Memory+import Obsidian.Names
+ Obsidian/Array.hs view
@@ -0,0 +1,179 @@+{-# LANGUAGE MultiParamTypeClasses, + FlexibleInstances, FlexibleContexts,+ GADTs, + TypeFamilies,+ RankNTypes #-} ++{- Joel Svensson 2012++ Notes:+ 2013-01-08: Removed number-of-blocks field from Distribs+ 2012-12-10: Drastically shortened. +-}++module Obsidian.Array where++import Obsidian.Exp +import Obsidian.Types+import Obsidian.Globs+import Obsidian.Program+import Obsidian.Names++import Data.List+import Data.Word++---------------------------------------------------------------------------+-- Aliases+---------------------------------------------------------------------------+type SPull = Pull Word32+type DPull = Pull EWord32++type SPush t a = Push t Word32 a+type DPush t a = Push t EWord32 a +---------------------------------------------------------------------------+-- Create arrats+---------------------------------------------------------------------------+undefinedGlobal n = Pull n $ \gix -> undefined+namedGlobal name n = Pull n $ \gix -> index name gix+namedPull name n = Pull n $ \gix -> index name gix++---------------------------------------------------------------------------+-- Class ArraySize+--------------------------------------------------------------------------- +class (Integral a, Num a) => ASize a where+ sizeConv :: a -> Exp Word32++instance ASize Word32 where+ sizeConv = fromIntegral++instance ASize (Exp Word32) where+ sizeConv = id ++---------------------------------------------------------------------------+-- Push and Pull arrays+---------------------------------------------------------------------------+data Push p s a =+ Push s (forall b. ((a -> Exp Word32 -> TProgram NameInfo) -> Program p NameInfo))++data Pull s a = Pull {pullLen :: s, + pullFun :: Exp Word32 -> a}++mkPushArray :: s -> (forall b. ((a -> Exp Word32 -> TProgram NameInfo)+ -> Program t NameInfo)) -> Push t s a+mkPushArray n p = Push n p +mkPullArray n p = Pull n p ++class Array a where+ resize :: r -> a s e -> a r e+ len :: ASize s => a s e -> s+ aMap :: (e -> e') -> a s e -> a s e'+ ixMap :: (Exp Word32 -> Exp Word32)+ -> a s e -> a s e+ +instance Array Pull where + resize m (Pull _ ixf) = Pull m ixf+ len (Pull s _) = s+ aMap f (Pull n ixf) = Pull n (f . ixf)+ ixMap f (Pull n ixf) = Pull n (ixf . f) + +instance Array (Push t) where + resize m (Push _ p) = Push m p+ len (Push s _) = s+ 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))+ ++class Indexible a where + access :: a s e -> Exp Word32 -> e + +instance Indexible Pull where+ access p ix = pullFun p ix++---------------------------------------------------------------------------+-- Functor instance Pull/Push arrays+---------------------------------------------------------------------------+instance Array arr => Functor (arr w) where + fmap = aMap+++---------------------------------------------------------------------------+-- Pushable+---------------------------------------------------------------------------+class Pushable a where + push :: ASize s => PT t -> a s e -> Push t s e+ -- Push using m threads+ -- m must be a divisor of nm (TODO: error otherwise) + --pushN :: Word32 -> a e -> Push e++ -- push grouped elements to adjacent indices using+ -- one thread per group. + --pushF :: a [e] -> Push e + +instance Pushable (Push Thread) where + push Thread = id+ push Block = error "not implemented: Program transformation!"+ push Grid = error "not implemented: Program transformation!" +instance Pushable (Push Block) where + push Block = id + push Thread = error "not implemented: Program transformations!"+ push Grid = error "not implemented: Program transformations!" +instance Pushable (Push Grid) where + push Grid = id+ push Thread = error "not implemented: Program transformations!"+ push Block = error "not implemented: Program transformations!" + +instance Pushable Pull where+ push Thread (Pull n ixf) =+ Push n $ \wf -> seqFor (sizeConv n) $ \i -> wf (ixf i) i+ push Block (Pull n ixf) =+ Push n $ \wf -> ForAll (sizeConv n) $ \i -> wf (ixf i) i + push Grid (Pull n ixf) =+ Push n $ \wf -> ForAllThreads (sizeConv n) $ \i -> wf (ixf i) i ++{- pushN m (Pull nm ixf) =+ Push nm -- There are still nm elements (info for alloc) + $ \wf ->+ ForAll (Just m) + $ \i ->+ sequence_ [wf (ixf (i + fromIntegral (j * n))) (i + (fromIntegral (j * n)))+ | j <- [0..n]] + -- Force can still Allocate n elements for this Push array.+ where+ n = fromIntegral (nm `div` m)++ pushF (Pull n ixf) =+ Push (n * m) $ \wf ->+ ForAll (Just n) $ \i ->+ sequence_ [wf ((ixf i) !! fromIntegral j) (i * fromIntegral m + fromIntegral j)+ | j <- [0..m]]+ where + m = fromIntegral$ length $ ixf 0+-} ++{- ------------------------------------------------------------------------++ m m m m m + |-----|-----|-----|-----|-----| (n * m)++ 01234 01234 01234 01234 01234 k ++ 0 1 2 3 4 j++ m threads, each writing n elements:+ [(tid + (j*m) | j <- [0..n]] ++ n threads, each writing m elements:+ [(tid * m + k | k <- [0..m]] ++------------------------------------------------------------------------ -} +---------------------------------------------------------------------------+-- Indexing, array creation.+---------------------------------------------------------------------------+namedArray name n = mkPullArray n (\ix -> index name ix)+indexArray n = mkPullArray n (\ix -> ix)++pushApp (Push _ p) a = p a ++infixl 9 ! +(!) :: Indexible a => a s e -> Exp Word32 -> e +(!) = access
+ Obsidian/Atomic.hs view
@@ -0,0 +1,21 @@+{-# LANGUAGE GADTs #-}++{- Joel Svensson,+ Josef Svenningsson+ 2012 -} +module Obsidian.Atomic where++import Obsidian.Exp+import Data.Word+ ++---------------------------------------------------------------------------+-- Atomic operations +---------------------------------------------------------------------------+data Atomic a where++ -- Cuda only allows AtomicInc on the Int type+ -- (todo: figure out if CUDA int is 32 or 64 bit) + AtomicInc :: Atomic Word32 ++printAtomic AtomicInc = "atomicInc"
+ Obsidian/CodeGen/CUDA.hs view
@@ -0,0 +1,209 @@+{- Joel Svensson 2012, 2013 -}++{-# LANGUAGE GADTs #-} +module Obsidian.CodeGen.CUDA + (genKernel) where ++import Data.List+import Data.Word +import Data.Monoid+import qualified Data.Map as Map+import Control.Monad.State++import Obsidian.Array+import Obsidian.Exp ++import Obsidian.Types+import Obsidian.Globs+import Obsidian.Atomic ++import Obsidian.CodeGen.PP+import Obsidian.CodeGen.Common+import Obsidian.CodeGen.InOut +import Obsidian.CodeGen.Memory+import Obsidian.CodeGen.Liveness++-- New imports+import Obsidian.CodeGen.Program +import qualified Obsidian.Program as P ++import Obsidian.CodeGen.SPMDC++---------------------------------------------------------------------------+-- a gc+--------------------------------------------------------------------------- +gc = genConfig "" ""+++---------------------------------------------------------------------------+-- C style function "header"+---------------------------------------------------------------------------++kernelHead :: Name -> + [(String,Type)] -> + [(String,Type)] -> + PP () +kernelHead name ins outs = + do + line ("__global__ void " ++ name ++ "(" ++ types ++ ")" ) + where + types = concat (intersperse "," (typeList (ins ++ outs)))+ typeList :: [(String,Type)] -> [String] + typeList [] = [] + typeList ((a,t):xs) = (genType gc t ++ a) : typeList xs+ +---------------------------------------------------------------------------+-- genKernel +---------------------------------------------------------------------------+ +genKernel :: ToProgram a b => String -> (a -> b) -> Ips a b -> String +genKernel name kernel a = proto ++ ts ++ cuda + where+ (ins,im) = toProgram 0 kernel a++ outs = getOutputs im+ + lc = computeLiveness im + + -- Creates (name -> memory address) map + (m,mm) = mmIM lc sharedMem Map.empty+ + -- What if its Right ??? (I DONT KNOW!) + (Left threadBudget) = numThreads im+ ts = "/* number of threads needed " ++ show threadBudget ++ "*/\n"++ spmd = imToSPMDC threadBudget im+ + + body' = (if size m > 0 then (shared :) else id) $ mmSPMDC mm spmd++ em = snd $ execState (collectExps body') ( 0, Map.empty)+ (decls,body'') = replacePass em body'+ spdecls = declsToSPMDC decls ++ body = spdecls ++ body''+ + swap (x,y) = (y,x)+ inputs = map ((\(t,n) -> (typeToCType t,n)) . swap) ins+ outputs = map ((\(t,n) -> (typeToCType t,n)) . swap) outs + + ckernel = CKernel CQualifyerKernel CVoid name (inputs++outputs) body+ shared = CDecl (CQualified CQualifyerExtern (CQualified CQualifyerShared ((CQualified (CQualifyerAttrib (CAttribAligned 16)) (CArray [] (CWord8)))))) "sbase"++ proto = getProto name ins outs + cuda = printCKernel (PPConfig "__global__" "" "" "__syncthreads()") ckernel +++---------------------------------------------------------------------------+-- Generate a function prototype+--------------------------------------------------------------------------- +getProto :: Name -> [(String,Type)] -> [(String,Type)] -> String+getProto name ins outs =+ runPP (+ do + line "extern \"C\" "+ kernelHead name ins outs+ line ";"+ newline) 0 ++---------------------------------------------------------------------------+-- generate a sbase CExpr+---------------------------------------------------------------------------+sbaseCExpr 0 = cVar "sbase" (CPointer CWord8) +sbaseCExpr addr = cBinOp CAdd (cVar "sbase" (CPointer CWord8)) + (cLiteral (Word32Val addr) CWord32) + (CPointer CWord8) +---------------------------------------------------------------------------+-- Memory map the arrays in an SPMDC+---------------------------------------------------------------------------+mmSPMDC :: MemMap -> [SPMDC] -> [SPMDC] +mmSPMDC mm [] = [] +mmSPMDC mm (x:xs) = mmSPMDC' mm x : mmSPMDC mm xs++mmSPMDC' :: MemMap -> SPMDC -> SPMDC+mmSPMDC' mm (CAssign e1 es e2) = + cAssign (mmCExpr mm e1) + (map (mmCExpr mm) es) + (mmCExpr mm e2)+mmSPMDC' mm (CAtomic op e1 e2 e3) = cAtomic op (mmCExpr mm e1)+ (mmCExpr mm e2)+ (mmCExpr mm e3) +mmSPMDC' mm (CFunc name es) = cFunc name (map (mmCExpr mm) es) +mmSPMDC' mm CSync = CSync+mmSPMDC' mm (CIf e s1 s2) = cIf (mmCExpr mm e) (mmSPMDC mm s1) (mmSPMDC mm s2)+mmSPMDC' mm (CFor name e s) = cFor name (mmCExpr mm e) (mmSPMDC mm s)+mmSPMDC' mm (CDeclAssign t nom e) = cDeclAssign t nom (mmCExpr mm e)+mmSPMDC' mm a@(CDecl t nom) = a+mmSPMDC' mm a = error $ "mmSPMDC': " ++ show a+---------------------------------------------------------------------------+-- Memory map the arrays in an CExpr+---------------------------------------------------------------------------+mmCExpr mm (CExpr (CVar nom t)) = + case Map.lookup nom mm of + Just (addr,t) -> + let core = sbaseCExpr addr + cast c = cCast c (typeToCType t)+ in cast core+ + Nothing -> cVar nom t+mmCExpr mm (CExpr (CIndex (e1,es) t)) = cIndex (mmCExpr mm e1, map (mmCExpr mm) es) t+mmCExpr mm (CExpr (CBinOp op e1 e2 t)) = cBinOp op (mmCExpr mm e1) (mmCExpr mm e2) t+mmCExpr mm (CExpr (CUnOp op e t)) = cUnOp op (mmCExpr mm e) t +mmCExpr mm (CExpr (CFuncExpr nom exprs t)) = cFuncExpr nom (map (mmCExpr mm) exprs) t+mmCExpr mm (CExpr (CCast e t)) = cCast (mmCExpr mm e) t+mmCExpr mm (CExpr (CCond e1 e2 e3 t)) = cCond (mmCExpr mm e1)+ (mmCExpr mm e2)+ (mmCExpr mm e3)+ t+mmCExpr mm a = a + + +---------------------------------------------------------------------------+-- New IM to SPCMD+---------------------------------------------------------------------------+atomicOpToCAtomicOp AtomicInc = CAtomicInc++imToSPMDC :: Word32 -> IMList a -> [SPMDC]+imToSPMDC nt im = concatMap (process nt) im+ where+ process nt (SAssign name [] e,_) =+ [cAssign (cVar name (typeToCType (typeOf e))) [] (expToCExp e)]++ process nt (SAssign name [ix] e,_) = + [cAssign (cVar name (typeToCType (Pointer (typeOf e)))) [expToCExp ix] (expToCExp e)]++ process nt (SAtomicOp res arr e op,_) = + [cAtomic (atomicOpToCAtomicOp op)+ (cVar res (typeToCType (typeOf e)))+ (cVar arr (typeToCType (Pointer (typeOf e))))+ (expToCExp e)]++ process nt (SCond bexp im,_) =+ [cIf (expToCExp bexp) (imToSPMDC nt im) []]++ process nt (SSeqFor name e im,_) =+ [cFor name (expToCExp e) (imToSPMDC nt im)]++ process nt (SForAll (Literal n) im,_) =+ if (n < nt) + then + [cIf (cBinOp CLt (cThreadIdx X) (cLiteral (Word32Val n) CWord32) CInt)+ code []]+ else + code + where + code = imToSPMDC nt im++ -- This one is tricky (since no corresponding CUDA construct exists) + process nt (SForAllBlocks n im,_) =+ -- TODO: there should be "number of blocks"-related conditionals here (possibly) + imToSPMDC nt im+ -- This one is even more tricky+ process nt (SForAllThreads n im,_) =+ imToSPMDC nt im + process nt (SAllocate name size t,_) = []+ process nt (SDeclare name t,_) =+ [cDecl (typeToCType t) name]+ process nt (SOutput name t,_) = [] -- RIGHT!+ process nt (SSynchronize,_) = [CSync]+
+ Obsidian/CodeGen/Common.hs view
@@ -0,0 +1,176 @@+{-# LANGUAGE RankNTypes, GADTs #-}++{- Joel Svensson 2012 -} +module Obsidian.CodeGen.Common where ++import Data.List+import Data.Word+import qualified Data.Map as Map +++import Obsidian.Exp +import Obsidian.Types+import Obsidian.Globs++import Obsidian.CodeGen.PP+import Obsidian.CodeGen.Memory++---------------------------------------------------------------------------+data GenConfig = GenConfig { global :: String,+ local :: String };+ +genConfig = GenConfig+++---------------------------------------------------------------------------+-- Helpers++mappedName :: Name -> Bool +mappedName name = isPrefixOf "arr" name++tid :: Exp Word32+tid = ThreadIdx X++genType _ Int = "int "+genType _ Int8 = "int8_t "+genType _ Int16 = "int16_t "+genType _ Int32 = "int32_t "+genType _ Int64 = "int64_t "+genType _ Float = "float "+genType _ Double = "double "+genType _ Bool = "int " +genType _ Word8 = "uint8_t "+genType _ Word16 = "uint16_t "+genType _ Word32 = "uint32_t "+genType _ Word64 = "uint64_t "++genType gc (Pointer t) = genType gc t ++ "*"+genType gc (Global t) = global gc ++" "++ genType gc t -- "__global " ++ genType t+genType gc (Local t) = local gc ++" "++ genType gc t ++genCast gc t = "(" ++ genType gc t ++ ")"++parens s = '(' : s ++ ")"+ +---------------------------------------------------------------------------+-- genExp C-style +genExp :: Scalar a => GenConfig -> MemMap -> Exp a -> [String]++-- Cheat and do CUDA printing here as well+genExp gc _ (BlockDim X) = ["blockDim.x"]+genExp gc _ (BlockIdx X) = ["blockIdx.x"]+genExp gc _ (BlockIdx Y) = ["blockIdx.y"]+genExp gc _ (BlockIdx Z) = ["blockIdx.z"]+genExp gc _ (ThreadIdx X) = ["threadIdx.x"]+genExp gc _ (ThreadIdx Y) = ["threadIdx.y"]+genExp gc _ (ThreadIdx Z) = ["threadIdx.z"]+++genExp gc _ (Literal a) = [show a] +genExp gc _ (Index (name,[])) = [name]+genExp gc mm exp@(Index (name,es)) = + [name' ++ genIndices gc mm es]+ where + (offs,t) = + case Map.lookup name mm of + Nothing -> error "array does not excist in map" + (Just x) -> x+ name' = if mappedName name + then parens$ genCast gc t ++ + if offs > 0 + then "(sbase+" ++ show offs ++ ")" + else "sbase"+ else name++ +genExp gc mm (BinOp op e1 e2) = + [genOp op (genExp gc mm e1 ++ genExp gc mm e2)]++genExp gc mm (UnOp op e) = + [genOp op (genExp gc mm e)] + +genExp gc mm (If b e1 e2) = + [genIf (genExp gc mm b ++ + genExp gc mm e1 ++ + genExp gc mm e2 )] ++---------------------------------------------------------------------------+--+genIndices gc mm es = concatMap (pIndex mm) es + where + pIndex mm e = "[" ++ concat (genExp gc mm e) ++ "]"+++genIf [b,e1,e2] = "(" ++ b ++ " ? " ++ e1 ++ " : " ++ e2 ++ ")"++---------------------------------------------------------------------------+-- genOp+genOp :: Op a -> [String] -> String+genOp Add [a,b] = oper "+" a b +genOp Sub [a,b] = oper "-" a b +genOp Mul [a,b] = oper "*" a b +genOp Div [a,b] = oper "/" a b ++genOp Mod [a,b] = oper "%" a b ++genOp Sin [a] = func "sin" a +genOp Cos [a] = func "cos" a +-- Bool ops+genOp Eq [a,b] = oper "==" a b +genOp Lt [a,b] = oper "<" a b +genOp LEq [a,b] = oper "<=" a b +genOp Gt [a,b] = oper ">" a b+genOp GEq [a,b] = oper ">=" a b++-- Bitwise ops+genOp BitwiseAnd [a,b] = oper "&" a b +genOp BitwiseOr [a,b] = oper "|" a b +genOp BitwiseXor [a,b] = oper "^" a b +genOp BitwiseNeg [a] = unOp "~" a +genOp ShiftL [a,b] = oper "<<" a b +genOp ShiftR [a,b] = oper ">>" a b +++-- built-ins +genOp Min [a,b] = func "min" (a ++ "," ++ b) +genOp Max [a,b] = func "max" (a ++ "," ++ b) ++genOp Int32ToWord32 [a] = func "(uint32_t)" a+genOp Word32ToInt32 [a] = func "(int32_t)" a ++func f a = f ++ "(" ++ a ++ ")" +oper f a b = "(" ++ a ++ f ++ b ++ ")" +unOp f a = "(" ++ f ++ a ++ ")"++---------------------------------------------------------------------------+-- Configurations, threads,memorymap ++data Config = Config {configThreads :: NumThreads, + configMM :: MemMap,+ configLocalMem :: Word32} +config = Config+++assign :: Scalar a => GenConfig -> MemMap -> Exp a -> Exp a -> PP () +assign gc mm name val = line ((concat (genExp gc mm name)) ++ + " = " ++ concat (genExp gc mm val) ++ + ";") + +cond :: GenConfig -> MemMap -> Exp Bool -> PP () +cond gc mm e = line ("if " ++ concat (genExp gc mm e)) ++++-- used in both OpenCL and CUDA generation+potentialCond gc mm n nt pp+ | n < nt = + do+ cond gc mm (tid <* (fromIntegral n))+ begin+ pp + end + | n == nt = pp+ + | otherwise = error "potentialCond: should not happen"++
+ Obsidian/CodeGen/InOut.hs view
@@ -0,0 +1,155 @@+{-# LANGUAGE FlexibleInstances,+ OverlappingInstances,+ UndecidableInstances,+ FlexibleContexts,+ MultiParamTypeClasses,+ TypeOperators,+ TypeFamilies ,+ ScopedTypeVariables+ #-}++{- Joel Svensson 2012, 2013+ Niklas Ulvinge 2013++ Notes:++ 2013-01-24: Changes with the new Array types in mind+ 2013-01-08: Edited+ 2012-12-10: Edited++-} +++module Obsidian.CodeGen.InOut where ++import Obsidian.Exp +import Obsidian.Array++import Obsidian.Types+import Obsidian.Globs +import Obsidian.Program+import Obsidian.Force+import Obsidian.Memory++import qualified Obsidian.CodeGen.Program as CG ++import Data.Word+import Data.Int+ +---------------------------------------------------------------------------+-- New approach (hopefully)+---------------------------------------------------------------------------+-- "reify" Haskell functions into CG.Programs++{-+ Blocks needs to be of specific sizes (a design choice we've made).+ Because of this a prototypical input array needs to be provided+ that has a static block size (the number of blocks is dynamic).++ To make things somewhat general a heterogeneous list of input arrays+ that has same shape as the actual parameter list of the function+ is passed into toProgram (the reifyer). ++-} + +type Inputs = [(Name,Type)] ++class ToProgram a b where+ toProgram :: Int -> (a -> b) -> Ips a b -> (Inputs,CG.IM)+ +typeOf_ a = typeOf (Literal a)+++---------------------------------------------------------------------------+-- Base cases+--------------------------------------------------------------------------- +instance (Scalar t) => ToProgram (Exp t) (GProgram b) where+ toProgram i f a = ([(nom,t)],CG.compileStep1 (f input))+ where nom = "s" ++ show i+ input = variable nom+ t = typeOf_ (undefined :: t)++instance (Scalar t) => ToProgram (Pull (Exp Word32) (Exp t)) (GProgram a) where+ toProgram i f (Pull n ixf) = ([(nom,Pointer t),(n,Word32)],CG.compileStep1 (f input)) + where nom = "input" ++ show i+ n = "n" ++ show i + lengthVar = variable n+ input = namedGlobal nom lengthVar+ t = typeOf_ (undefined :: t)++instance (Scalar t) => ToProgram (Pull Word32 (Exp t)) (GProgram a) where+ toProgram i f (Pull n ixf) = ([(nom,Pointer t){-,(n,Word32)-}],CG.compileStep1 (f input)) + where nom = "input" ++ show i+ --n = "n" ++ show i + --lengthVar = variable n+ input = namedGlobal nom n -- lengthVar+ t = typeOf_ (undefined :: t)++---------------------------------------------------------------------------+-- More natural to work with these in some cases+---------------------------------------------------------------------------+instance (ToProgram b (GProgram ()),+ GlobalMemoryOps a)+ => ToProgram b (Push Grid Word32 a) where+ toProgram i f arr = toProgram i (forceG . f) arr++instance (ToProgram b (GProgram ()),+ GlobalMemoryOps a)+ => ToProgram b (Push Grid EWord32 a) where+ toProgram i f arr = toProgram i (forceG . f) arr+++---------------------------------------------------------------------------+-- Recursive cases+--------------------------------------------------------------------------- +instance (Scalar t, ToProgram b c) => ToProgram (Exp t) (b -> c) where+ toProgram i f (a :-> rest) = ((nom,t):ins,prg)+ where+ (ins,prg) = toProgram (i+1) (f input) rest+ nom = "s" ++ show i+ input = variable nom+ t = typeOf_ (undefined :: t)++instance (Scalar t, ToProgram b c) => ToProgram (Pull (Exp Word32) (Exp t)) (b -> c) where+ toProgram i f ((Pull n ixf) :-> rest) = ((nom,Pointer t):(n,Word32):ins,prg)+ where+ (ins,prg) = toProgram (i+1) (f input) rest+ nom = "input" ++ show i+ n = "n" ++ show i+ lengthVar = variable n+ input = namedGlobal nom lengthVar+ t = typeOf_ (undefined :: t)+++instance (Scalar t, ToProgram b c) => ToProgram (Pull Word32 (Exp t)) (b -> c) where+ toProgram i f ((Pull n ixf) :-> rest) = ((nom,Pointer t){-:(n,Word32)-}:ins,prg)+ where+ (ins,prg) = toProgram (i+1) (f input) rest+ nom = "input" ++ show i+ --n = "n" ++ show i+ --lengthVar = variable n+ input = namedGlobal nom n --lengthVar+ t = typeOf_ (undefined :: t)+++---------------------------------------------------------------------------+-- heterogeneous lists of inputs +---------------------------------------------------------------------------+data head :-> tail = head :-> tail++infixr 5 :->+++---------------------------------------------------------------------------+-- Function types to input list types. +--------------------------------------------------------------------------- +type family Ips a b+ +-- type instance Ips a (GlobArray b) = Ips' a -- added Now 26+-- type instance Ips a (Final (GProgram b)) = a+type instance Ips a (Push Grid l b) = a+type instance Ips a (Pull l b) = a+type instance Ips a (GProgram b) = a+type instance Ips a (b -> c) = a :-> Ips b c++
+ Obsidian/CodeGen/Liveness.hs view
@@ -0,0 +1,144 @@+++{- Joel Svensson 2012, 2013++ notes:+ added case for SeqFor Jan-21-2013++-} +module Obsidian.CodeGen.Liveness where++import qualified Data.Set as Set++import Obsidian.Exp+import Obsidian.Globs+import Obsidian.CodeGen.Program++import Control.Monad.State+++---------------------------------------------------------------------------+--+---------------------------------------------------------------------------+type Liveness = Set.Set Name++--------------------------------------------------------------------------- +--+--------------------------------------------------------------------------- +type IML = [(Statement Liveness,Liveness)] +++{- Plan:+ # Step through program from end to start+ # as soon as a new name is encountered, add it to the living set+ # when an "Allocate" is found, delete the name it allocated from the living set.++ Requirements:+ # All names are unique! ++ TODO: Think more carefully about the ForAllBlocks case+ TODO: Can ixs contain array names ?+ (Most likely yes! think about the counting sort example)++-} + ++-- Nice type +computeLiveness :: IMList a -> IML+computeLiveness im = reverse $ evalState (cl (reverse im)) Set.empty++-- Nice Type +computeLiveness1 :: Liveness -> IMList a -> IML+computeLiveness1 l im = reverse $ evalState (cl (reverse im)) l++-- cl :: IM -> State Liveness IML +cl im = mapM process im+ where+ safeHead [] = Set.empty+ safeHead (x:xs) = snd x++ -- Horrific type + process :: (Statement a,a) -> State Liveness (Statement Liveness,Liveness)+ process (SAssign nom ixs e,_) =+ do+ s <- get+ let arrays = collectArrays e+ living = Set.fromList (nom:arrays) `Set.union` s+ + put living -- update state + return (SAssign nom ixs e,living)+ + process (SAtomicOp n1 n2 ixs op,_) =+ do+ s <- get+ return (SAtomicOp n1 n2 ixs op,s)+ + process (SAllocate name size t,_) =+ do+ modify (name `Set.delete`)+ s <- get + return (SAllocate name size t,s) + + process (SDeclare name t,_) = + do + s <- get + return (SDeclare name t,s)++ process (SOutput name t,_) = + do + s <- get + return (SOutput name t,s)++ process (SSynchronize,_) = + do + s <- get+ return (SSynchronize,s) ++ process (SCond bexp im,_) = + do+ -- TODO: What should really happen here ?+ s <- get + let iml = computeLiveness1 s im + l = safeHead iml + ns = s `Set.union` l+ put ns + -- Is this correct ? Same question, all below+ return (SCond bexp iml,ns)++ process (SSeqFor nom n im,_) = + do + s <- get+ let iml = computeLiveness1 s im + l = safeHead iml+ ns = s `Set.union` l+ put ns+ return (SSeqFor nom n iml,ns) +++ process (SForAll n im,_) = + do + s <- get + let iml = computeLiveness1 s im + l = safeHead iml + ns = s `Set.union` l+ put ns+ return (SForAll n iml,ns) + + process (SForAllBlocks n im,_) = + do + s <- get + let iml = computeLiveness1 s im + l = safeHead iml + ns = s `Set.union` l+ put ns+ return (SForAllBlocks n iml,ns)++ process (SForAllThreads n im,_) = + do + s <- get + let iml = computeLiveness1 s im + l = safeHead iml + ns = s `Set.union` l+ put ns + return (SForAllThreads n iml,ns)+
+ Obsidian/CodeGen/Memory.hs view
@@ -0,0 +1,152 @@++{- 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,+ --mapMemory,+ -- NEW+ mmIM) + 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)+ + +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')+ + 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+ -- Yet another tricky case.+ process (SForAll n im,_) m mm = mmIM im m mm + -- The worst of them all.+ process (SForAllThreads n im,_) m mm = mmIM im m mm++ 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+
+ Obsidian/CodeGen/PP.hs view
@@ -0,0 +1,55 @@++{- Joel Svensson 2012 -}+module Obsidian.CodeGen.PP where +++import Control.Monad.State+------------------------------------------------------------------------------+-- print and indent and stuff... +-- This is probably very ugly ++-- TODO: There is a chapter about this pretty printing in "implementing functional lang..." +-- Look at that and learn +++type PP a = State (Int,String) a ++indent :: PP ()+indent = + do + (i,s) <- get + put (i+1,s) + +unindent :: PP () +unindent = + do + (i,s) <- get + if i <= 0 then error "Whats going on" else put (i-1,s) ++line :: String -> PP () +line str = + do + (i,s) <- get + put (i,s ++ str) ++ +newline :: PP () +newline = + do + (i,s) <- get + let ind = replicate (i*2) ' '+ put (i,s ++ "\n" ++ ind)+ +runPP :: PP a -> Int -> String+runPP pp i = snd$ execState pp (i,"")++begin :: PP () +begin = line "{" >> indent >> newline++end :: PP () +end = unindent >> newline >> line "}" >> newline++space = line " " +cTermLn = line ";" >> newline++wrap s e p = line s >> p >> line e
+ Obsidian/CodeGen/Program.hs view
@@ -0,0 +1,216 @@+{-# LANGUAGE GADTs,+ ExistentialQuantification,+ FlexibleInstances #-}++{- CodeGen.Program.++ Joel Svensson 2012, 2013++ Notes:+ 2013-03-17: Codegeneration is changing+++++-} +++module Obsidian.CodeGen.Program where++import Obsidian.Exp+import Obsidian.Globs+import Obsidian.Types+import Obsidian.Atomic++import qualified Obsidian.Program as P ++import Data.Word+import Data.Supply+import Data.List++import System.IO.Unsafe++---------------------------------------------------------------------------+-- New Intermediate representation+---------------------------------------------------------------------------++type IMList a = [(Statement a,a)]++type IM = IMList ()++-- out :: +out a = [(a,())]+++data Statement t = forall a. (Show a, Scalar a) => SAssign Name [Exp Word32] (Exp a)+ | forall a. (Show a, Scalar a) => SAtomicOp Name Name (Exp Word32) (Atomic a)+ | SCond (Exp Bool) (IMList t) + | SSeqFor String (Exp Word32) (IMList t)+ -- See if it is possible to get away+ -- with only one kind of ForAll (plus maybe some flag) + | SForAll (Exp Word32) (IMList t) + | SForAllBlocks (Exp Word32) (IMList t)+ -- a special loop over all threads..+ | SForAllThreads (Exp Word32) (IMList t)++ -- Memory Allocation..+ | SAllocate Name Word32 Type+ | SDeclare Name Type+ | SOutput Name Type++ -- Synchronisation+ | SSynchronize++ -- ProgramPar and ProgramSeq does not exist+ -- at this level (the par or seq info is lost!)+ ++compileStep1 :: P.Program t a -> IM+compileStep1 p = snd $ cs1 ns p+ where+ ns = unsafePerformIO$ newEnumSupply++cs1 :: Supply Int -> P.Program t a -> (a,IM) +cs1 i P.Identifier = (supplyValue i, [])++cs1 i (P.Assign name ix e) =+ ((),out (SAssign name ix e))+ +cs1 i (P.AtomicOp name ix at) = (v,out im)+ where + nom = "a" ++ show (supplyValue i)+ v = variable nom+ im = SAtomicOp nom name ix at+ +cs1 i (P.Cond bexp p) = ((),out (SCond bexp im)) + where ((),im) = cs1 i p++cs1 i (P.SeqFor n f) = (a,out (SSeqFor nom n im))+ where+ (i1,i2) = split2 i+ nom = "i" ++ show (supplyValue i1)+ v = variable nom+ p = f v+ (a,im) = cs1 i2 p + +cs1 i (P.ForAll n f) = (a,out (SForAll n im))+ where+ p = f (ThreadIdx X) + (a,im) = cs1 i p ++cs1 i (P.ForAllBlocks n f) = (a,out (SForAllBlocks n im)) + where+ p = f (BlockIdx X)+ (a,im) = cs1 i p+++-- Warning: Every thread will ALWAYS need to perform a conditional+-- (Only in special case is the conditional not needed) +-- TRY To express all library functions using ForAllBlocks + ForAll+-- For more flexibility and probably in the end performance. +cs1 i (P.ForAllThreads n f) = (a,out (SForAllThreads n im)) + where+ p = f (BlockIdx X * BlockDim X + ThreadIdx X)+ (a,im) = cs1 i p+++cs1 i (P.Allocate id n t) = ((),out (SAllocate id n t))+cs1 i (P.Declare id t) = ((),out (SDeclare id t))+-- Output works in a different way! (FIX THIS!)+-- Uniformity! (Allocate Declare Output) +cs1 i (P.Output t) = (nom,out (SOutput nom t))+ where nom = "output" ++ show (supplyValue i) +cs1 i (P.Sync) = ((),out (SSynchronize))+++cs1 i (P.Bind p f) = (b,im1 ++ im2) + where+ (s1,s2) = split2 i+ (a,im1) = cs1 s1 p+ (b,im2) = cs1 s2 (f a)++cs1 i (P.Return a) = (a,[])+++---------------------------------------------------------------------------+-- Analysis+--------------------------------------------------------------------------- +numThreads :: IMList a -> Either Word32 (EWord32)+numThreads im = foldl maxCheck (Left 0) $ map process im+ where+ process (SCond bexp im,_) = numThreads im+ process (SSeqFor _ _ _,_) = Left 1+ process (SForAll (Literal n) _,_) = Left n+ process (SForAll n _,_) = Right n+ process (SForAllBlocks _ im,_) = numThreads im+ process (SForAllThreads n im,_) = Right (variable "UNKNOWN") --fix this!+ process a = Left 0 -- ok ? ++ maxCheck (Left a) (Right b) = Right $ max (fromIntegral a) b+ maxCheck (Right a) (Left b) = Right $ max a (fromIntegral b)+ maxCheck (Left a) (Left b) = Left $ max a b+ maxCheck (Right a) (Right b) = Right $ max a b+++getOutputs :: IMList a -> [(Name,Type)]+getOutputs im = concatMap process im+ where+ process (SOutput name t,_) = [(name,t)]+ process (SSeqFor _ _ im,_) = getOutputs im+ process (SForAll _ im,_) = getOutputs im+ process (SForAllBlocks _ im,_) = getOutputs im+ process (SForAllThreads _ im,_) = getOutputs im+ process a = []+ ++---------------------------------------------------------------------------+-- Turning IM to strings+---------------------------------------------------------------------------++printIM :: Show a => IMList a -> String +printIM im = concatMap printStm im+ +-- Print a Statement with metadata +printStm :: Show a => (Statement a,a) -> String+printStm (SAssign name [] e,m) =+ name ++ " = " ++ printExp e ++ ";" ++ meta m+printStm (SAssign name ix e,m) =+ name ++ "[" ++ concat (intersperse "," (map printExp ix)) ++ "]" +++ " = " ++ printExp e ++ ";" ++ meta m+printStm (SAtomicOp res arr ix op,m) =+ res ++ " = " +++ printAtomic op ++ "(" ++ arr ++ "[" ++ printExp ix ++ "]);" ++ meta m+printStm (SAllocate name n t,m) =+ name ++ " = malloc(" ++ show n ++ ");" ++ meta m+printStm (SDeclare name t,m) =+ show t ++ " " ++ name ++ ";" ++ meta m+printStm (SOutput name t,m) =+ show t ++ " " ++ name ++ ";" ++ meta m+printStm (SCond bexp im,m) =+ "if " ++ show bexp ++ "{\n" ++ + concatMap printStm im ++ "\n};" ++ meta m++printStm (SSynchronize,m) =+ "sync();" ++ meta m+ +printStm (SSeqFor name n im,m) =+ "for " ++ name ++ " in [0.." ++ show n ++"] do" ++ meta m ++ + concatMap printStm im ++ "\ndone;\n"++printStm (SForAll n im,m) =+ "forAll i in [0.." ++ show n ++"] do" ++ meta m +++ concatMap printStm im ++ "\ndone;\n"++printStm (SForAllBlocks n im,m) =+ "forAllBlocks i in [0.." ++ show n ++"] do" ++ meta m +++ concatMap printStm im ++ "\ndone;\n"+printStm (SForAllThreads n im,m) =+ "forAllThreads i in [0.." ++ show n ++"] do" ++ meta m ++ + concatMap printStm im ++ "\ndone;\n"+++ +-- printStm (a,m) = error $ show m ++meta :: Show a => a -> String+meta m = "\t//" ++ show m ++ "\n"
+ Obsidian/CodeGen/SPMDC.hs view
@@ -0,0 +1,590 @@++{- Joel Svensson 2012,2013 -} +module Obsidian.CodeGen.SPMDC where++import Obsidian.Globs+import Obsidian.DimSpec++import Obsidian.CodeGen.PP++import Data.Word+import Data.Int++import qualified Data.List as L+import qualified Data.Map as M+import qualified Data.Set as S ++import Control.Monad.State++import Data.Maybe++-- TODO: Add Atomic ops ++---------------------------------------------------------------------------+-- A C LIKE AST (SPMDC - Single Program Multiple Data C) +--------------------------------------------------------------------------- +data Value = IntVal Int -- allow ? + | Int8Val Int8+ | Int16Val Int16+ | Int32Val Int32+ | Int64Val Int64+ | FloatVal Float + | DoubleVal Double+ | WordVal Word -- allow ? + | Word8Val Word8+ | Word16Val Word16+ | Word32Val Word32+ | Word64Val Word64+ deriving (Eq,Ord,Show)+ +data CType = CVoid | CInt | CFloat | CDouble+ | CInt8 | CInt16 | CInt32 | CInt64 + | CWord | CWord8 | CWord16 | CWord32 | CWord64+ | CPointer CType -- *type+ | CArray [CExpr] CType -- type[e1][e2][e3]..[en] or type[] + | CQualified CQualifyer CType + deriving (Eq,Ord,Show)+ +data CQualifyer = CQualifyerGlobal -- CUDA: "" OpenCL: "__global" + | CQualifyerLocal -- CUDA: "" OpenCL: "__local"+ | CQualifyerKernel -- CUDA: "__global__" OpenCL: "__kernel" + | CQualifyerShared -- CUDA: "__shared__" OpenCL: "__local" + | CQualifyerExtern -- extern + | CQualifyerAttrib CQAttribute+ deriving (Eq,Ord,Show)++data CQAttribute = CAttribAligned Word32+ deriving (Eq,Ord,Show)+++data CExprP e = CVar Name CType + -- Threads, Blocks, Grids (All of type Word32) + | CBlockIdx DimSpec + | CThreadIdx DimSpec+ | CBlockDim DimSpec+ | CGridDim DimSpec+ + | CLiteral Value CType+ | CIndex (e,[e]) CType+ | CCond e e e CType+ | CBinOp CBinOp e e CType+ | CUnOp CUnOp e CType+ | CFuncExpr Name [e] CType -- min, max, sin, cos + | CCast e CType -- cast expr to type + deriving (Eq,Ord,Show)+cTypeOfP (CVar _ t) = t+cTypeOfP (CBlockIdx d) = CWord32+cTypeOfP (CThreadIdx d) = CWord32+cTypeOfP (CBlockDim d) = CWord32+cTypeOfP (CGridDim d) = CWord32+cTypeOfP (CLiteral _ t) = t+cTypeOfP (CIndex _ t) = t+cTypeOfP (CCond _ _ _ t) = t+cTypeOfP (CBinOp _ _ _ t) = t+cTypeOfP (CUnOp _ _ t) = t+cTypeOfP (CFuncExpr _ _ t) = t+cTypeOfP (CCast _ t) = t++cSizeOf (CExpr (CIndex (e,es) _)) = 1 + max (cSizeOf e) (maximum (map cSizeOf es))+cSizeOf (CExpr (CCond e1 e2 e3 _)) = 1 + maximum [cSizeOf e1, cSizeOf e2, cSizeOf e3] +cSizeOf (CExpr (CFuncExpr _ es _)) = 1 + maximum (map cSizeOf es) +cSizeOf (CExpr (CUnOp _ e _)) = 1 + cSizeOf e +cSizeOf (CExpr (CBinOp _ e1 e2 _ )) = 1+ cSizeOf e1 + cSizeOf e2 +cSizeOf e = 0+++data CBinOp = CAdd | CSub | CMul | CDiv | CMod + | CEq | CNotEq | CLt | CLEq | CGt | CGEq + | CAnd | COr+ | CPow+ | CBitwiseAnd | CBitwiseOr | CBitwiseXor + | CShiftL | CShiftR + deriving (Eq,Ord,Show) + +data CUnOp = CBitwiseNeg+ deriving (Eq,Ord,Show)++data CAtomicOp = CAtomicAdd | CAtomicInc+ deriving (Eq, Ord, Show) ++---------------------------------------------------------------------------+-- SPMDC+---------------------------------------------------------------------------+data SPMDC = CAssign CExpr [CExpr] CExpr -- array or scalar assign+ | CAtomic CAtomicOp CExpr CExpr CExpr + | CDecl CType Name -- Declare but no assign+ | CDeclAssign CType Name CExpr -- declare variable and assign a value + | CFunc Name [CExpr] + | CSync -- CUDA: "__syncthreads()" OpenCL: "barrier(CLK_LOCAL_MEM_FENCE)"+ | CThreadFence+ | CThreadFenceBlock -- these could be taken care of with a simple+ -- application of the CFunc constructor+ -- but since sync,threadfence etc are special+ -- and might need attention during code gen+ -- I give them specific constructors. + | CFor Name CExpr [SPMDC] -- very simple loop for now.+ | CIf CExpr [SPMDC] [SPMDC]+ deriving (Eq,Ord,Show)+ +-- ret_t param list body+data CKernel = CKernel CQualifyer CType Name [(CType,Name)] [SPMDC] + deriving (Eq,Show)+ +----------------------------------------------------------------------------+-- CExpr +newtype CExpr = CExpr (CExprP CExpr)+ deriving (Eq,Ord,Show)++cTypeOf (CExpr e) = cTypeOfP e + +---------------------------------------------------------------------------- +-- DAGs+type NodeID = Integer +newtype CENode = CENode (CExprP NodeID) + deriving Show+ +----------------------------------------------------------------------------+-- Helpers ++cexpr1 exp a = CExpr $ exp a +cexpr2 exp a b = CExpr $ exp a b +cexpr3 exp a b c = CExpr $ exp a b c +cexpr4 exp a b c d = CExpr $ exp a b c d ++cWarpSize = CExpr $ CVar "warpSize" CWord32 +cBlockIdx = cexpr1 CBlockIdx+cThreadIdx = cexpr1 CThreadIdx+cBlockDim = cexpr1 CBlockDim+cGridDim = cexpr1 CGridDim +cVar = cexpr2 CVar +cLiteral = cexpr2 CLiteral +cIndex = cexpr2 CIndex +cCond = cexpr4 CCond +cFuncExpr = cexpr3 CFuncExpr +cBinOp = cexpr4 CBinOp +cUnOp = cexpr3 CUnOp +cCast = cexpr2 CCast ++cAssign = CAssign+cAtomic = CAtomic +cFunc = CFunc +cDecl = CDecl+cSync = CSync+cThreadFence = CThreadFence+cThreadFenceBlock = CThreadFenceBlock+cDeclAssign = CDeclAssign +cIf = CIf +cFor = CFor +--------------------------------------------------------------------------+-- Printing +data PPConfig = PPConfig {ppKernelQ :: String, + ppGlobalQ :: String, + ppLocalQ :: String,+ ppSyncLine :: String} ++printCKernel :: PPConfig -> CKernel -> String +printCKernel ppc kern = runPP (ppCKernel ppc kern ) 0 ++ppCKernel :: PPConfig -> CKernel -> PP () +ppCKernel ppc (CKernel q t nom ins body) = + ppCQual ppc q >> space >> ppCType ppc t >> space >> line nom >> ppCommaSepList ppIns "(" ")" ins >> + begin >> indent >> newline >> + ppSPMDCList ppc body >> unindent >> newline >>+ end + where + ppIns (t,nom) = ppCType ppc t >> space >> line nom+ +----------------------------------------------------------------------------+ppCQual ppc CQualifyerGlobal = line$ ppGlobalQ ppc +ppCQual ppc CQualifyerLocal = line$ ppLocalQ ppc +ppCQual ppc CQualifyerKernel = line$ ppKernelQ ppc +ppCQual ppc CQualifyerExtern = line$ "extern" +ppCQual ppc CQualifyerShared = line$ "__shared__" -- should this be same as local ?+ppCQual ppc (CQualifyerAttrib a) = ppCAttrib ppc a++ppCAttrib ppc (CAttribAligned x) = line$ "__attribute__ ((aligned(" ++ show x ++ ")))" +----------------------------------------------------------------------------+ppCType ppc CVoid = line "void"+ppCType ppc CInt = line "int"+ppCType ppc CInt8 = line "int8_t"+ppCType ppc CInt16 = line "int16_t"+ppCType ppc CInt32 = line "int32_t"+ppCType ppc CInt64 = line "int64_t"+ppCType ppc CFloat = line "float"+ppCType ppc CDouble = line "double" +ppCType ppc CWord8 = line "uint8_t"+ppCType ppc CWord16 = line "uint16_t"+ppCType ppc CWord32 = line "uint32_t"+ppCType ppc CWord64 = line "uint64_t" +ppCType ppc (CPointer t) = ppCType ppc t >> line "*"+ppCType ppc (CQualified q t) = ppCQual ppc q >> space >> ppCType ppc t++-- a hack (whats the correct way to handle C's t[] ?)+-- Breaks down already for a[][], i think.+ppCTypedName ppc CVoid nom = line "void" >> space >> line nom+ppCTypedName ppc CInt nom = line "int" >> space >> line nom+ppCTypedName ppc CFloat nom = line "float" >> space >> line nom+ppCTypedName ppc CDouble nom = line "double" >> space >> line nom +ppCTypedName ppc CWord8 nom = line "uint8_t" >> space >> line nom+ppCTypedName ppc CWord16 nom = line "uint16_t" >> space >> line nom+ppCTypedName ppc CWord32 nom = line "uint32_t" >> space >> line nom+ppCTypedName ppc CWord64 nom = line "uint64_t" >> space >> line nom+ppCTypedName ppc (CPointer t) nom = ppCType ppc t >> line "*" >> line nom+ppCTypedName ppc (CArray [] t) nom = ppCType ppc t >> space >> line nom >> line "[]"+ppCTypedName ppc (CQualified q t) nom = ppCQual ppc q >> space >> ppCTypedName ppc t nom ++----------------------------------------------------------------------------+ppValue (IntVal i) = line$ show i+ppValue (Int8Val i) = line$ show i+ppValue (Int16Val i) = line$ show i+ppValue (Int32Val i) = line$ show i+ppValue (Int64Val i) = line$ show i+ppValue (FloatVal f) = line$ show f +ppValue (DoubleVal d) = line$ show d+ppValue (Word8Val w) = line$ show w +ppValue (Word16Val w) = line$ show w+ppValue (Word32Val w) = line$ show w+ppValue (Word64Val w) = line$ show w ++----------------------------------------------------------------------------+ppBinOp CAdd = line$ "+"+ppBinOp CSub = line$ "-"+ppBinOp CMul = line$ "*"+ppBinOp CDiv = line$ "/"+ppBinOp CMod = line$ "%" +ppBinOp CEq = line$ "=="+ppBinOp CLt = line$ "<" +ppBinOp CLEq = line$ "<="+ppBinOp CGt = line$ ">" +ppBinOp CGEq = line$ ">="+ppBinOp CNotEq = line$ "/=" +ppBinOp CAnd = line$ "&&"+ppBinOp COr = line$ "||" +ppBinOp CBitwiseAnd = line$ "&" +ppBinOp CBitwiseOr = line$ "|" +ppBinOp CBitwiseXor = line$ "^" +ppBinOp CShiftL = line$ "<<" +ppBinOp CShiftR = line$ ">>"+ +ppUnOp CBitwiseNeg = line$ "~" +-- May be incorrect.+--ppUnOp CInt32ToWord32 = line$ "(uint32_t)"+--ppUnOp CWord32ToInt32 = line$ "(int32_t)" ++---------------------------------------------------------------------------+--+---------------------------------------------------------------------------+ppCommaSepList ppElt s e xs = + line s >> + sequence_ (L.intersperse (line ",") (commaSepList' xs)) >> line e+ where + commaSepList' [] = [] + commaSepList' (x:xs) = ppElt x : commaSepList' xs+ +---------------------------------------------------------------------------+--+---------------------------------------------------------------------------+ppSPMDCList ppc xs = sequence_ (map (ppSPMDC ppc) xs) +++ppSPMDC :: PPConfig -> SPMDC -> PP () +ppSPMDC ppc (CAssign e [] expr) =+ ppCExpr ppc e >> + line " = " >> + ppCExpr ppc expr >> + cTermLn+ppSPMDC ppc (CAssign e exprs expr) =+ ppCExpr ppc e >> + ppCommaSepList (ppCExpr ppc) "[" "]" exprs >> + line " = " >> + ppCExpr ppc expr >> + cTermLn+ppSPMDC ppc (CAtomic op res arr e) =+ --ppCExpr ppc res >>+ --line " = " >>+ ppAtomicOp ppc op >>+ wrap "(" ")" (ppCExpr ppc arr >> line ", " >> ppCExpr ppc e ) >>+ cTermLn ++ppSPMDC ppc (CDecl t n) = ppCTypedName ppc t n >> cTermLn+ppSPMDC ppc (CDeclAssign t n e) =+ ppCTypedName ppc t n >>+ line " = " >>+ ppCExpr ppc e >> cTermLn+ppSPMDC ppc (CFunc nom args) =+ line nom >>+ ppCommaSepList (ppCExpr ppc) "(" ")" args >> cTermLn+ppSPMDC ppc CSync = line (ppSyncLine ppc) >> cTermLn +ppSPMDC ppc (CIf e [] []) = return ()+ppSPMDC ppc (CIf e xs []) =+ line "if " >> + wrap "(" ")" (ppCExpr ppc e) >> + begin >> indent >> newline >> + ppSPMDCList ppc xs >> unindent >> end+ppSPMDC ppc (CIf e xs ys) =+ line "if " >> + wrap "(" ")" (ppCExpr ppc e) >> + begin >> indent >> newline >> + ppSPMDCList ppc xs >> unindent >> end >> + line "else " >> begin >> indent >> newline >> + ppSPMDCList ppc ys >> unindent >> end+-- TODO: Clean up here+ppSPMDC ppc (CFor name e s) =+ line "for " >>+ wrap "(" ")" (line ("int " ++ name ++ " = 0;") >>+ line (name ++ " < ") >> (ppCExpr ppc e) >>+ line (";") >> line (name ++ "++")) >>+ begin >> indent >> newline >> + ppSPMDCList ppc s >> unindent >> end+++ppAtomicOp :: PPConfig -> CAtomicOp -> PP ()+ppAtomicOp ppc CAtomicInc = line "atomicInc" ++----------------------------------------------------------------------------+--+ppCExpr :: PPConfig -> CExpr -> PP () +-- Cheat and do CUDA print for now!+ -- should do lookup in PPConfig and figure out how to + -- print these for CUDA/OpenCL+ppCExpr ppc (CExpr (CBlockIdx X)) = line "blockIdx.x" +ppCExpr ppc (CExpr (CBlockIdx Y)) = line "blockIdx.y" +ppCExpr ppc (CExpr (CBlockIdx Z)) = line "blockIdx.z" +ppCExpr ppc (CExpr (CThreadIdx X)) = line "threadIdx.x" +ppCExpr ppc (CExpr (CThreadIdx Y)) = line "threadIdx.y" +ppCExpr ppc (CExpr (CThreadIdx Z)) = line "threadIdx.z" +ppCExpr ppc (CExpr (CBlockDim X)) = line "blockDim.x" +ppCExpr ppc (CExpr (CBlockDim Y)) = line "blockDim.y" +ppCExpr ppc (CExpr (CBlockDim Z)) = line "blockDim.z" +ppCExpr ppc (CExpr (CGridDim X)) = line "gridDim.x" +ppCExpr ppc (CExpr (CGridDim Y)) = line "gridDim.y" +ppCExpr ppc (CExpr (CGridDim Z)) = line "gridDim.z" ++ppCExpr ppc (CExpr (CVar nom _)) = line nom+ppCExpr ppc (CExpr (CLiteral v _)) = ppValue v +ppCExpr ppc (CExpr (CIndex (e,[]) _)) = ppCExpr ppc e +ppCExpr ppc (CExpr (CIndex (e,xs) _)) =+ ppCExpr ppc e >> + ppCommaSepList (ppCExpr ppc) "[" "]" xs+ppCExpr ppc (CExpr (CCond e1 e2 e3 _)) =+ wrap "(" ")" + (ppCExpr ppc e1 >> + line " ? " >> + ppCExpr ppc e2 >> + line " : " >> + ppCExpr ppc e3+ )+ppCExpr ppc (CExpr (CBinOp bop e1 e2 _)) =+ wrap "(" ")"+ (+ ppCExpr ppc e1 >> + ppBinOp bop >> + ppCExpr ppc e2 + ) +ppCExpr ppc (CExpr (CUnOp uop e _)) =+ wrap "(" ")" + (+ ppUnOp uop >> + ppCExpr ppc e + )+ppCExpr ppc (CExpr (CFuncExpr nom args _)) =+ line nom >> + ppCommaSepList (ppCExpr ppc) "(" ")" args+ppCExpr ppc (CExpr (CCast e t)) =+ line "((" >> + ppCType ppc t >> + line ")" >> + ppCExpr ppc e >> + line ")"++---------------------------------------------------------------------------+-- Optimize for complicated indexing expressions+---------------------------------------------------------------------------++-- TODO: #1: Discover all expressions that represent an index into an array+-- #2: Count usages of them+-- #3: For "Complicated" expressions used more than once+-- declare a new name for the index and compute it once. (if not data dependent) +--+-- Possible approach is two passes over the SPMDC structure.+-- The first discovers expressions+-- The in-between create small SPMDC code that declares variables. +-- The second replaces some of them by a variable+--++-- Assign with all expressions an integer +type ExpMap = M.Map CExpr (Int,Int) ++-- Insert, but only if number of occurances and size is right! +insert :: CExpr -> State (Int,ExpMap) () +insert e | cSizeOf e >= 2 =+ do+ (i,m) <- get+ case M.lookup e m of+ (Just (id,count)) ->+ do+ let m' = M.insert e (id,count+1) m+ put (i,m')+ Nothing ->+ do+ let m' = M.insert e (i,1) m+ put (i+1,m')+insert e = return () ++-- Decide if an expression is safe or not to move to+-- function prelude.+-- Simply put it checks for any data dependency.+-- (This code is unused! ) +safeExp :: S.Set Name -> CExpr -> Bool+safeExp s (CExpr (CVar name _)) = S.member name s+safeExp s (CExpr (CIndex (e,es) _)) = safeExp s e && all (safeExp s) es+safeExp s (CExpr (CCond e1 e2 e3 _)) = safeExp s e1 && safeExp s e2 && safeExp s e3+safeExp s (CExpr (CBinOp _ e1 e2 _)) = safeExp s e2 && safeExp s e2+safeExp s (CExpr (CUnOp _ e _)) = safeExp s e+safeExp s (CExpr (CFuncExpr _ es _)) = all (safeExp s) es+safeExp s (CExpr (CCast e _)) = safeExp s e +safeExp _ _ = True + +collectExps :: [SPMDC] -> State (Int,ExpMap) () +collectExps sp = mapM_ process sp+ where+ process (CAssign _ ixs e) =+ do+ mapM_ processE ixs + processE e+ process (CDeclAssign _ _ e) = processE e+ process (CFunc _ es) = mapM_ processE es+ process (CFor _ e sp) =+ do + processE e+ collectExps sp+ process (CIf bexp sp1 sp2) =+ do+ processE bexp+ collectExps sp1+ collectExps sp2 + process a = return () +++ processE (CExpr (CVar _ _)) = return () -- too simple+ processE (CExpr (CBlockIdx d)) = return () + processE (CExpr (CThreadIdx d)) = return ()+ processE (CExpr (CBlockDim d)) = return ()+ processE (CExpr (CGridDim d)) = return ()+ processE (CExpr (CLiteral _ _)) = return ()+ processE e@(CExpr (CIndex (e1,es) _)) =+ do + -- insert e+ processE e1+ mapM_ processE es+ processE e@(CExpr (CCond e1 e2 e3 _)) =+ do+ insert e+ mapM_ processE [e1,e2,e3]+ processE e@(CExpr (CBinOp _ e1 e2 _)) =+ do+ insert e+ processE e1+ processE e2+ processE e@(CExpr (CUnOp _ e1 _)) =+ do+ insert e+ processE e1+ processE e@(CExpr (CFuncExpr _ es _)) =+ do+ insert e+ mapM_ processE es+ processE e@(CExpr (CCast e1 _)) =+ do+ -- refine this step. Only insert if e1 is nonsimple+ insert e+ processE e1+ ++-- REMEMBER TO KEEP IT SIMPLE.+replacePass :: ExpMap -> [SPMDC] -> ([(Int,CExpr)],[SPMDC])+replacePass _ [] = ([],[])+replacePass m (x:xs) = let (decls,x') = process m x+ (rest, xs') = replacePass m xs+ + in (L.nubBy fstEq (decls ++ rest), x':xs')+ where+ fstEq :: (Int,a) -> (Int,a) -> Bool+ fstEq a b = fst a == fst b++ process m (CFor name e sp) = (decls,CFor name e' sp')+ where+ (decls1, e') = processE m e+ (decls2, sp') = replacePass m sp+ decls = L.nubBy fstEq (decls1++decls2) + process m (CAssign name es e) = (decls,CAssign name es' e') + where+ (decls1,es') = processEList m es+ (decls2,e') = processE m e+ decls = L.nubBy fstEq (decls1 ++ decls2)+ process m s = ([],s) ++ processEList m [] = ([],[])+ processEList m (e:es) =+ let (decls1,e') = processE m e+ (decls2,es') = processEList m es+ in (L.nubBy fstEq (decls1 ++ decls2),e':es')+++ processE m e@(CExpr (CIndex (e1,es) t)) =+ case M.lookup e m of+ Nothing ->+ let (d1,es') = processEList m es+ in (L.nubBy fstEq d1, CExpr (CIndex (e1,es') t))+ + (Just _) -> error "Just in CIndex case"++ processE m e@(CExpr (CCond e1 e2 e3 t)) =+ case M.lookup e m of+ Nothing ->+ let + (d1,e1') = processE m e1+ (d2,e2') = processE m e2+ (d3,e3') = processE m e3+ in (L.nubBy fstEq (d1++d2++d3), CExpr (CCond e1' e2' e3' t))+ Just (id,1) ->+ let + (d1,e1') = processE m e1+ (d2,e2') = processE m e2+ (d3,e3') = processE m e3+ in (L.nubBy fstEq (d1++d2++d3), CExpr (CCond e1' e2' e3' t))+ Just (id,n) -> -- error "SERIOUS FLAW. FIX THIS"+ ([(id,e)],CExpr (CVar ("t" ++ show id) (cTypeOf e)))+ + processE m e@(CExpr (CBinOp op e1 e2 t)) =+ case M.lookup e m of+ Nothing -> + let (d1,e1') = processE m e1+ (d2,e2') = processE m e2+ in (L.nubBy fstEq (d1++d2), CExpr (CBinOp op e1' e2' t))+ + (Just (id,1)) -> + let (d1,e1') = processE m e1+ (d2,e2') = processE m e2+ in (L.nubBy fstEq (d1++d2), CExpr (CBinOp op e1' e2' t))+ + (Just (id,n)) -> + ([(id,e)],CExpr (CVar ("t" ++ show id) (cTypeOf e)))+ + processE m e@(CExpr (CCast e1 t)) = (id,CExpr (CCast e1' t))+ where+ (id,e1') = processE m e1 ++ processE m e =+ case M.lookup e m of+ Nothing -> ([],e)+ (Just (id,1)) -> ([],e)+ (Just (id,n)) -> ([(id,e)],CExpr (CVar ("t" ++ show id) (cTypeOf e)))+ ++++declsToSPMDC :: [(Int,CExpr)] -> [SPMDC]+declsToSPMDC decls = map process decls+ where+ process (i,e) = CDeclAssign (cTypeOf e) ("t" ++ show i) e
+ Obsidian/DimSpec.hs view
@@ -0,0 +1,7 @@++{- Joel Svensson 2012 -} ++module Obsidian.DimSpec (DimSpec (..)) where ++data DimSpec = X | Y | Z+ deriving (Eq,Ord,Show)
+ Obsidian/Exp.hs view
@@ -0,0 +1,816 @@+{-# LANGUAGE GADTs, + TypeFamilies, + FlexibleContexts,+ FlexibleInstances, + UndecidableInstances,+ RankNTypes #-} ++{- Joel Svensson 2012 -} ++module Obsidian.Exp + (module Obsidian.Exp,+ module Obsidian.DimSpec) where ++++import Data.List+import Data.Word+import Data.Int+import Data.Bits++import qualified Foreign.Storable as Storable++import Obsidian.DimSpec++---------------------------------------------------------------------------+-- Obsidian imports+import Obsidian.Types+import Obsidian.Globs++import Obsidian.CodeGen.SPMDC++---------------------------------------------------------------------------+-- some synonyms+type Data a = Exp a +++type EInt = Exp Int +type EWord = Exp Word++type EInt8 = Exp Int8+type EInt16 = Exp Int16+type EInt32 = Exp Int32+type EInt64 = Exp Int64++type EWord8 = Exp Word8 +type EWord16 = Exp Word16 +type EWord32 = Exp Word32 +type EWord64 = Exp Word64 ++type EFloat = Exp Float +type EDouble = Exp Double +type EBool = Exp Bool ++++---------------------------------------------------------------------------+-- Class Scalar. All the things we can handle code generation for ++class (Eq a, ExpToCExp a, Show a) => Scalar a where + sizeOf :: Exp a -> Int -- + typeOf :: Exp a -> Type -- Good enough for me ... +++instance Scalar Bool where + sizeOf _ = Storable.sizeOf (undefined :: Int)+ typeOf _ = Bool ++instance Scalar Int where + sizeOf _ = Storable.sizeOf (undefined :: Int)+ typeOf _ = Int++instance Scalar Int8 where + sizeOf _ = 1+ typeOf _ = Int8++instance Scalar Int16 where + sizeOf _ = 2+ typeOf _ = Int16++instance Scalar Int32 where + sizeOf _ = 4+ typeOf _ = Int32++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++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+ +instance Scalar Word64 where + sizeOf _ = 8 + typeOf _ = Word64++++---------------------------------------------------------------------------+-- Expressions +data Exp a where+ Literal :: Scalar a + => a + -> Exp a + + {- + Add more specific constructors for block,thread variables+ (these concepts excist in both OpenCL and CUDA + but are accessed differently so it could be a good + idea to add them as constructors here. These + can be translated into the CUDA/OpenCL specific + concept later in the codegeneration + -}+ WarpSize :: Exp Word32+ + BlockDim :: DimSpec -> Exp Word32+ + BlockIdx :: DimSpec + -> Exp Word32+ ThreadIdx :: DimSpec+ -> Exp Word32+ + Index :: Scalar a => + (Name,[Exp Word32]) + -> Exp a + + If :: Scalar a + => Exp Bool+ -> Exp a + -> Exp a + -> Exp a + + BinOp :: (Scalar a,+ Scalar b, + Scalar c) + => Op ((a,b) -> c) + -> Exp a + -> Exp b + -> Exp c + + UnOp :: (Scalar a, + Scalar b)+ => Op (a -> b) + -> Exp a + -> Exp b + ++ +---------------------------------------------------------------------------+-- Operations +-- TODO: needs conversion operations.. (Int -> Word) etc. +data Op a where + Add :: Num a => Op ((a,a) -> a) + Sub :: Num a => Op ((a,a) -> a) + Mul :: Num a => Op ((a,a) -> a) + Div :: Num a => Op ((a,a) -> a) + -- If :: Op ((Bool,a,a) -> a) + + Mod :: Integral a => Op ((a,a) -> a)+ + -- Trig+ Sin :: Floating a => Op (a -> a) + Cos :: Floating a => Op (a -> a)+ + -- Comparisons+ Eq :: Ord a => Op ((a,a) -> Bool)+ NotEq :: Ord a => Op ((a,a) -> Bool) + Lt :: Ord a => Op ((a,a) -> Bool) + LEq :: Ord a => Op ((a,a) -> Bool) + Gt :: Ord a => Op ((a,a) -> Bool) + GEq :: Ord a => Op ((a,a) -> Bool) + + -- Boolean + And :: Op ((Bool,Bool) -> Bool) + Or :: Op ((Bool,Bool) -> Bool)+ + -- Bitwise + BitwiseAnd :: Bits a => Op ((a,a) -> a) + BitwiseOr :: Bits a => Op ((a,a) -> a)+ BitwiseXor :: Bits a => Op ((a,a) -> a) + BitwiseNeg :: Bits a => Op (a -> a)++ -- I DO NOT EVEN KNOW WHAT THIS MEANS: work around it! + ShiftL :: forall a b. (Num b, Bits a) => Op ((a, b) -> a) + ShiftR :: forall a b .(Num b, Bits a) => Op ((a, b) -> a) + + -- built-ins+ Min :: Ord a => Op ((a,a) -> a) + Max :: Ord a => Op ((a,a) -> a) ++ -- Floating (different CUDA functions for float and double, issue maybe?) + Exp :: Floating a => Op (a -> a) -- "expf" + Sqrt :: Floating a => Op (a -> a) -- "sqrtf" + --RSqrt :: Floating a => Op (a -> a) -- "rsqrtf"+ Log :: Floating a => Op (a -> a) -- "logf"+ Log2 :: Floating a => Op (a -> a) -- "log2f"+ Log10 :: Floating a => Op (a -> a) -- "log10f"+ Pow :: Floating a => Op ((a, a) -> a) -- "powf"+ -- Floating Trig+ Tan :: Floating a => Op (a -> a) -- "tanf"+ ASin :: Floating a => Op (a -> a) -- "asinf"+ ATan :: Floating a => Op (a -> a) -- "atanf"+ ACos :: Floating a => Op (a -> a) -- "acosf"+ SinH :: Floating a => Op (a -> a) -- "sinhf"+ TanH :: Floating a => Op (a -> a) -- "tanhf"+ CosH :: Floating a => Op (a -> a) -- "coshf"+ ASinH :: Floating a => Op (a -> a) -- "asinhf" + ATanH :: Floating a => Op (a -> a) -- "atanhf"+ ACosH :: Floating a => Op (a -> a) -- "acoshf"+ -- There is no "div" in "Num" but it's already defined above. + FDiv :: Floating a => Op ((a, a) -> a) -- "acoshf"++ Int32ToWord32 :: Op (Int32 -> Word32)+ Word32ToInt32 :: Op (Word32 -> Int32) +++---------------------------------------------------------------------------+-- helpers ++variable name = Index (name,[])+index name ix = Index (name,[ix])++warpSize :: Exp Word32+warpSize = WarpSize++---------------------------------------------------------------------------+-- Collect array names++collectArrays :: Scalar a => Exp a -> [Name]+collectArrays (Literal _) = []+collectArrays (ThreadIdx _) = []+collectArrays (BlockIdx _) = [] +collectArrays (Index (name,[])) = []+collectArrays (Index (name,_)) = [name]+collectArrays (BinOp _ e1 e2) = collectArrays e1 ++ collectArrays e2+collectArrays (UnOp _ e) = collectArrays e+collectArrays (If b e1 e2) = collectArrays b ++ + collectArrays e1 ++ + collectArrays e2+-- collectArrays a = error $ show a++collectArrayIndexPairs :: Scalar a => Exp a -> [(Name,Exp Word32)]+collectArrayIndexPairs (Literal _) = []+collectArrayIndexPairs (Index (name,[])) = []+collectArrayIndexPairs (Index (name,[ix])) = [(name,ix)]+collectArrayIndexPairs (BinOp _ e1 e2) = collectArrayIndexPairs e1 ++ collectArrayIndexPairs e2+collectArrayIndexPairs (UnOp _ e) = collectArrayIndexPairs e+collectArrayIndexPairs (If b e1 e2) = collectArrayIndexPairs b ++ + collectArrayIndexPairs e1 ++ + collectArrayIndexPairs e2+++---------------------------------------------------------------------------+-- Typecasts+---------------------------------------------------------------------------+int32ToWord32 = UnOp Int32ToWord32+word32ToInt32 = UnOp Word32ToInt32++---------------------------------------------------------------------------+-- +instance Scalar a => Show (Exp a) where + show = printExp ++-- Look this over. Do I really need a types expression data type ?+-- (No real need for a Exp GADT I think. Go back to keeping it simple!) +instance (Eq a, Scalar a) => Eq (Exp a) where+ (==) a b = -- error $ "equality test between exps: " ++ show a ++ " " ++ show b --+ expToCExp a == expToCExp b+ -- Maybe not efficient! But simple.++ +instance (Scalar a, Ord a) => Ord (Exp a) where + min a b = BinOp Min a b+ max a b = BinOp Max a b++---------------------------------------------------------------------------+-- INT Instances+---------------------------------------------------------------------------+instance Num (Exp Int) where + (+) a (Literal 0) = a+ (+) (Literal 0) a = a+ (+) (Literal a) (Literal b) = Literal (a+b)+ -- Added 2 Oct 2012+ (+) (BinOp Sub b (Literal a)) (Literal c) | a == c = b + (+) (Literal b) (BinOp Sub a (Literal c)) | b == c = a + (+) a b = BinOp Add a b + + (-) a (Literal 0) = a + (-) (Literal a) (Literal b) = Literal (a - b) + (-) a b = BinOp Sub a b + + (*) a (Literal 1) = a + (*) (Literal 1) a = a+ (*) _ (Literal 0) = Literal 0+ (*) (Literal 0) _ = Literal 0+ (*) (Literal a) (Literal b) = Literal (a*b) + (*) a b = BinOp Mul a b + + signum = error "signum: not implemented for Exp Int" + abs = error "abs: not implemented for Exp Int" + fromInteger a = Literal (fromInteger a) + +-- Added new cases for literal 0 (2012/09/25)+instance Bits (Exp Int) where + (.&.) x (Literal 0) = Literal 0+ (.&.) (Literal 0) x = Literal 0 + (.&.) (Literal a) (Literal b) = Literal (a .&. b) + (.&.) a b = BinOp BitwiseAnd a b+ (.|.) (Literal a) (Literal b) = Literal (a .|. b)+ (.|.) a b = BinOp BitwiseOr a b+ xor (Literal a) (Literal b) = Literal (a `xor` b) + xor a b = BinOp BitwiseXor a b + + --TODO: See that this is not breaking something (32/64 bit, CUDA/Haskell)+ complement (Literal i) = Literal (complement i)+ + complement a = UnOp BitwiseNeg a+ shiftL a i = BinOp ShiftL a (Literal i)+ shiftR a i = BinOp ShiftR a (Literal i)+ bitSize a = sizeOf a * 8+ isSigned a = True++ bit = error "bit: is undefined for Exp Int"+ testBit = error "testBit: is undefined for Exp Int"+ popCount = error "popCoint: is undefined for Exp Int"+++-- TODO: change undefined to some specific error.+instance Real (Exp Int) where+ toRational = error "toRational: not implemented for Exp Int)" ++instance Enum (Exp Int) where+ toEnum = error "toEnum: not implemented for Exp Int" + fromEnum = error "fromEnum: not implemented for Exp Int"+ +instance Integral (Exp Int) where+ mod (Literal a) (Literal b) = Literal (a `mod` b) + mod a b = BinOp Mod a b+ div _ (Literal 0) = error "Division by zero in expression" + div a b = BinOp Div a b+ quotRem = error "quotRem: not implemented for Exp Int" + toInteger = error "toInteger: not implemented for Exp Int" ++---------------------------------------------------------------------------+-- Int32+---------------------------------------------------------------------------+instance Num (Exp Int32) where + (+) a (Literal 0) = a+ (+) (Literal 0) a = a+ (+) (Literal a) (Literal b) = Literal (a+b)+ -- Added 2 Oct 2012+ (+) (BinOp Sub b (Literal a)) (Literal c) | a == c = b + (+) (Literal b) (BinOp Sub a (Literal c)) | b == c = a + (+) a b = BinOp Add a b + + (-) a (Literal 0) = a + (-) (Literal a) (Literal b) = Literal (a - b) + (-) a b = BinOp Sub a b + + (*) a (Literal 1) = a + (*) (Literal 1) a = a+ (*) _ (Literal 0) = 0+ (*) (Literal 0) _ = 0 + (*) (Literal a) (Literal b) = Literal (a*b) + (*) a b = BinOp Mul a b + + signum = error "signum: not implemented for Exp Int32"+ abs = error "abs: not implemented for Exp Int32" + fromInteger a = Literal (fromInteger a) + +-- Added new cases for literal 0 (2012/09/25)+instance Bits (Exp Int32) where + (.&.) x (Literal 0) = Literal 0+ (.&.) (Literal 0) x = Literal 0 + (.&.) (Literal a) (Literal b) = Literal (a .&. b) + (.&.) a b = BinOp BitwiseAnd a b+ (.|.) (Literal a) (Literal b) = Literal (a .|. b)+ (.|.) a b = BinOp BitwiseOr a b+ xor (Literal a) (Literal b) = Literal (a `xor` b) + xor a b = BinOp BitwiseXor a b + + --TODO: See that this is not breaking something (32/64 bit, CUDA/Haskell)+ complement (Literal i) = Literal (complement i)+ + complement a = UnOp BitwiseNeg a+ shiftL a i = BinOp ShiftL a (Literal i)+ shiftR a i = BinOp ShiftR a (Literal i)+ bitSize a = 32 -- sizeeOf a * 8+ isSigned a = True++ bit = error "bit: is undefined for Exp Int32"+ testBit = error "testBit: is undefined for Exp Int32"+ popCount = error "popCoint: is undefined for Exp Int32"+++-- TODO: change undefined to some specific error.+instance Real (Exp Int32) where+ toRational = error "toRational: not implemented for Exp Int32"++instance Enum (Exp Int32) where+ toEnum = error "toEnum: not implemented for Exp Int32" + fromEnum = error "fromEnum: not implemented for Exp Int32" + +instance Integral (Exp Int32) where+ mod (Literal a) (Literal b) = Literal (a `mod` b) + mod a b = BinOp Mod a b+ div _ (Literal 0) = error "Division by zero in expression" + div a b = BinOp Div a b+ quotRem = error "quotRem: not implemented for Exp Int32" + toInteger = error "toInteger: not implemented for Exp Int32" +++---------------------------------------------------------------------------+-- Word32 Instances+---------------------------------------------------------------------------+instance Num (Exp Word32) where + (+) a (Literal 0) = a+ (+) (Literal 0) a = a+ (+) (Literal a) (Literal b) = Literal (a+b)++ -- Added 15 Jan 2013+ (+) (BinOp Mul (BinOp Div x (Literal a)) (Literal b))+ (BinOp Mod y (Literal c))+ | x == y && a == b && b == c = x + -- This spots the kind of indexing that occurs from + -- converting a bix tix view to and from gix view+ + -- Added 2 oct 2012+ (+) (BinOp Sub b (Literal a)) (Literal c) | a == c = b + (+) (Literal b) (BinOp Sub a (Literal c)) | b == c = a + + (+) a b = BinOp Add a b + + (-) a (Literal 0) = a + (-) (Literal a) (Literal b) = Literal (a - b) + (-) a b = BinOp Sub a b + + (*) a (Literal 1) = a + (*) (Literal 1) a = a+ (*) _ (Literal 0) = Literal 0+ (*) (Literal 0) _ = Literal 0+ (*) (Literal a) (Literal b) = Literal (a*b) + (*) a b = BinOp Mul a b + + signum = error "signum: not implemented for Exp Word32"+ abs = error "abs: not implemented for Exp Word32" + fromInteger a = Literal (fromInteger a) + ++-- adding special shift operators for when both inputs are +-- runtime values (2013-01-08) +(<<*) :: (Scalar b, Scalar a, Bits a, Num b ) => Exp a -> Exp b -> Exp a +(<<*) a b = BinOp ShiftL a b ++(>>*) :: (Scalar b, Scalar a, Bits a, Num b ) => Exp a -> Exp b -> Exp a +(>>*) a b = BinOp ShiftR a b +++ -- Added new cases for literal 0 (2012/09/25)+instance Bits (Exp Word32) where + (.&.) x (Literal 0) = Literal 0+ (.&.) (Literal 0) x = Literal 0 + (.&.) (Literal a) (Literal b) = Literal (a .&. b) + (.&.) a b = BinOp BitwiseAnd a b + (.|.) (Literal a) (Literal b) = Literal (a .|. b) + (.|.) a b = BinOp BitwiseOr a b+ xor (Literal a) (Literal b) = Literal (a `xor` b) + xor a b = BinOp BitwiseXor a b + complement (Literal i) = Literal (complement i) + complement a = UnOp BitwiseNeg a+ + shiftL (Literal j) i = Literal (j `shiftL` i) + shiftL a i = BinOp ShiftL a (Literal i)+ + shiftR (Literal j) i = Literal (j `shiftL` i)+ shiftR a i = BinOp ShiftR a (Literal i)+ bitSize a = 32+ isSigned a = False++ bit = error "bit: is undefined for Exp Word32"+ testBit = error "testBit: is undefined for Exp Word32"+ popCount = error "popCoint: is undefined for Exp Word32"++instance Real (Exp Word32) where + toRational = error "toRational: not implemented for Exp Word32" + ++instance Enum (Exp Word32) where+ toEnum = error "toEnum: not implemented for Exp Word32" + fromEnum = error "fromEnum: not implemented for Exp Word32" ++instance Integral (Exp Word32) where+ mod (Literal a) (Literal b) = Literal (a `mod` b) + mod a b = BinOp Mod a b+ div _ (Literal 0) = error "Division by zero in expression" + div a b = BinOp Div a b+ quotRem = error "quotRem: not implemented for Exp Word32" + toInteger = error "toInteger: not implemented for Exp Word32"+ +instance Num (Exp Float) where+ (+) a (Literal 0) = a+ (+) (Literal 0) a = a+ (+) (Literal a) (Literal b) = Literal (a + b)+ (+) a b = BinOp Add a b+ + (-) a (Literal 0) = a+ (-) (Literal a) (Literal b) = Literal (a - b)+ (-) a b = BinOp Sub a b+ + (*) a (Literal 1) = a+ (*) (Literal 1) a = a+ (*) _ (Literal 0) = Literal 0+ (*) (Literal 0) _ = Literal 0+ (*) (Literal a) (Literal b) = Literal (a * b)+ (*) a b = BinOp Mul a b+ + signum = undefined+ abs = undefined+ fromInteger a = Literal (fromInteger a)++instance Fractional (Exp Float) where+ (/) a b = BinOp FDiv a b+ recip a = (Literal 1) / a+ fromRational a = Literal (fromRational a)++instance Floating (Exp Float) where+ pi = Literal pi+ exp a = UnOp Exp a+ sqrt a = UnOp Sqrt a+ log a = UnOp Log a+ (**) a b = BinOp Pow a b+ + -- log_b(x) = log_e(x) / log_e(b)+ logBase (Literal 2) b = UnOp Log2 b+ logBase (Literal 10) b = UnOp Log10 b+ logBase a b = (UnOp Log b) / (UnOp Log a)+ + sin (Literal 0) = Literal 0+ sin a = UnOp Sin a+ tan (Literal 0) = Literal 0+ tan a = UnOp Tan a+ cos (Literal 0) = Literal 1+ cos a = UnOp Cos a+ + asin (Literal 0) = Literal 0+ asin a = UnOp ASin a+ atan (Literal 0) = Literal 0+ atan a = UnOp ATan a+ acos (Literal 1) = Literal 0+ acos a = UnOp ACos a+ + sinh (Literal 0) = Literal 0+ sinh a = UnOp Sin a+ tanh (Literal 0) = Literal 0+ tanh a = UnOp Tan a+ cosh (Literal 0) = Literal 1+ cosh a = UnOp Cos a+ + asinh a = UnOp ASinH a+ atanh a = UnOp ATanH a+ acosh a = UnOp ACosH a+ + -- Y-Less's comment+ -- Don't second guess the CUDA compiler (or, more accurately, assume that+ -- other compilers have this).+ --(/) (Literal 1) (UnOp Sqrt b) = UnOp RSqrt b -- Optimisation.++ + + +---------------------------------------------------------------------------+ +infix 4 ==*, /=*, <*, >*, >=*, <=* + +(==*) (Literal a) (Literal b) = Literal (a == b) +(==*) a b = BinOp Eq a b+(/=*) a b = BinOp NotEq a b +(<*) (Literal a) (Literal b) = Literal (a < b) +(<*) a b = BinOp Lt a b+(<=*) (Literal a) (Literal b) = Literal (a <= b) +(<=*) a b = BinOp LEq a b+(>*) a b = BinOp Gt a b+(>=*) a b = BinOp GEq a b++infixr 3 &&*+infixr 2 ||* +(&&*) a b = BinOp And a b +(||*) a b = BinOp Or a b ++---------------------------------------------------------------------------+-- Choice class+---------------------------------------------------------------------------+class Choice a where + ifThenElse :: Exp Bool -> a -> a -> a ++instance Scalar a => Choice (Exp a) where + ifThenElse (Literal False) e1 e2 = e2+ ifThenElse (Literal True) e1 e2 = e1+ ifThenElse b e1 e2 = If b e1 e2+ +instance (Choice a, Choice b) => Choice (a,b) where+ ifThenElse b (e1,e1') (e2,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 (Literal a) = show a +printExp (Index (name,[])) = name+printExp (Index (name,es)) = + name ++ "[" ++ ((concat . intersperse "," . map printExp) es) ++ "]"+printExp (BinOp op e1 e2) = "(" ++ printOp op ++ " " ++ printExp e1 ++ " " ++ printExp e2 ++ " )"+printExp (UnOp op e) = "(" ++ printOp op ++ " " ++ printExp e ++ " )"+printExp (If b e1 e2) = "(" ++ printExp b ++ " ? " ++ printExp e1 ++ " : " ++ printExp e2 ++ ")"+++printOp :: Op a -> String+printOp Add = " + " +printOp Sub = " - " +printOp Mul = " * "+printOp Div = " / "+printOp Mod = " % "++-- printOp If = " if "++printOp Eq = " == "+printOp NotEq = " /= " +printOp Lt = " < " +printOp LEq = " <= " +printOp Gt = " > "+printOp GEq = " >= " ++printOp And = " && "+printOp Or = " || " ++printOp Min = " Min "+printOp Max = " Max " ++printOp Sin = " Sin " +printOp Cos = " Cos "++printOp BitwiseAnd = " & "+printOp BitwiseOr = " | " +printOp BitwiseXor = " ^ " +printOp BitwiseNeg = " ~ " ++---------------------------------------------------------------------------+-- Turn expressions into backend-expressions+---------------------------------------------------------------------------+class ExpToCExp a where + expToCExp :: Exp a -> CExpr +++instance ExpToCExp Bool where + expToCExp (Literal True) = cLiteral (IntVal 1) CInt + expToCExp (Literal False) = cLiteral (IntVal 0) CInt+ expToCExp a = expToCExpGeneral a ++instance ExpToCExp Int where + expToCExp (Literal a) = cLiteral (IntVal a) CInt+ expToCExp a = expToCExpGeneral a ++instance ExpToCExp Int8 where + expToCExp (Literal a) = cLiteral (Int8Val a) CInt8+ expToCExp a = expToCExpGeneral a ++instance ExpToCExp Int16 where + expToCExp (Literal a) = cLiteral (Int16Val a) CInt16+ expToCExp a = expToCExpGeneral a ++instance ExpToCExp Int32 where + expToCExp (Literal a) = cLiteral (Int32Val a) CInt32+ expToCExp a = expToCExpGeneral a ++instance ExpToCExp Int64 where + expToCExp (Literal a) = cLiteral (Int64Val a) CInt64+ expToCExp a = expToCExpGeneral a ++instance ExpToCExp Float where + expToCExp (Literal a) = cLiteral (FloatVal a) CFloat+ expToCExp a = expToCExpGeneral a ++instance ExpToCExp Double where + expToCExp (Literal a) = cLiteral (DoubleVal a) CDouble+ expToCExp a = expToCExpGeneral a ++instance ExpToCExp Word where + expToCExp (Literal a) = cLiteral (WordVal a) CWord+ expToCExp a = expToCExpGeneral a ++instance ExpToCExp Word8 where + expToCExp (Literal a) = cLiteral (Word8Val a) CWord8+ expToCExp a = expToCExpGeneral a ++instance ExpToCExp Word16 where + expToCExp (Literal a) = cLiteral (Word16Val a) CWord16+ expToCExp a = expToCExpGeneral a ++instance ExpToCExp Word32 where + expToCExp (Literal a) = cLiteral (Word32Val a) CWord32+ expToCExp a = expToCExpGeneral a ++instance ExpToCExp Word64 where + expToCExp (Literal a) = cLiteral (Word64Val a) CWord64+ expToCExp a = expToCExpGeneral a ++ +expToCExpGeneral :: ExpToCExp a => Exp a -> CExpr+expToCExpGeneral WarpSize = cWarpSize +expToCExpGeneral (BlockIdx d) = cBlockIdx d+expToCExpGeneral (BlockDim d) = cBlockDim d +expToCExpGeneral (ThreadIdx d) = cThreadIdx d++expToCExpGeneral e@(Index (name,[])) = cVar name (typeToCType (typeOf e))+expToCExpGeneral e@(Index (name,xs)) = cIndex (cVar name (CPointer (typeToCType (typeOf e))),map expToCExp xs) (typeToCType (typeOf e)) +expToCExpGeneral e@(If b e1 e2) = cCond (expToCExp b) (expToCExp e1) (expToCExp e2) (typeToCType (typeOf e))++expToCExpGeneral (UnOp Word32ToInt32 e) = cCast (expToCExp e) CInt32+expToCExpGeneral (UnOp Int32ToWord32 e) = cCast (expToCExp e) CWord32++expToCExpGeneral e@(BinOp Min e1 e2) = cFuncExpr "min" [expToCExp e1, expToCExp e2] (typeToCType (typeOf e)) +expToCExpGeneral e@(BinOp Max e1 e2) = cFuncExpr "max" [expToCExp e1, expToCExp e2] (typeToCType (typeOf e)) +expToCExpGeneral e@(BinOp op e1 e2) = cBinOp (binOpToCBinOp op) (expToCExp e1) (expToCExp e2) (typeToCType (typeOf e)) ++expToCExpGeneral (UnOp Exp e) = cFuncExpr "exp" [expToCExp e] (typeToCType (typeOf e))+expToCExpGeneral (UnOp Sqrt e) = cFuncExpr "sqrt" [expToCExp e] (typeToCType (typeOf e))+expToCExpGeneral (UnOp Log e) = cFuncExpr "log" [expToCExp e] (typeToCType (typeOf e))+expToCExpGeneral (UnOp Log2 e) = cFuncExpr "log2" [expToCExp e] (typeToCType (typeOf e))+expToCExpGeneral (UnOp Log10 e) = cFuncExpr "log10" [expToCExp e] (typeToCType (typeOf e))+ +-- Floating trig+expToCExpGeneral (UnOp Sin e) = cFuncExpr "sin" [expToCExp e] (typeToCType (typeOf e))+expToCExpGeneral (UnOp Cos e) = cFuncExpr "cos" [expToCExp e] (typeToCType (typeOf e))+expToCExpGeneral (UnOp Tan e) = cFuncExpr "tan" [expToCExp e] (typeToCType (typeOf e))+expToCExpGeneral (UnOp ASin e) = cFuncExpr "asin" [expToCExp e] (typeToCType (typeOf e))+expToCExpGeneral (UnOp ACos e) = cFuncExpr "acos" [expToCExp e] (typeToCType (typeOf e))+expToCExpGeneral (UnOp ATan e) = cFuncExpr "atan" [expToCExp e] (typeToCType (typeOf e))+expToCExpGeneral (UnOp SinH e) = cFuncExpr "sinh" [expToCExp e] (typeToCType (typeOf e))+expToCExpGeneral (UnOp CosH e) = cFuncExpr "cosh" [expToCExp e] (typeToCType (typeOf e))+expToCExpGeneral (UnOp TanH e) = cFuncExpr "tanh" [expToCExp e] (typeToCType (typeOf e))+expToCExpGeneral (UnOp ASinH e) = cFuncExpr "asinh" [expToCExp e] (typeToCType (typeOf e))+expToCExpGeneral (UnOp ACosH e) = cFuncExpr "acosh" [expToCExp e] (typeToCType (typeOf e))+expToCExpGeneral (UnOp ATanH e) = cFuncExpr "atanh" [expToCExp e] (typeToCType (typeOf e))+ +expToCExpGeneral e@(UnOp op e1) = cUnOp (unOpToCUnOp op) (expToCExp e1) (typeToCType (typeOf e))+ ++typeToCType Bool = CInt +typeToCType Int = CInt+typeToCType Int8 = CInt8+typeToCType Int16 = CInt16+typeToCType Int32 = CInt32+typeToCType Int64 = CInt64+typeToCType Float = CFloat+typeToCType Double = CDouble+typeToCType Word8 = CWord8+typeToCType Word16 = CWord16+typeToCType Word32 = CWord32+typeToCType Word64 = CWord64+typeToCType (Pointer t) = CPointer (typeToCType t)+typeToCType (Global t) = CQualified CQualifyerGlobal (typeToCType t) +typeToCType (Local t) = CQualified CQualifyerLocal (typeToCType t) ++-- maybe unnecessary+binOpToCBinOp Add = CAdd+binOpToCBinOp Sub = CSub+binOpToCBinOp Mul = CMul+binOpToCBinOp Div = CDiv +binOpToCBinOp FDiv = CDiv -- (???) +binOpToCBinOp Mod = CMod++binOpToCBinOp Eq = CEq +binOpToCBinOp NotEq = CNotEq+binOpToCBinOp Lt = CLt +binOpToCBinOp LEq = CLEq+binOpToCBinOp Gt = CGt +binOpToCBinOp GEq = CGEq ++binOpToCBinOp And = CAnd +binOpToCBinOp Or = COr ++binOpToCBinOp Pow = CPow++binOpToCBinOp BitwiseAnd = CBitwiseAnd+binOpToCBinOp BitwiseOr = CBitwiseOr+binOpToCBinOp BitwiseXor = CBitwiseXor+binOpToCBinOp ShiftL = CShiftL +binOpToCBinOp ShiftR = CShiftR+-- notice min and max is not here ! ++unOpToCUnOp BitwiseNeg = CBitwiseNeg+
+ Obsidian/Force.hs view
@@ -0,0 +1,68 @@++{-# LANGUAGE MultiParamTypeClasses,+ FlexibleInstances,+ ScopedTypeVariables,+ TypeFamilies,+ GADTs #-}++{- Joel Svensson 2012, 2013 ++ Notes:+ 2013-04-10: Looking at force and threads+ 2013-01-27: globArrays nolonger exist+ 2013-01-02: Added simple forceG for globArrays+ 2012-12-10: Edited ++-}++-- write_ should be internal use only+module Obsidian.Force (write,+ force,+ forceG+ ) where+++import Obsidian.Program+import Obsidian.Exp+import Obsidian.Array+import Obsidian.Types+import Obsidian.Globs+import Obsidian.Memory+import Obsidian.Names++import Data.Word+---------------------------------------------------------------------------+-- Force local (requires static lengths!) +---------------------------------------------------------------------------++write :: forall p a. (Array p, Pushable p, MemoryOps a) => p Word32 a -> BProgram (Pull Word32 a)+write arr = do + -- snames <- names "arr" (undefined :: a)++ -- Here I know that this pattern match will succeed+ let n = len arr+ + -- allocateArray snames (undefined :: a) n++ let (Push m p) = push Block arr++ snames <- p (assignArrayN n) + + return $ pullFromS snames n++ +force :: forall p a. (Array p, Pushable p, MemoryOps a) => p Word32 a -> BProgram (Pull Word32 a)+force arr = do+ rval <- write arr+ Sync+ return rval+++forceG :: forall l a. GlobalMemoryOps a+ => Push Grid l a+ -> GProgram () +forceG (Push _ p) =+ do+ output <- outputs (undefined :: a) + p (\a e -> do {assignOut output a e; return (Single (Var,""))}) + return ()
+ Obsidian/Globs.hs view
@@ -0,0 +1,16 @@++{- Joel Svensson 2012 -} ++module Obsidian.Globs where+++import Data.Word++---------------------------------------------------------------------------+-- Aliases++type Name = String++++type NumThreads = Word32
+ Obsidian/Library.hs view
@@ -0,0 +1,606 @@+{- Joel Svensson 2012, 2013 + Mary Sheeran 2012++ Notes:+ 2013-01-24: GlobPull nolonger exists+ GlobPush is Push Grid+ + 2013-01-08: Renamed GlobArray to GlobPush + 2013-01-02: Added toGlobArray and toGlobArrayN+ 2012-12-10: Refactoring+ (adherence to new Array types and program types) +-}++{-# LANGUAGE FlexibleInstances,+ TypeSynonymInstances,+ ScopedTypeVariables,+ TypeFamilies,+ GADTs #-}++module Obsidian.Library where ++import Obsidian.Array +import Obsidian.Exp +import Obsidian.Program+import Obsidian.Types++import Obsidian.Force++-- needed for threadsPerBlock analysis +import qualified Obsidian.CodeGen.Program as P ++import Data.Bits +import Data.Word++import Prelude hiding (splitAt,zipWith,replicate)+++++---------------------------------------------------------------------------+-- Reverse an array by indexing in it backwards+---------------------------------------------------------------------------+ +rev :: ASize l => Pull l a -> Pull l a +rev arr = mkPullArray n (\ix -> arr ! ((sizeConv m) - ix)) + where m = n-1+ n = len arr+ +---------------------------------------------------------------------------+-- splitAt (name clashes with Prelude.splitAt)+---------------------------------------------------------------------------+splitAt :: (Integral i, ASize l) => i -> Pull l a -> (Pull l a, Pull l a) +splitAt n arr = (mkPullArray m (\ix -> arr ! ix), + mkPullArray (len arr - m) (\ix -> arr ! (ix + pos)))+ where pos = fromIntegral n+ m = fromIntegral n+++halve arr = splitAt n2 arr+ where + n = len arr+ n2 = n `div` 2++---------------------------------------------------------------------------+-- replicate +---------------------------------------------------------------------------+replicate n a = mkPullArray n (\ix -> a)++singleton a = replicate 1 a ++---------------------------------------------------------------------------+-- last+---------------------------------------------------------------------------++last arr = arr ! fromIntegral ( len arr - 1)+++---------------------------------------------------------------------------+-- Take and Drop (what about strange sizes ? fix) +---------------------------------------------------------------------------+take :: ASize l => l -> Pull l a -> Pull l a+take n arr = resize n arr++drop :: ASize l => l -> Pull l a -> Pull l a+drop n arr = resize (len arr - n) $ ixMap (\ix -> ix + sizeConv n) arr++---------------------------------------------------------------------------+-- Shift arrays+---------------------------------------------------------------------------+shiftRight :: (ASize l, Choice a) => Word32 -> a -> Pull l a -> Pull l a+shiftRight dist elt arr = resize (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 = mkPullArray (len arr)+ $ \ix -> (arr `conc` replicate (fromIntegral dist) elt)+ ! (ix + fromIntegral dist) + +---------------------------------------------------------------------------+-- elements at even indices to fst output, odd to snd.+---------------------------------------------------------------------------+evenOdds :: ASize l => Pull l a -> (Pull l a, Pull l a)+evenOdds arr = (mkPullArray (n-n2) (\ix -> arr ! (2*ix)) ,+ mkPullArray n2 (\ix -> arr ! (2*ix + 1)))+ where+ n = len arr+ n2 = div n 2+ +evens, odds :: ASize l => Pull l a -> Pull l a+evens = fst . evenOdds+odds = snd . evenOdds++-- 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+---------------------------------------------------------------------------+conc :: (ASize l, Choice a) => Pull l a -> Pull l a -> Pull l a +conc a1 a2 = mkPullArray (n1+n2)+ $ \ix -> ifThenElse (ix <* (fromIntegral n1)) + (a1 ! ix) + (a2 ! (ix - (fromIntegral n1)))+ where + n1 = len a1+ n2 = len a2 ++ +---------------------------------------------------------------------------+-- zipp unzipp+---------------------------------------------------------------------------+unzipp :: ASize l => Pull l (a,b) -> (Pull l a, Pull l b) +unzipp arr = (mkPullArray (len arr) (\ix -> fst (arr ! ix)) ,+ mkPullArray (len arr) (\ix -> snd (arr ! ix)) )+ +zipp :: ASize l => (Pull l a, Pull l b) -> Pull l (a, b) +zipp (arr1,arr2) = Pull (min (len arr1) (len arr2))+ $ \ix -> (arr1 ! ix, arr2 ! ix) ++unzipp3 :: ASize l => Pull l (a,b,c) + -> (Pull l a, Pull l b, Pull l c) +unzipp3 arr = (fmap (\(x,_,_) -> x) arr,+ fmap (\(_,y,_) -> y) arr,+ fmap (\(_,_,z) -> z) arr) +++zipp3 :: ASize l => (Pull l a, Pull l b, Pull l c) + -> Pull l (a,b,c) +zipp3 (arr1,arr2,arr3) = + mkPullArray (minimum [len arr1, len arr2, len arr3])+ (\ix -> (arr1 ! ix, arr2 ! ix, arr3 ! ix))+ +++zipWith :: ASize l => (a -> b -> c) -> Pull l a -> Pull l b -> Pull l c+zipWith op a1 a2 = + mkPullArray (min (len a1) (len a2))+ (\ix -> (a1 ! ix) `op` (a2 ! ix))+ +---------------------------------------------------------------------------+-- pair +---------------------------------------------------------------------------+pair :: ASize l => Pull l a -> Pull l (a,a)+pair (Pull n ixf) = + mkPullArray n' (\ix -> (ixf (ix*2),ixf (ix*2+1))) + where + n' = n `div` 2 ++++unpair :: ASize l => Choice a => Pull l (a,a) -> Pull l a+unpair arr = + let n = len arr+ in mkPullArray (2*n) (\ix -> ifThenElse ((mod ix 2) ==* 0) + (fst (arr ! (ix `shiftR` 1)))+ (snd (arr ! (ix `shiftR` 1)))) +++---------------------------------------------------------------------------+-- twoK (untested for proper functionality) +---------------------------------------------------------------------------++binSplit = twoK++-- 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+twoK n f = (\arr -> + let arr' = mkPullArray lt (\i -> (f (mkPullArray m (\j -> (arr ! (g i j)))) ! (h i))) + m = (len arr `shiftR` n) --pow of two + g i j = i .&. (fromIntegral (complement (m-1))) .|. j + h i = i .&. (fromIntegral (nl2-1)) -- optimize ++ nl2 = (len (f (mkPullArray m (\j -> arr ! variable "X"))))+ lt = nl2 `shiftL` n + in arr')+ ++---------------------------------------------------------------------------+-- *** PUSHY LIBRARY *** ---+---------------------------------------------------------------------------++---------------------------------------------------------------------------+-- Concatenate on Push arrays +---------------------------------------------------------------------------+concP :: (Array arr1, Array arr2, ASize l,+ Pushable arr1, Pushable arr2)+ => PT t -> arr1 l a -> arr2 l a -> Push t l a +concP pt arr1 arr2 = + mkPushArray (n1 + n2)+ $ \wf ->+ do+ parr1 wf+ parr2 $ \a i -> wf a (sizeConv n1 + i)+ + where+ n1 = len arr1 + n2 = len arr2 + (Push _ parr1) = push pt arr1+ (Push _ parr2) = push pt arr2+++-- More general versions of this can be imagined +mergeL :: (EWord32 -> a -> a -> a) -> Pull Word32 a -> Pull Word32 a -> Push Block Word32 a+mergeL _ arr1 arr2 | len arr1 /= len arr2 = error "incorrect lengths" +mergeL f arr1 arr2 =+ Push (len arr1) $ \wf ->+ do+ ForAll (sizeConv (len arr1)) $+ \ tid -> wf (f tid (arr1 ! tid) (arr2 ! tid)) tid +++----------------------------------------------------------------------------+--+--unpairP :: Pushable arr => arr (a,a) -> Push a+--unpairP arr =+-- Push n $ \k -> pushf (everyOther k)+-- where+-- parr@(Push n pushf) = push arr++--everyOther :: (a -> Exp Word32 -> TProgram ()) +-- -> (a,a) -> Exp Word32 -> TProgram ()+--everyOther wf (a,b) ix = wf a (ix * 2) >> wf b (ix * 2 + 1) ++---------------------------------------------------------------------------+-- zipP+---------------------------------------------------------------------------+--zipP :: Pushable arr => arr a -> arr a -> Push a +--zipP arr1 arr2 =+-- Push (n1+n2)+-- $ \func -> p1 (\a ix -> func a (2*ix)) >>+-- p2 (\a ix -> func a (2*ix + 1))+-- +-- where +-- (Push n1 p1) = push arr1+-- (Push n2 p2) = push arr2+++---------------------------------------------------------------------------+-- *** GLOBAL ARRAYS *** --+---------------------------------------------------------------------------+{- +-- A very experimenental instance of mapG +mapG' :: (Pull a -> BProgram (Pull b))+ -> Word32+ -> GlobPull a+ -> GlobPush b+mapG' f n (GlobPull ixf) =+ GlobPush + $ \wf ->+ ForAllBlocks + $ \bix ->+ do -- BProgram do block + let pully = Pull n (\ix -> ixf (bix * fromIntegral n + ix))++ let res' = f pully + res <- res' + let numThreads = P.threadsPerBlock $ P.convPrg res'+ elemsPerThread = len res `div` numThreads++ if (numThreads == 0 || len res `mod` numThreads /= 0)+ then + ForAll (Just n) $ \ix -> wf (res ! ix) (bix * fromIntegral n + ix)+ else+ ForAll (Just n) $ \ix ->+ sequence_ [wf (res ! (ix + fromIntegral (numThreads * j)))+ (bix * fromIntegral n + (ix + fromIntegral (numThreads * j )))+ | j <- [0..elemsPerThread]]+++-- Old fasioned mapG+mapG :: (Pull a -> BProgram (Pull b))+ -> Word32 -- BlockSize ! + -> GlobPull a+ -> GlobPush b+mapG f n (GlobPull ixf) =+ GlobPush + $ \wf ->+ ForAllBlocks + $ \bix ->+ do -- BProgram do block + let pully = Pull n (\ix -> ixf (bix * fromIntegral n + ix))+ res <- f pully+ ForAll (Just (len res)) $ \ix -> wf (res ! ix) (bix * fromIntegral n + ix)+++-- I Think this one has more potential for generalisations. +mapD :: (Pull a -> BProgram b)+ -> Word32+ -> GlobPull a+ -> DistPull b+mapD f n (GlobPull ixf) =+ DistPull $ \bix ->+ do+ let pully = Pull n (\ix -> ixf (bix * fromIntegral n + ix))+ f pully ++mapDist :: (Pull a -> BProgram b)+ -> Word32+ -> GlobPull a+ -> (Exp Word32 -> BProgram b)+mapDist f n (GlobPull ixf) bix =+ let pully = Pull n (\ix -> ixf (bix * fromIntegral n + ix))+ in f pully ++-} ++---------------------------------------------------------------------------+-- See if having an Assignable/Allocable class is enough to generalise+-- the code below to pairs etc of Exp+---------------------------------------------------------------------------+-- Experimental (Improve upon this if it works)+-- I think this code looks quite horrific right now.+-- The potentially unnecessary assignments are pretty bad.++{- + +class ToGProgram a where+ type Global a+ toGProgram :: (Exp Word32 -> BProgram a) -> GProgram (Global a)++instance Scalar a => ToGProgram (Pull (Exp a)) where+ type Global (Pull (Exp a)) = GlobPush (Exp a)+ toGProgram f =+ do + let (pulla,_) = runPrg 0 $ f (BlockIdx X)+ let n = len pulla ++ shared <- uniqueSM++ ForAllBlocks $ \bix ->+ do+ res <- f bix -- compute res.++ -- Sync+ + Allocate shared (n * fromIntegral (sizeOf (undefined :: Exp a)))+ (Pointer (typeOf (undefined :: Exp a)))++ ForAll (Just n) $ \tix ->+ -- potentially unnessecary assignment...+ -- if the last thing the local computation does is force. + Assign shared tix (res ! tix)+++ Sync+ + return $+ GlobPush $ \wf ->+ do+ ForAllBlocks $ \bix-> + ForAll (Just n) $ \tix ->+ wf (index shared tix)+ (bix * fromIntegral n + tix)+ + + + +instance (Scalar a, Scalar b) => ToGProgram (Pull (Exp a), Pull (Exp b)) where+ type Global (Pull (Exp a),Pull (Exp b)) = (GlobPush (Exp a), GlobPush (Exp b))+ toGProgram f = + do + let ((pulla,pullb),_) = runPrg 0 $ f (BlockIdx X)+ let n1 = len pulla + let n2 = len pullb +++ shared1 <- uniqueSM+ shared2 <- uniqueSM ++ ForAllBlocks $ \bix ->+ do+ -- This is the heart of the matter. I want the f bix Program + -- to only occur once in the generated complete Program.+ (res1,res2) <- f bix -- compute results.++ -- Sync+ + ------------------------------------------------------------------+ Allocate shared1 (n1 * fromIntegral (sizeOf (undefined :: Exp a)))+ (Pointer (typeOf (undefined :: Exp a)))+ + ForAll (Just n1) $ \tix ->+ -- potentially unnessecary assignment...+ -- if the last thing the local computation does is force. + Assign shared1 tix (res1 ! tix)++ ------------------------------------------------------------------+ Allocate shared2 (n2 * fromIntegral (sizeOf (undefined :: Exp b)))+ (Pointer (typeOf (undefined :: Exp b)))+ + ForAll (Just n2) $ \tix ->+ -- potentially unnessecary assignment...+ -- if the last thing the local computation does is force. + Assign shared2 tix (res2 ! tix)++ Sync+ + let gp1 = + GlobPush $ \wf ->+ do+ ForAllBlocks $ \bix-> + ForAll (Just n1) $ \tix ->+ wf (index shared1 tix)+ (bix * fromIntegral n1 + tix) ++ let gp2 =+ GlobPush $ \wf ->+ do+ ForAllBlocks $ \bix-> + ForAll (Just n2) $ \tix ->+ wf (index shared2 tix)+ (bix * fromIntegral n2 + tix)++ + return (gp1,gp2)+-} +--------------------------------------------------------------------------- + ++-- The Problem is that I cannot share computations that+-- take place on the gridlevel (I think). +-- Experiment+{- +mapE :: forall a b . (Scalar a, Scalar b)+ => (Pull (Exp a) -> BProgram (Pull (Exp b)))+ -> Word32+ -> GlobPull (Exp a)+ -> GProgram (Pull (Exp b)) +mapE f n (GlobPull ixf) =+ do+++ shared <- uniqueSM + -- Allocate bytes in all shared memories and obtain a name+ + --return undefined + ForAllBlocks $ \bix ->+ do + let pully = Pull n (\ix -> ixf (bix * fromIntegral n + ix))+ res <- f pully++ Allocate shared (n * fromIntegral (sizeOf (undefined :: Exp b)))+ (Pointer (typeOf (undefined :: Exp b)) )+ ForAll (Just n) $ \ tid -> + do + Assign shared tid (res ! tid)+ + return $ Pull 32 $ \ix -> (index shared (ix `mod` (fromIntegral n)))+-} + + -- Assume Pull a here is one the special distributed pulls from above+{- +pushBlocks :: Pull a -> GlobPush a+pushBlocks (Pull n ixf) =+ GlobPush $ \wf ->+ ForAllBlocks $ \bix ->+ ForAll (Just n) $ \tix -> wf (ixf tix) (bix * fromIntegral n + tix) +++experiment :: GlobPull (Exp Int) -> GProgram (Pull (Exp Int), Pull (Exp Int))+experiment input =+ do + arr <- mapE force 32 input+ return (arr, singleton (arr ! 31))++experiment2 :: GlobPull (Exp Int) -> GProgram (GlobPush (Exp Int), GlobPush (Exp Int))+experiment2 input =+ do+ (arr1,arr2) <- experiment input+ return (pushBlocks arr1, pushBlocks arr2)+-} + +{-+ I think something like mapE above is needed to expressed shared computations.+ Bad things about mapE is its very specific type and+ that it allocates a new shared memory array and potentially performs a+ completely unnecessary copy from the old shared memory array.++ To generalise mapE typeclasses are probably needed.++ A way to skip the unnecessary copy would be if it was possible+ to hand the name over to the local computation. "compute this and+ store the result here".+ This sounds like some notion of a mutable array...+ What would it mean if we could express such local computations ?++ sklansky :: Int -> (Exp a -> Exp a -> Exp a) -> Pull (Exp a) -> Name -> BProgram (Pull (Exp a) + sklansky 0 op arr res =+ forAllN (len arr) $ \ix ->+ Assign res ix (arr ! ix)+ sklansky n op arr res =+ let arr1 = twoK (n-1) (fan op) arr+ arr2 <- force arr1+ sklansky (n-1) op arr2 ++ Not pure... But then, who says our dsl must be ? (isn't that part of the beauty+ of Haskell and embedded languages?). We are already in a wierd Program Monad+ and allow really dangerous push arrays... ++ Going even further then maybe force should take a mutable array and a+ push/pull array and they push the elements into the mutable array..+ This means that creation of mutable arrays must also be in the hands of+ the programmer.+ If array creation and force targets are in the hand of the programmer then+ in-placeness is also in the hands of the programmer + about a trillion+ new ways to shoot ones foot off. ++ I imagine that a mutable array could be simply a:+ data MArray a = MArray Word32 Name -- length and identifier.+-} ++++-- mapG2 is really hard to get right. +{- +mapG2 :: (Pull a -> BProgram (Pull b, Pull c))+ -> Word32+ -> GlobPull a+ -> (GlobPush b, GlobPush c) -- hard to get this right!+ -- without repeating computations. +mapG2 f n (GlobPull ixf) =+ (GlobPush+ $ \wf -> ForAllBlock+ $ \bix ->+ do+ let pully = Pull n (\ixf (bix * fromIntegral n + ix))+ res = f pully + , GlobPush +-} ++---------------------------------------------------------------------------+-- From Distributed array of blocks to a global push array+---------------------------------------------------------------------------+--toGlobPush :: Distrib (BProgram (Pull a))+-- -> GlobPush a +--toGlobPush inp@(Distrib bixf) =+-- GlobPush bs $+-- \wf -> ForAllBlocks +-- $ \bix ->+-- do -- BProgram do block +-- arr <- bixf bix +-- ForAll bs $ \ix -> wf (arr ! ix) bix ix +-- where+-- -- Is this Ok?! +-- bs = len $ fst $ runPrg 0 $ bixf 0++---------------------------------------------------------------------------+-- Create a global array that pushes to global+-- memory N elements per thread. +--------------------------------------------------------------------------- +--toGlobPushN :: Word32+-- -> Distrib (BProgram (Pull a))+-- -> GlobPush a+--toGlobPushN n dist =+-- GlobPush bs $ +-- \wf -> ForAllBlocks +-- $ \bix ->+-- do -- BProgram do block+-- arr <- getBlock dist bix+-- ForAll (bs `div` n) $+-- \ix ->+-- sequence_ +-- -- correct indexing ? +-- [wf (arr ! (ix * n' + i')) bix (ix * n' + i')+-- | i <- [0..n-1]+-- , let n' = fromIntegral n+-- , let i' = fromIntegral i]+-- +-- +-- where+-- bs = len $ fst $ runPrg 0 $ getBlock dist 0 +-- -- nb = numBlocks dist++ +---------------------------------------------------------------------------+--+---------------------------------------------------------------------------
+ Obsidian/LibraryG.hs view
@@ -0,0 +1,129 @@++{-# LANGUAGE ScopedTypeVariables #-}+module Obsidian.LibraryG where++import Obsidian.Array+import Obsidian.Program+import Obsidian.Exp+import Obsidian.Memory++import Data.Word+++---------------------------------------------------------------------------+-- +--------------------------------------------------------------------------- +mapG :: ASize l => (SPull a -> BProgram (SPull b))+ -> Pull l (SPull a)+ -> Push Grid l b+mapG kern as =+ Push (blocks * fromIntegral n) $+ \wf ->+ do+ forAllBlocks (sizeConv blocks) $ \bix -> do+ res <- kern (as ! bix)+ let (Push _ p) = push Block res+ wf' a ix = wf a (bix * sizeConv n + ix)+ p wf'+ where+ blocks = len as+ n = len (as ! 0)+++mapT :: (SPull a -> TProgram (SPull b))+ -> SPull (SPull a)+ -> SPush Block b+mapT threadf as =+ Push (n * m) $+ \wf ->+ do+ forAll (sizeConv n) $ \tix -> do+ res <- threadf (as ! tix)+ let (Push _ p) = push Thread res+ wf' a ix = wf a (tix * sizeConv m + ix)+ p wf' ++ where+ n = len as+ m = len (as ! 0) +---------------------------------------------------------------------------+-- +--------------------------------------------------------------------------- +zipWithG :: ASize l => (SPull a -> SPull b -> BProgram (SPull c))+ -> Pull l (SPull a)+ -> Pull l (SPull b)+ -> Push Grid l c+zipWithG kern as bs =+ Push (blocks * fromIntegral rn) $+ \wf ->+ do+ forAllBlocks (sizeConv blocks) $ \bix -> do+ res <- kern (as ! bix) (bs ! bix)+ let (Push _ p) = push Block res+ wf' a ix = wf a (bix * sizeConv n + ix)+ p wf'+ where+ -- Is this ok?! (does it break?) + rn = len $ fst $ runPrg 0 (kern (as ! 0) (bs ! 0))+ n = min m k+ -- assume uniformity+ m = len (as ! 0)+ k = len (bs ! 0)+ blocks = min (len as) (len bs) + ++zipWithT :: (SPull a -> SPull b -> TProgram (SPull c))+ -> SPull (SPull a)+ -> SPull (SPull b) + -> SPush Block c+zipWithT threadf as bs =+ Push (threads * rn) $+ \wf ->+ do+ forAll (sizeConv threads) $ \tix -> do+ res <- threadf (as ! tix) (bs ! tix) + let (Push _ p) = push Thread res+ wf' a ix = wf a (tix * sizeConv n + ix)+ p wf' ++ where++ -- Is this ok?! (does it break?) + rn = len $ fst $ runPrg 0 (threadf (as ! 0) (bs ! 0))+ n = min m k ++ m = len (as ! 0)+ k = len (bs ! 0)+ threads = min (len as) (len bs) ++---------------------------------------------------------------------------+-- Experimental+---------------------------------------------------------------------------+zipWithG' :: forall a b c l. (ASize l, MemoryOps c)+ => (SPull a -> SPull b -> BProgram (SPull c))+ -> Pull l (SPull a)+ -> Pull l (SPull b)+ -> GProgram (Push Grid l c)+zipWithG' kern as bs =+ do+ snames <- forAllBlocks (sizeConv n) $ \bix ->+ do+ res <- kern (as ! bix) (bs ! bix)+ let (Push _ p) = push Block res+ p (assignArrayN n)+ let pully = Pull blocks $ \bix -> (pullFromS snames n :: Pull Word32 c)+ + return $ Push (blocks * fromIntegral n) $+ \wf ->+ do+ forAllBlocks (sizeConv blocks) $ \bix ->+ forAll (sizeConv n) $ \tix ->+ do+ wf ((pully ! bix) ! tix) (bix * (sizeConv n) + tix)+ + where+ n = min m k + -- Assume uniformity+ m = len (as ! 0)+ k = len (bs ! 0)+ blocks = (min (len as) (len bs))
+ Obsidian/Memory.hs view
@@ -0,0 +1,184 @@+{-# LANGUAGE ScopedTypeVariables,+ GADTs #-}++{- Joel Svensson 2013++ This Module became quite messy.+ TODO: CLEAN IT UP! ++-} ++module Obsidian.Memory (MemoryOps(..),GlobalMemoryOps(..),assignArrayN) where+++import Obsidian.Program+import Obsidian.Exp+import Obsidian.Types+import Obsidian.Globs+import Obsidian.Array -- Importing this feels a bit strange.+import Obsidian.Names++import Data.Word++class Inspect a where+ inspect :: a -> Tree (Either (Kind,Name) Type)++instance Scalar a => Inspect (Exp a) where+ inspect (Index (name,[])) = Single (Left (Var,name))+ inspect (Index (name,[ix])) | isid ix = Single (Left (Arr,name))+ inspect _ = Single $ Right (typeOf (undefined :: Exp a)) ++instance (Inspect a, Inspect b) => Inspect (a,b) where+ inspect (a,b) = Tuple [inspect a, inspect b]++instance (Inspect a, Inspect b, Inspect c) => Inspect (a,b,c) where+ inspect (a,b,c) = Tuple [inspect a, inspect b, inspect c]++isid (ThreadIdx X) = True+isid _ = False++---------------------------------------------------------------------------+-- Local Memory+---------------------------------------------------------------------------+class Inspect a => MemoryOps a where+ names :: String -> a -> Program t Names+ allocateArray :: Names -> a -> Word32 -> Program t ()+ allocateScalar :: Names -> a -> Program t () + assignArray :: Names -> a -> Exp Word32 -> TProgram ()+ assignArrayS :: Tree (Either (Kind,Name) (Kind,Name))+ -> a+ -> Exp Word32+ -> TProgram (Tree (Kind,Name)) + assignScalar :: Names -> a -> TProgram () + pullFrom :: Names -> Word32 -> Pull Word32 a+ readFrom :: Names -> a++ pullFromS :: Tree (Kind,Name) -> Word32 -> Pull Word32 a++---------------------------------------------------------------------------+-- Derived+---------------------------------------------------------------------------+assignArrayN :: MemoryOps a =>+ Word32 -> a -> Exp Word32 -> TProgram (Tree (Kind,Name))+assignArrayN n a ix =+ do+ names <- allocateNeeded n insp+ assignArrayS names a ix+ -- return names+ where+ insp = inspect a+++allocateNeeded :: Word32 -> Tree (Either (Kind,Name) Type) -> TProgram (Tree (Either (Kind,Name) (Kind,Name)))+allocateNeeded n None = return None+allocateNeeded n a@(Single (Left (k,nom))) = return $ Single (Left (k,nom))+allocateNeeded n (Single (Right t)) =+ do+ name <- uniqueNamed "arr"+ Allocate name (n * typeSize t) (Pointer t)+ return $ Single (Right (Arr,name))+allocateNeeded n (Tuple xs) =+ do+ ns <- mapM (allocateNeeded n) xs+ return $ Tuple ns + ++---------------------------------------------------------------------------+-- Instances+---------------------------------------------------------------------------+instance Scalar a => MemoryOps (Exp a) where+ names pre a = do {i <- uniqueNamed pre; return (Single i)}+ allocateArray (Single name) a n = + Allocate name (n * fromIntegral (sizeOf a))+ (Pointer (typeOf a))+ allocateScalar (Single name) a =+ Declare name (typeOf a) + assignArray (Single name) a ix = Assign name [ix] a+ assignArrayS (Single (Left (k,name))) a ix = return $ Single (k,name)+ assignArrayS (Single (Right (k,name))) a ix =+ do + Assign name [ix] a -- ? + return $ Single (k,name)+ + ++ ++ assignScalar (Single name) a = Assign name [] a + pullFrom (Single name) n = Pull n (\i -> index name i) + readFrom (Single name) = variable name++ pullFromS (Single (Var,name)) n = Pull n $ \_ -> variable name+ pullFromS (Single (Arr,name)) n = Pull n (\i -> index name i) ++instance (MemoryOps a, MemoryOps b) => MemoryOps (a, b) where+ names pre (a,b) =+ do+ a' <- names pre a+ b' <- names pre b+ return $ Tuple [a', b']+ allocateArray (Tuple [ns1,ns2]) (a,b) n =+ do + allocateArray ns1 a n+ allocateArray ns2 b n+ allocateScalar (Tuple [ns1,ns2]) (a,b) =+ do+ allocateScalar ns1 a+ allocateScalar ns2 b + assignArray (Tuple [ns1,ns2]) (a,b) ix =+ do+ assignArray ns1 a ix + assignArray ns2 b ix++ assignArrayS (Tuple [ns1,ns2]) (a,b) ix =+ do+ nas1 <- assignArrayS ns1 a ix + nas2 <- assignArrayS ns2 b ix+ return $ Tuple [nas1,nas2]+ assignScalar (Tuple [ns1,ns2]) (a,b) =+ do+ assignScalar ns1 a + assignScalar ns2 b + pullFrom (Tuple [ns1,ns2]) n =+ let p1 = pullFrom ns1 n+ p2 = pullFrom ns2 n+ in Pull n (\ix -> (p1 ! ix, p2 ! ix))+ readFrom (Tuple [ns1,ns2]) =+ let p1 = readFrom ns1+ p2 = readFrom ns2+ in (p1,p2)++ pullFromS (Tuple [ns1,ns2]) n =+ let p1 = pullFromS ns1 n+ p2 = pullFromS ns2 n+ in Pull n (\ix -> (p1 ! ix, p2 ! ix))+++---------------------------------------------------------------------------+-- Global Memory+---------------------------------------------------------------------------++class GlobalMemoryOps a where+ outputs :: a -> GProgram Names+ assignOut :: Names -> a -> Exp Word32 -> Program Thread ()+++instance Scalar a => GlobalMemoryOps (Exp a) where+ outputs a =+ do+ name <- Output $ Pointer $ typeOf a+ return (Single name) + assignOut (Single name) a ix = Assign name [ix] a+++instance (GlobalMemoryOps a, GlobalMemoryOps b)+ => GlobalMemoryOps (a,b) where+ outputs (a,b) =+ do+ na <- outputs a+ nb <- outputs b+ return (Tuple [na,nb]) + assignOut (Tuple [n1,n2]) (a,b) ix =+ do+ assignOut n1 a ix + assignOut n2 b ix
+ Obsidian/Names.hs view
@@ -0,0 +1,20 @@+++module Obsidian.Names where++import Obsidian.Globs++data Tree a = None+ | Single a+ | Tuple [Tree a]++type Names = Tree Name++type NameInfo = Tree (Kind,Name)++data Kind = Var | Arr++instance Functor Tree where+ fmap f None = None+ fmap f (Single a) = Single $ f a+ fmap f (Tuple ts) = Tuple $ map (fmap f) ts
+ Obsidian/Program.hs view
@@ -0,0 +1,263 @@+{- Joel Svensson 2012,2013++ Notes:+ 2013-01-08: removed number-of-blocks field from ForAllBlocks++-}++{-# LANGUAGE GADTs,+ FlexibleInstances,+ EmptyDataDecls #-} +++module Obsidian.Program 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++---------------------------------------------------------------------------+-- Thread/Block/Grid +---------------------------------------------------------------------------+data Thread+data Block+data Grid ++data PT a where+ Thread :: PT Thread+ Block :: PT Block+ Grid :: PT Grid ++type Identifier = Int ++---------------------------------------------------------------------------+-- Obsidian+---------------------------------------------------------------------------+data Obsidian a = Obsidian (Program Grid a)+ ++---------------------------------------------------------------------------+-- Program datatype+--------------------------------------------------------------------------+data Program t a where+ + Identifier :: Program t Identifier + + Assign :: Scalar a+ => Name+ -> [Exp Word32]+ -> (Exp a)+ -> Program Thread ()+ + + AtomicOp :: Scalar a+ => Name + -> Exp Word32+ -> Atomic a+ -> Program Thread (Exp a)++ Cond :: Exp Bool+ -> Program t ()+ -> Program t ()+ + -- DONE: Code generation for this.+ -- TODO: Generalize this loop! (Replace Thread with t) + SeqFor :: Exp Word32 -> (Exp Word32 -> Program t a)+ -> Program t a+ + ForAll :: (Exp Word32) + -> (Exp Word32 -> Program Thread a)+ -> Program Block a ++ {-+ I'm not sure about this constructor.+ As I see it programs from which we generate a kernel+ must be wrapped in one of these ForAllBlocks.+ Programs with sequences of 'ForAllBlocks' are problematic.++ Maybe a (ForAllBlocks n f *>* ForAllBlocks m g) Program+ should be split into two kernels. + -} + ForAllBlocks :: (Exp Word32) -> (Exp Word32 -> Program Block a) + -> Program Grid a++ ForAllThreads :: (Exp Word32) -> (Exp Word32 -> Program Thread a)+ -> Program Grid a ++ -- Allocate shared memory in each MP++ -- TODO: Change the Liveness analysis to a two-pass algo+ -- and remove the Allocate constructor. + Allocate :: Name -> Word32 -> Type -> Program t () ++ -- Automatic Variables+ Declare :: Name -> Type -> Program t () + + {- About Output (Creates a named output array). + This is similar to Allocate but concerning global arrays.++ Since we cannot synchronize writes to a global array inside of an+ kernel, global arrays will only be written as outputs of the kernel+ -} + Output :: Type -> Program Grid Name+ -- (Output may be replaced by AllocateG) + + Sync :: Program Block ()+ -- Two very experimental threadfence constructs.+ -- should correspond to cuda __threadfence();+ -- and __threadfenceBlock(); + ThreadFence :: Program Grid ()+ ThreadFenceBlock :: Program Block () ++ -- Parallel composition of Programs+ -- TODO: Will I use this ? + Par :: Program p () ->+ Program p () ->+ Program p () ++ -- Monad+ Return :: a -> Program t a+ Bind :: Program t a -> (a -> Program t b) -> Program t b++---------------------------------------------------------------------------+-- Helpers +--------------------------------------------------------------------------- +uniqueSM = do+ id <- Identifier+ return $ "arr" ++ show id ++uniqueNamed pre = do+ id <- Identifier+ return $ pre ++ show id ++---------------------------------------------------------------------------+-- forAll and forAllN+---------------------------------------------------------------------------+--forAll :: (Exp Word32 -> Program Thread ()) -> Program Block () +--forAll f = ForAll Nothing f++forAll :: Exp Word32 -> (Exp Word32 -> Program Thread a) -> Program Block a+forAll n f = ForAll n f++(*||*) = Par++---------------------------------------------------------------------------+-- SeqFor+---------------------------------------------------------------------------+seqFor :: Exp Word32 -> (Exp Word32 -> Program t a)+ -> Program t a+seqFor (Literal 1) f = f 0+seqFor n f = SeqFor n f+++---------------------------------------------------------------------------+-- forAllT+--------------------------------------------------------------------------- +-- When we know that all threads are independent and+-- independent of "blocksize".+-- Also any allocation of local storage is impossible.+-- Composition of something using forAllT with something+-- that performs local computations is impossible.+-- Using the hardcoded BlockDim may turn out to be a problem when+-- we want to compute more than one thing per thread (may be fine though). +forAllT :: (Exp Word32) -> (Exp Word32 -> Program Thread a)+ -> Program Grid a+forAllT n f = ForAllThreads n + $ \gtid -> f gtid ++++forAllBlocks = ForAllBlocks++---------------------------------------------------------------------------+-- Monad+--------------------------------------------------------------------------+instance Monad (Program t) where+ return = Return+ (>>=) = Bind++---------------------------------------------------------------------------+-- Aliases +---------------------------------------------------------------------------+type TProgram = Program Thread+type BProgram = Program Block+type GProgram = Program Grid ++---------------------------------------------------------------------------+-- runPrg (fix types here, Integer!)+---------------------------------------------------------------------------+runPrg :: Int -> Program t a -> (a,Int)+runPrg i Identifier = (i,i+1)+runPrg i (Return a) = (a,i)+runPrg i (Bind m f) =+ let (a,i') = runPrg i m+ in runPrg i' (f a) +runPrg i (Sync) = ((),i)+runPrg i (ForAll n ixf) =+ let (p,i') = runPrg i (ixf (variable "tid")) + in (p,i') +runPrg i (Allocate id _ _ ) = ((),i)+runPrg i (Assign _ _ a) = ((),i) -- Probaby wrong.. +runPrg i (AtomicOp _ _ _) = (variable ("new"++show i),i+1)+ +---------------------------------------------------------------------------+-- printPrg+---------------------------------------------------------------------------+printPrg prg = (\(_,x,_) -> x) $ printPrg' 0 prg++printPrg' :: Int -> Program t a -> (a,String,Int)+printPrg' i Identifier = (i,"getId;\n",i+1) +-- printPrg' i Skip = ((),";\n", i)+printPrg' i (Assign n ix e) =+ ((),n ++ "[" ++ show ix ++ "] = " ++ show e ++ ";\n", i) +printPrg' i (AtomicOp n ix e) =+ let newname = "r" ++ show i+ in (variable newname,+ newname ++ " = " ++ printAtomic e +++ "( " ++ n ++ "[" ++ show ix ++ "])\n",i+1)+printPrg' i (Allocate id n t) =+ let newname = id -- "arr" ++ show id+ in ((),newname ++ " = malloc(" ++ show n ++ ");\n",i+1)+printPrg' i (Declare id t) =+ let newname = id -- "arr" ++ show id+ in ((),show t ++ " " ++ newname ++ "\n",i+1)+printPrg' i (Output t) =+ let newname = "globalOut" ++ show i+ in (newname,newname ++ " = new Global output;\n",i+1)+printPrg' i (SeqFor n f) =+ let (a,prg2,i') = printPrg' i (f (variable "i"))+ + in ( a, + "for (i in 0.." ++ show n ++ ")" +++ "{\n" ++ prg2 ++ "\n}",+ i')+ +printPrg' i (ForAll n f) =+ let (a,prg2,i') = printPrg' i (f (variable "i"))+ + in ( a, + "par (i in 0.." ++ show n ++ ")" +++ "{\n" ++ prg2 ++ "\n}",+ i')+printPrg' i (ForAllBlocks n f) =+ let (d,prg2,i') = printPrg' i (f (variable "BIX"))+ in (d, + "blocks (i)" +++ "{\n" ++ prg2 ++ "\n}",+ i')+printPrg' i (Return a) = (a,"MonadReturn;\n",i)+printPrg' i (Bind m f) =+ let (a1, str1,i1) = printPrg' i m+ (a2,str2,i2) = printPrg' i1 (f a1)+ in (a2,str1 ++ str2, i2)+printPrg' i Sync = ((),"Sync;\n",i)+
+ Obsidian/SeqLoop.hs view
@@ -0,0 +1,102 @@+{-# LANGUAGE ScopedTypeVariables #-} +{-++ sequential loops with state+ 2013 : Joel Svensson ++-}++module Obsidian.SeqLoop where+++import Obsidian.Program+import Obsidian.Exp+import Obsidian.Array+import Obsidian.Memory+import Obsidian.Names++-- TODO: Add suitable allocs+-- TODO: Rename module to something better++---------------------------------------------------------------------------+-- seqFold (actually reduce) +---------------------------------------------------------------------------++seqFold :: forall l a. (ASize l, MemoryOps a)+ => (a -> a -> a)+ -> a+ -> Pull l a+ -> Program Thread a+seqFold op init arr = do+ ns <- names "v" (undefined :: a) + allocateScalar ns (undefined :: a)++ assignScalar ns init + -- Assign nom [] init + SeqFor n $ \ ix ->+ do+ assignScalar ns (readFrom ns `op` (arr ! ix))+ return None++ + return $ readFrom ns+ where + n = sizeConv$ len arr+++---------------------------------------------------------------------------+-- Sequential scan+---------------------------------------------------------------------------++seqScan :: forall l a. (ASize l, MemoryOps a)+ => (a -> a -> a)+ -> Pull l a+ -> Push Thread l a+seqScan op (Pull n ixf) =+ Push n $ \wf -> do+ ns <- names "v" (undefined :: a) + allocateScalar ns (undefined :: a)+ assignScalar ns (ixf 0)+ wf (readFrom ns) 0 + SeqFor (sizeConv (n-1)) $ \ix -> do+ wf (readFrom ns) ix + assignScalar ns $ readFrom ns `op` (ixf (ix + 1))+ return None+ +---------------------------------------------------------------------------+-- Sequential Map (here for uniformity) +---------------------------------------------------------------------------++seqMap :: forall l a b. ASize l+ => (a -> b)+ -> Pull l a+ -> Push Thread l b+seqMap f arr =+ Push (len arr) $ \wf -> do+ SeqFor (sizeConv (len arr)) $ \ix ->+ wf (f (arr ! ix)) ix +++---------------------------------------------------------------------------+-- Sequential Map and scan (generalisation of map + accum) +---------------------------------------------------------------------------++seqMapScan :: forall l a b acc. (ASize l, MemoryOps acc, MemoryOps b)+ => (acc -> a -> (acc,b))+ -> acc + -> Pull l a+ -> Push Thread l (acc,b)+seqMapScan op acc (Pull n ixf) =+ Push n $ \wf -> do+ ns <- names "v" (undefined :: b) + allocateScalar ns (undefined :: b)+ nacc <- names "v" (undefined :: acc)+ allocateScalar nacc (undefined :: acc)+ + assignScalar nacc acc++ SeqFor (sizeConv n) $ \ix -> do+ let (a,b) = op (readFrom nacc) (ixf ix)+ wf (a,b) ix + assignScalar nacc a+ return None
+ Obsidian/Types.hs view
@@ -0,0 +1,33 @@+module Obsidian.Types where +++---------------------------------------------------------------------------+-- Types+---------------------------------------------------------------------------++data Type +-- The allowed scalar types+ = Bool+ | Int | Word -- A bit problematic since the size of+ -- of these are platform dependent+ | Int8 | Int16 | Int32 | Int64 + | Word8 | Word16 | Word32 | Word64 + | Float | Double + +-- Used by CUDA, C And OpenCL generators + | Pointer Type -- C thing + | Global Type -- OpenCL thing+ | Local Type -- OpenCL thing+ deriving (Eq, Show)+ +typeSize Int8 = 1+typeSize Int16 = 2+typeSize Int32 = 4+typeSize Int64 = 8+typeSize Word8 = 1+typeSize Word16 = 2+typeSize Word32 = 4+typeSize Word64 = 8+typeSize Bool = 4+typeSize Float = 4+typeSize Double = 8
+ Setup.hs view
@@ -0,0 +1,2 @@+import Distribution.Simple+main = defaultMain