Obsidian-0.0.0.5: Obsidian/Exp.hs
{-# 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)
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