diff --git a/LICENSE b/LICENSE new file mode 100644 --- /dev/null +++ b/LICENSE @@ -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 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 BE LIABLE FOR ANY +DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES +(INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; +LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND +ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS +SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. diff --git a/Obsidian.cabal b/Obsidian.cabal new file mode 100644 --- /dev/null +++ b/Obsidian.cabal @@ -0,0 +1,79 @@ +Name: Obsidian +Version: 0.0.0.1 + +License: BSD3 +License-file: LICENSE +Stability: Beta +Maintainer: Joel Svensson +Author: Joel Svensson + +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: + + diff --git a/Obsidian.hs b/Obsidian.hs new file mode 100644 --- /dev/null +++ b/Obsidian.hs @@ -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 diff --git a/Obsidian/Array.hs b/Obsidian/Array.hs new file mode 100644 --- /dev/null +++ b/Obsidian/Array.hs @@ -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 diff --git a/Obsidian/Atomic.hs b/Obsidian/Atomic.hs new file mode 100644 --- /dev/null +++ b/Obsidian/Atomic.hs @@ -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" diff --git a/Obsidian/CodeGen/CUDA.hs b/Obsidian/CodeGen/CUDA.hs new file mode 100644 --- /dev/null +++ b/Obsidian/CodeGen/CUDA.hs @@ -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] + diff --git a/Obsidian/CodeGen/Common.hs b/Obsidian/CodeGen/Common.hs new file mode 100644 --- /dev/null +++ b/Obsidian/CodeGen/Common.hs @@ -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" + + diff --git a/Obsidian/CodeGen/InOut.hs b/Obsidian/CodeGen/InOut.hs new file mode 100644 --- /dev/null +++ b/Obsidian/CodeGen/InOut.hs @@ -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 + + diff --git a/Obsidian/CodeGen/Liveness.hs b/Obsidian/CodeGen/Liveness.hs new file mode 100644 --- /dev/null +++ b/Obsidian/CodeGen/Liveness.hs @@ -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) + diff --git a/Obsidian/CodeGen/Memory.hs b/Obsidian/CodeGen/Memory.hs new file mode 100644 --- /dev/null +++ b/Obsidian/CodeGen/Memory.hs @@ -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 + diff --git a/Obsidian/CodeGen/PP.hs b/Obsidian/CodeGen/PP.hs new file mode 100644 --- /dev/null +++ b/Obsidian/CodeGen/PP.hs @@ -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 diff --git a/Obsidian/CodeGen/Program.hs b/Obsidian/CodeGen/Program.hs new file mode 100644 --- /dev/null +++ b/Obsidian/CodeGen/Program.hs @@ -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" diff --git a/Obsidian/CodeGen/SPMDC.hs b/Obsidian/CodeGen/SPMDC.hs new file mode 100644 --- /dev/null +++ b/Obsidian/CodeGen/SPMDC.hs @@ -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 diff --git a/Obsidian/DimSpec.hs b/Obsidian/DimSpec.hs new file mode 100644 --- /dev/null +++ b/Obsidian/DimSpec.hs @@ -0,0 +1,7 @@ + +{- Joel Svensson 2012 -} + +module Obsidian.DimSpec (DimSpec (..)) where + +data DimSpec = X | Y | Z + deriving (Eq,Ord,Show) diff --git a/Obsidian/Exp.hs b/Obsidian/Exp.hs new file mode 100644 --- /dev/null +++ b/Obsidian/Exp.hs @@ -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 + diff --git a/Obsidian/Force.hs b/Obsidian/Force.hs new file mode 100644 --- /dev/null +++ b/Obsidian/Force.hs @@ -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 () diff --git a/Obsidian/Globs.hs b/Obsidian/Globs.hs new file mode 100644 --- /dev/null +++ b/Obsidian/Globs.hs @@ -0,0 +1,16 @@ + +{- Joel Svensson 2012 -} + +module Obsidian.Globs where + + +import Data.Word + +--------------------------------------------------------------------------- +-- Aliases + +type Name = String + + + +type NumThreads = Word32 diff --git a/Obsidian/Library.hs b/Obsidian/Library.hs new file mode 100644 --- /dev/null +++ b/Obsidian/Library.hs @@ -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 + + +--------------------------------------------------------------------------- +-- +--------------------------------------------------------------------------- diff --git a/Obsidian/LibraryG.hs b/Obsidian/LibraryG.hs new file mode 100644 --- /dev/null +++ b/Obsidian/LibraryG.hs @@ -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)) diff --git a/Obsidian/Memory.hs b/Obsidian/Memory.hs new file mode 100644 --- /dev/null +++ b/Obsidian/Memory.hs @@ -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 diff --git a/Obsidian/Names.hs b/Obsidian/Names.hs new file mode 100644 --- /dev/null +++ b/Obsidian/Names.hs @@ -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 diff --git a/Obsidian/Program.hs b/Obsidian/Program.hs new file mode 100644 --- /dev/null +++ b/Obsidian/Program.hs @@ -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) + diff --git a/Obsidian/SeqLoop.hs b/Obsidian/SeqLoop.hs new file mode 100644 --- /dev/null +++ b/Obsidian/SeqLoop.hs @@ -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 diff --git a/Obsidian/Types.hs b/Obsidian/Types.hs new file mode 100644 --- /dev/null +++ b/Obsidian/Types.hs @@ -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 diff --git a/Setup.hs b/Setup.hs new file mode 100644 --- /dev/null +++ b/Setup.hs @@ -0,0 +1,2 @@ +import Distribution.Simple +main = defaultMain