accelerate 0.7.1.0 → 0.8.0.0
raw patch · 57 files changed
+2934/−1232 lines, 57 filesdep +QuickCheckdep ~fclabelsPVP ok
version bump matches the API change (PVP)
Dependencies added: QuickCheck
Dependency ranges changed: fclabels
API changes (from Hackage documentation)
+ Data.Array.Accelerate: Clamp :: Boundary a
+ Data.Array.Accelerate: Constant :: a -> Boundary a
+ Data.Array.Accelerate: Mirror :: Boundary a
+ Data.Array.Accelerate: Wrap :: Boundary a
+ Data.Array.Accelerate: class IsBounded a
+ Data.Array.Accelerate: class (Floating a, IsScalar a, IsNum a) => IsFloating a
+ Data.Array.Accelerate: class (IsScalar a, IsNum a, IsBounded a) => IsIntegral a
+ Data.Array.Accelerate: class IsNonNum a
+ Data.Array.Accelerate: class (Num a, IsScalar a) => IsNum a
+ Data.Array.Accelerate: class (Typeable a) => IsScalar a
+ Data.Array.Accelerate: class (Elem (StencilRepr dim stencil), Stencil dim a (StencilRepr dim stencil)) => Stencil dim a stencil
+ Data.Array.Accelerate: data Boundary a
+ Data.Array.Accelerate: stencil :: (Ix dim, Elem a, Elem b, Stencil dim a stencil) => (stencil -> Exp b) -> Boundary a -> Acc (Array dim a) -> Acc (Array dim b)
+ Data.Array.Accelerate: stencil2 :: (Ix dim, Elem a, Elem b, Elem c, Stencil dim a stencil1, Stencil dim b stencil2) => (stencil1 -> stencil2 -> Exp c) -> Boundary a -> Acc (Array dim a) -> Boundary b -> Acc (Array dim b) -> Acc (Array dim c)
+ Data.Array.Accelerate: type Stencil3 a = (Exp a, Exp a, Exp a)
+ Data.Array.Accelerate: type Stencil3x3 a = (Stencil3 a, Stencil3 a, Stencil3 a)
+ Data.Array.Accelerate: type Stencil3x3x3 a = (Stencil3x3 a, Stencil3x3 a, Stencil3x3 a)
+ Data.Array.Accelerate: type Stencil3x3x5 a = (Stencil3x3 a, Stencil3x3 a, Stencil3x3 a, Stencil3x3 a, Stencil3x3 a)
+ Data.Array.Accelerate: type Stencil3x5 a = (Stencil3 a, Stencil3 a, Stencil3 a, Stencil3 a, Stencil3 a)
+ Data.Array.Accelerate: type Stencil3x5x3 a = (Stencil3x5 a, Stencil3x5 a, Stencil3x5 a)
+ Data.Array.Accelerate: type Stencil3x5x5 a = (Stencil3x5 a, Stencil3x5 a, Stencil3x5 a, Stencil3x5 a, Stencil3x5 a)
+ Data.Array.Accelerate: type Stencil5 a = (Exp a, Exp a, Exp a, Exp a, Exp a)
+ Data.Array.Accelerate: type Stencil5x3 a = (Stencil5 a, Stencil5 a, Stencil5 a)
+ Data.Array.Accelerate: type Stencil5x3x3 a = (Stencil5x3 a, Stencil5x3 a, Stencil5x3 a)
+ Data.Array.Accelerate: type Stencil5x3x5 a = (Stencil5x3 a, Stencil5x3 a, Stencil5x3 a, Stencil5x3 a, Stencil5x3 a)
+ Data.Array.Accelerate: type Stencil5x5 a = (Stencil5 a, Stencil5 a, Stencil5 a, Stencil5 a, Stencil5 a)
+ Data.Array.Accelerate: type Stencil5x5x3 a = (Stencil5x5 a, Stencil5x5 a, Stencil5x5 a)
+ Data.Array.Accelerate: type Stencil5x5x5 a = (Stencil5x5 a, Stencil5x5 a, Stencil5x5 a, Stencil5x5 a, Stencil5x5 a)
+ Data.Array.Accelerate: type Stencil7 a = (Exp a, Exp a, Exp a, Exp a, Exp a, Exp a, Exp a)
+ Data.Array.Accelerate: type Stencil9 a = (Exp a, Exp a, Exp a, Exp a, Exp a, Exp a, Exp a, Exp a, Exp a)
Files
- Data/Array/Accelerate.hs +8/−1
- Data/Array/Accelerate/AST.hs +210/−55
- Data/Array/Accelerate/Analysis/Type.hs +32/−1
- Data/Array/Accelerate/Array/Representation.hs +43/−9
- Data/Array/Accelerate/Array/Sugar.hs +126/−9
- Data/Array/Accelerate/CUDA.hs +13/−3
- Data/Array/Accelerate/CUDA/Analysis/Hash.hs +28/−17
- Data/Array/Accelerate/CUDA/Analysis/Launch.hs +23/−10
- Data/Array/Accelerate/CUDA/Array/Data.hs +43/−29
- Data/Array/Accelerate/CUDA/CodeGen.hs +141/−66
- Data/Array/Accelerate/CUDA/CodeGen/Skeleton.hs +58/−35
- Data/Array/Accelerate/CUDA/CodeGen/Tuple.hs +4/−4
- Data/Array/Accelerate/CUDA/CodeGen/Util.hs +71/−50
- Data/Array/Accelerate/CUDA/Compile.hs +11/−6
- Data/Array/Accelerate/CUDA/Execute.hs +122/−28
- Data/Array/Accelerate/CUDA/Smart.hs +77/−0
- Data/Array/Accelerate/CUDA/State.hs +67/−71
- Data/Array/Accelerate/Internal/Check.hs +107/−0
- Data/Array/Accelerate/Interpreter.hs +62/−0
- Data/Array/Accelerate/Language.hs +83/−13
- Data/Array/Accelerate/Pretty.hs +22/−1
- Data/Array/Accelerate/Smart.hs +279/−18
- Data/Array/Accelerate/Test.hs +18/−0
- Data/Array/Accelerate/Test/QuickCheck.hs +18/−0
- Data/Array/Accelerate/Test/QuickCheck/Arbitrary.hs +236/−0
- Data/Array/Accelerate/Tuple.hs +33/−10
- Data/Array/Accelerate/Type.hs +56/−9
- LICENSE +2/−2
- accelerate.cabal +67/−21
- cubits/accelerate_cuda_extras.h +167/−30
- cubits/backpermute.inl +8/−4
- cubits/fold.inl +28/−22
- cubits/fold_segmented.inl +120/−0
- cubits/permute.inl +9/−7
- cubits/replicate.inl +30/−0
- cubits/slice.inl +31/−0
- cubits/thrust/scan_safe.inl +6/−31
- cubits/zipWith.inl +9/−4
- examples/simple/DotP.hs +0/−25
- examples/simple/Filter.hs +0/−39
- examples/simple/Main.hs +0/−372
- examples/simple/Makefile +13/−13
- examples/simple/SAXPY.hs +0/−24
- examples/simple/SMVM.hs +0/−39
- examples/simple/Square.hs +0/−23
- examples/simple/Sum.hs +0/−23
- examples/simple/Time.hs +0/−108
- examples/simple/src/DotP.hs +22/−0
- examples/simple/src/Filter.hs +40/−0
- examples/simple/src/Main.hs +123/−0
- examples/simple/src/Random.hs +60/−0
- examples/simple/src/SAXPY.hs +21/−0
- examples/simple/src/SMVM.hs +39/−0
- examples/simple/src/Square.hs +23/−0
- examples/simple/src/Stencil.hs +77/−0
- examples/simple/src/Sum.hs +23/−0
- include/accelerate.h +25/−0
Data/Array/Accelerate.hs view
@@ -1,6 +1,6 @@ -- | -- Module : Data.Array.Accelerate--- Copyright : [2008..2009] Manuel M T Chakravarty, Gabriele Keller, Sean Lee+-- Copyright : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee -- License : BSD3 -- -- Maintainer : Manuel M T Chakravarty <chak@cse.unsw.edu.au>@@ -27,6 +27,10 @@ -- -- * "Data.Array.Accelerate.Interpreter": simple interpreter in Haskell as a -- reference implementation defining the semantics of the Accelerate language+--+-- * "Data.Array.Accelerate.CUDA": an implementation supporting parallel+-- execution on CUDA-capable NVIDIA GPUs+-- module Data.Array.Accelerate ( @@ -35,6 +39,9 @@ CShort, CUShort, CInt, CUInt, CLong, CULong, CLLong, CULLong, Float, Double, CFloat, CDouble, Bool, Char, CChar, CSChar, CUChar,++ -- * Scalar type classes+ IsScalar, IsNum, IsBounded, IsIntegral, IsFloating, IsNonNum, -- * Array data types Array, Scalar, Vector, Segments,
Data/Array/Accelerate/AST.hs view
@@ -1,65 +1,66 @@-{-# LANGUAGE GADTs, EmptyDataDecls, FlexibleContexts, TypeFamilies #-}+{-# LANGUAGE GADTs, EmptyDataDecls, FlexibleContexts, TypeFamilies, TypeOperators #-}+{-# LANGUAGE FlexibleInstances, MultiParamTypeClasses, TypeSynonymInstances #-} --- |Embedded array processing language: accelerate AST with de Bruijn indices+-- Module : Data.Array.Accelerate.AST+-- Copyright : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee+-- License : BSD3 ----- Copyright (c) [2008..2009] Manuel M T Chakravarty, Gabriele Keller, Sean Lee+-- Maintainer : Manuel M T Chakravarty <chak@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions) ----- License: BSD3+-- /Scalar versus collective operations/ ------ Description ---------------------------------------------------------------+-- The embedded array processing language is a two-level language. It+-- combines a language of scalar expressions and functions with a language of+-- collective array operations. Scalar expressions are used to compute+-- arguments for collective operations and scalar functions are used to+-- parametrise higher-order, collective array operations. The two-level+-- structure, in particular, ensures that collective operations cannot be+-- parametrised with collective operations; hence, we are following a flat+-- data-parallel model. The collective operations manipulate+-- multi-dimensional arrays whose shape is explicitly tracked in their types.+-- In fact, collective operations cannot produce any values other than+-- multi-dimensional arrays; when they yield a scalar, this is in the form of+-- a 0-dimensional, singleton array. Similarly, scalar expression can -as+-- their name indicates- only produce tuples of scalar, but not arrays. ----- Scalar versus collective operations--- ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~--- The embedded array processing language is a two-level language. It--- combines a language of scalar expressions and functions with a language of--- collective array operations. Scalar expressions are used to compute--- arguments for collective operations and scalar functions are used to--- parametrise higher-order, collective array operations. The two-level--- structure, in particular, ensures that collective operations cannot be--- parametrised with collective operations; hence, we are following a flat--- data-parallel model. The collective operations manipulate--- multi-dimensional arrays whose shape is explicitly tracked in their types.--- In fact, collective operations cannot produce any values other than--- multi-dimensional arrays; when they yield a scalar, this is in the form of--- a 0-dimensional, singleton array. Similarly, scalar expression can -as--- their name indicates- only produce tuples of scalar, but not arrays. +-- There are, however, two expression forms that take arrays as arguments. As+-- a result scalar and array expressions are recursively dependent. As we+-- cannot and don't want to compute arrays in the middle of scalar+-- computations, array computations will always be hoisted out of scalar+-- expressions. So that this is always possible, these array expressions may+-- not contain any free scalar variables. To express that condition in the+-- type structure, we use separate environments for scalar and array variables. ----- There are, however, two expression forms that take arrays as arguments. As--- a result scalar and array expressions are recursively dependent. As we--- cannot and don't want to compute arrays in the middle of scalar--- computations, array computations will always be hoisted out of scalar--- expressions. So that this is always possible, these array expressions may--- not contain any free scalar variables. To express that condition in the--- type structure, we use separate environments for scalar and array variables.+-- /Programs/ ----- Programs--- ~~~~~~~~--- Collective array programs comprise closed expressions of array operations.--- There is no explicit sharing in the initial AST form, but sharing is--- introduced subsequently by common subexpression elimination and floating--- of array computations.+-- Collective array programs comprise closed expressions of array operations.+-- There is no explicit sharing in the initial AST form, but sharing is+-- introduced subsequently by common subexpression elimination and floating+-- of array computations. ----- Functions--- ~~~~~~~~~--- The array expression language is first-order and only provides limited--- control structures to ensure that it can be efficiently executed on--- compute-acceleration hardware, such as GPUs. To restrict functions to--- first-order, we separate function abstraction from the main expression--- type. Functions are represented using de Bruijn indices.+-- /Functions/ ----- Parametric and ad-hoc polymorphism--- ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~--- The array language features paramatric polymophism (e.g., pairing and--- projections) as well as ad-hoc polymorphism (e.g., arithmetic operations).--- All ad-hoc polymorphic constructs include reified dictionaries (c.f.,--- module 'Types'). Reified dictionaries also ensure that constants--- (constructor 'Const') are representable on compute acceleration hardware.+-- The array expression language is first-order and only provides limited+-- control structures to ensure that it can be efficiently executed on+-- compute-acceleration hardware, such as GPUs. To restrict functions to+-- first-order, we separate function abstraction from the main expression+-- type. Functions are represented using de Bruijn indices. ----- The AST contains both reified dictionaries and type class constraints. --- Type classes are used for array-related functionality that is uniformly--- available for all supported types. In contrast, reified dictionaries are--- used for functionality that is only available for certain types, such as--- arithmetic operations.+-- /Parametric and ad-hoc polymorphism/+--+-- The array language features paramatric polymophism (e.g., pairing and+-- projections) as well as ad-hoc polymorphism (e.g., arithmetic operations).+-- All ad-hoc polymorphic constructs include reified dictionaries (c.f.,+-- module 'Types'). Reified dictionaries also ensure that constants+-- (constructor 'Const') are representable on compute acceleration hardware.+--+-- The AST contains both reified dictionaries and type class constraints. +-- Type classes are used for array-related functionality that is uniformly+-- available for all supported types. In contrast, reified dictionaries are+-- used for functionality that is only available for certain types, such as+-- arithmetic operations. module Data.Array.Accelerate.AST ( @@ -70,7 +71,7 @@ Val(..), prj, -- * Accelerated array expressions- OpenAcc(..), Acc,+ OpenAcc(..), Acc, Stencil(..), -- * Scalar expressions OpenFun(..), Fun, OpenExp(..), Exp, PrimConst(..), PrimFun(..)@@ -154,7 +155,7 @@ -> OpenAcc aenv (Array dim' e') -- Variable bound by a 'Let', represented by a de Bruijn index - Avar :: Elem e+ Avar :: (Ix dim, Elem e) => Idx aenv (Array dim e) -> OpenAcc aenv (Array dim e) @@ -253,7 +254,7 @@ -- -- The target array is initialised from an array of default values (in case -- some positions in the target array are never picked by the permutation- -- functions). Moroever, we have a combination function (in case some+ -- functions). Moreover, we have a combination function (in case some -- positions on the target array are picked multiple times by the -- permutation functions). The combination function needs to be -- /associative/ and /commutative/ . We drop every element for which the@@ -272,10 +273,163 @@ -> OpenAcc aenv (Array dim e) -- source array -> OpenAcc aenv (Array dim' e) + -- Map a stencil over an array. In contrast to 'map', the domain of a stencil function is an+ -- entire /neighbourhood/ of each array element.+ Stencil :: (Elem e, Elem e', Stencil dim e stencil)+ => Fun aenv (stencil -> e') -- stencil function+ -> Boundary (ElemRepr e) -- boundary condition+ -> OpenAcc aenv (Array dim e) -- source array+ -> OpenAcc aenv (Array dim e')++ -- Map a binary stencil over an array.+ Stencil2 :: (Elem e1, Elem e2, Elem e', + Stencil dim e1 stencil1,+ Stencil dim e2 stencil2)+ => Fun aenv (stencil1 -> + stencil2 -> e') -- stencil function+ -> Boundary (ElemRepr e1) -- boundary condition #1+ -> OpenAcc aenv (Array dim e1) -- source array #1+ -> Boundary (ElemRepr e2) -- boundary condition #2+ -> OpenAcc aenv (Array dim e2) -- source array #2+ -> OpenAcc aenv (Array dim e')+ -- |Closed array expression aka an array program -- type Acc a = OpenAcc () a +class IsTuple stencil => Stencil dim e stencil where+ stencilAccess :: (dim -> e) -> dim -> stencil++-- DIM1+instance Stencil DIM1 a (a, a, a) where+ stencilAccess rf ix = (rf (ix - 1), rf ix, rf (ix + 1))+instance Stencil DIM1 a (a, a, a, a, a) where+ stencilAccess rf ix = (rf (ix - 2), rf (ix - 1), rf ix, rf (ix + 1), rf (ix + 2))+instance Stencil DIM1 a (a, a, a, a, a, a, a) where+ stencilAccess rf ix = (rf (ix - 3), rf (ix - 2), rf (ix - 1), rf ix, + rf (ix + 1), rf (ix + 2), rf (ix + 3))+instance Stencil DIM1 a (a, a, a, a, a, a, a, a, a) where+ stencilAccess rf ix = (rf (ix - 4), rf (ix - 3), rf (ix - 2), rf (ix - 1), rf ix, + rf (ix + 1), rf (ix + 2), rf (ix + 3), rf (ix + 4))++-- DIM2+instance (Stencil DIM1 a row2, + Stencil DIM1 a row1,+ Stencil DIM1 a row0) => Stencil DIM2 a (row2, row1, row0) where+ stencilAccess rf (x, y) = (stencilAccess (rf' (y - 1)) x, + stencilAccess (rf' y ) x,+ stencilAccess (rf' (y + 1)) x)+ where+ rf' y x = rf (x, y)+instance (Stencil DIM1 a row1,+ Stencil DIM1 a row2,+ Stencil DIM1 a row3,+ Stencil DIM1 a row4,+ Stencil DIM1 a row5) => Stencil DIM2 a (row1, row2, row3, row4, row5) where+ stencilAccess rf (x, y) = (stencilAccess (rf' (y - 2)) x, + stencilAccess (rf' (y - 1)) x, + stencilAccess (rf' y ) x,+ stencilAccess (rf' (y + 1)) x,+ stencilAccess (rf' (y + 2)) x)+ where+ rf' y x = rf (x, y)+instance (Stencil DIM1 a row1,+ Stencil DIM1 a row2,+ Stencil DIM1 a row3,+ Stencil DIM1 a row4,+ Stencil DIM1 a row5,+ Stencil DIM1 a row6,+ Stencil DIM1 a row7) => Stencil DIM2 a (row1, row2, row3, row4, row5, row6, row7) where+ stencilAccess rf (x, y) = (stencilAccess (rf' (y - 3)) x, + stencilAccess (rf' (y - 2)) x, + stencilAccess (rf' (y - 1)) x, + stencilAccess (rf' y ) x,+ stencilAccess (rf' (y + 1)) x,+ stencilAccess (rf' (y + 2)) x,+ stencilAccess (rf' (y + 3)) x)+ where+ rf' y x = rf (x, y)+instance (Stencil DIM1 a row1,+ Stencil DIM1 a row2,+ Stencil DIM1 a row3,+ Stencil DIM1 a row4,+ Stencil DIM1 a row5,+ Stencil DIM1 a row6,+ Stencil DIM1 a row7,+ Stencil DIM1 a row8,+ Stencil DIM1 a row9) + => Stencil DIM2 a (row1, row2, row3, row4, row5, row6, row7, row8, row9) where+ stencilAccess rf (x, y) = (stencilAccess (rf' (y - 4)) x, + stencilAccess (rf' (y - 3)) x, + stencilAccess (rf' (y - 2)) x, + stencilAccess (rf' (y - 1)) x, + stencilAccess (rf' y ) x,+ stencilAccess (rf' (y + 1)) x,+ stencilAccess (rf' (y + 2)) x,+ stencilAccess (rf' (y + 3)) x,+ stencilAccess (rf' (y + 4)) x)+ where+ rf' y x = rf (x, y)++-- DIM3+instance (Stencil DIM2 a row1, + Stencil DIM2 a row2,+ Stencil DIM2 a row3) => Stencil DIM3 a (row1, row2, row3) where+ stencilAccess rf (x, y, z) = (stencilAccess (rf' (z - 1)) (x, y), + stencilAccess (rf' z ) (x, y),+ stencilAccess (rf' (z + 1)) (x, y))+ where+ rf' z (x, y) = rf (x, y, z)+instance (Stencil DIM2 a row1,+ Stencil DIM2 a row2,+ Stencil DIM2 a row3,+ Stencil DIM2 a row4,+ Stencil DIM2 a row5) => Stencil DIM3 a (row1, row2, row3, row4, row5) where+ stencilAccess rf (x, y, z) = (stencilAccess (rf' (z - 2)) (x, y), + stencilAccess (rf' (z - 1)) (x, y), + stencilAccess (rf' z ) (x, y),+ stencilAccess (rf' (z + 1)) (x, y),+ stencilAccess (rf' (z + 2)) (x, y))+ where+ rf' z (x, y) = rf (x, y, z)+instance (Stencil DIM2 a row1,+ Stencil DIM2 a row2,+ Stencil DIM2 a row3,+ Stencil DIM2 a row4,+ Stencil DIM2 a row5,+ Stencil DIM2 a row6,+ Stencil DIM2 a row7) => Stencil DIM3 a (row1, row2, row3, row4, row5, row6, row7) where+ stencilAccess rf (x, y, z) = (stencilAccess (rf' (z - 3)) (x, y), + stencilAccess (rf' (z - 2)) (x, y), + stencilAccess (rf' (z - 1)) (x, y), + stencilAccess (rf' z ) (x, y),+ stencilAccess (rf' (z + 1)) (x, y),+ stencilAccess (rf' (z + 2)) (x, y),+ stencilAccess (rf' (z + 3)) (x, y))+ where+ rf' z (x, y) = rf (x, y, z)+instance (Stencil DIM2 a row1,+ Stencil DIM2 a row2,+ Stencil DIM2 a row3,+ Stencil DIM2 a row4,+ Stencil DIM2 a row5,+ Stencil DIM2 a row6,+ Stencil DIM2 a row7,+ Stencil DIM2 a row8,+ Stencil DIM2 a row9) + => Stencil DIM3 a (row1, row2, row3, row4, row5, row6, row7, row8, row9) where+ stencilAccess rf (x, y, z) = (stencilAccess (rf' (z - 4)) (x, y), + stencilAccess (rf' (z - 3)) (x, y), + stencilAccess (rf' (z - 2)) (x, y), + stencilAccess (rf' (z - 1)) (x, y), + stencilAccess (rf' z ) (x, y),+ stencilAccess (rf' (z + 1)) (x, y),+ stencilAccess (rf' (z + 2)) (x, y),+ stencilAccess (rf' (z + 3)) (x, y),+ stencilAccess (rf' (z + 4)) (x, y))+ where+ rf' z (x, y) = rf (x, y, z)+ -- Embedded expressions -- --------------------@@ -406,6 +560,7 @@ PrimLog :: FloatingType a -> PrimFun (a -> a) PrimFPow :: FloatingType a -> PrimFun ((a,a) -> a) PrimLogBase :: FloatingType a -> PrimFun ((a,a) -> a)+ PrimAtan2 :: FloatingType a -> PrimFun ((a,a) -> a) -- FIXME: add operations from Floating, RealFrac & RealFloat -- relational and equality operators
Data/Array/Accelerate/Analysis/Type.hs view
@@ -17,7 +17,7 @@ module Data.Array.Accelerate.Analysis.Type ( -- * Query AST types- accType, accType2, expType, sizeOf+ accType, accType2, accShapeType, accShapeType2, expType, sizeOf ) where @@ -52,6 +52,8 @@ accType (FoldSeg _ _ acc _) = accType acc accType (Permute _ _ _ acc) = accType acc accType (Backpermute _ _ acc) = accType acc+accType (Stencil _ _ _) = elemType (undefined::e)+accType (Stencil2 _ _ _ _ _) = elemType (undefined::e) -- |Reify the element types of the results of an array computation that yields -- two arrays.@@ -60,6 +62,35 @@ -> (TupleType (ElemRepr e1), TupleType (ElemRepr e2)) accType2 (Scanl _ e acc) = (accType acc, expType e) accType2 (Scanr _ e acc) = (accType acc, expType e)++-- |Reify the shape type of the result of an array computation+--+accShapeType :: forall aenv dim e.+ OpenAcc aenv (Array dim e) -> TupleType (ElemRepr dim)+accShapeType (Let _ acc) = accShapeType acc+accShapeType (Let2 _ acc) = accShapeType acc+accShapeType (Avar _) = elemType (undefined::dim)+accShapeType (Use (Array _ _)) = elemType (undefined::dim)+accShapeType (Unit _) = elemType (undefined::DIM0)+accShapeType (Reshape _ _) = elemType (undefined::dim)+accShapeType (Replicate _ _ _) = elemType (undefined::dim)+accShapeType (Index _ _ _) = elemType (undefined::dim)+accShapeType (Map _ acc) = accShapeType acc+accShapeType (ZipWith _ _ acc) = accShapeType acc+accShapeType (Fold _ _ _) = elemType (undefined::DIM0)+accShapeType (FoldSeg _ _ _ _) = elemType (undefined::DIM1)+accShapeType (Permute _ acc _ _) = accShapeType acc+accShapeType (Backpermute _ _ _) = elemType (undefined::dim)+accShapeType (Stencil _ _ acc) = accShapeType acc+accShapeType (Stencil2 _ _ acc _ _) = accShapeType acc++-- |Reify the shape type of the results of a computation that yields two arrays+--+accShapeType2 :: OpenAcc aenv (Array dim1 e1, Array dim2 e2)+ -> (TupleType (ElemRepr dim1), TupleType (ElemRepr dim2))+accShapeType2 (Scanl _ _ acc) = (accShapeType acc, elemType (undefined::()))+accShapeType2 (Scanr _ _ acc) = (accShapeType acc, elemType (undefined::()))+ -- |Reify the result type of a scalar expression. --
Data/Array/Accelerate/Array/Representation.hs view
@@ -16,9 +16,12 @@ ) where +-- friends +import Data.Array.Accelerate.Type + -- |Index representation--- -+-- -- |Class of index representations (which are nested pairs) --@@ -31,8 +34,10 @@ index :: ix -> ix -> Int -- yield the index position in a linear, -- row-major representation of the array -- (first argument is the shape)+ bound :: ix -> ix -> Boundary e -> Either e ix+ -- apply a boundary condition to an index - iter :: ix -> (ix -> a) -> (a -> a -> a) -> a -> a+ iter :: ix -> (ix -> a) -> (a -> a -> a) -> a -> a -- iterate through the entire shape, applying -- the function; third argument combines results -- and fourth is returned in case of an empty@@ -43,18 +48,28 @@ rangeToShape :: (ix, ix) -> ix -- convert a minpoint-maxpoint index -- into a shape shapeToRange :: ix -> (ix, ix) -- ...the converse+ + -- other conversions+ shapeToList :: ix -> [Int] -- convert a shape into its list of dimensions+ listToShape :: [Int] -> ix -- convert a list of dimensions into a shape+ instance Ix () where- dim () = 0- size () = 1- intersect () () = ()- ignore = ()- index () () = 0- iter () f _ _ = f ()+ dim () = 0+ size () = 1+ () `intersect` () = ()+ ignore = ()+ index () () = 0+ bound () () _ = Right ()+ iter () f _ _ = f () rangeToShape ((), ()) = () shapeToRange () = ((), ()) + shapeToList () = []+ listToShape [] = ()+ listToShape _ = error "Data.Array.Accelerate.Array: non-empty list when converting to unit"+ instance Ix ix => Ix (ix, Int) where dim (sh, _) = dim sh + 1 size (sh, sz) = size sh * sz@@ -64,6 +79,21 @@ | i >= 0 && i < sz = index sh ix + size sh * i | otherwise = error "Data.Array.Accelerate.Array: index out of bounds"+ bound (sh, sz) (ix, i) bndy+ | i < 0 = case bndy of+ Clamp -> bound sh ix bndy `addDim` 0+ Mirror -> bound sh ix bndy `addDim` (-i)+ Wrap -> bound sh ix bndy `addDim` (sz-i)+ Constant e -> Left e+ | i >= sz = case bndy of+ Clamp -> bound sh ix bndy `addDim` (sz-1)+ Mirror -> bound sh ix bndy `addDim` (sz-(i-sz+1))+ Wrap -> bound sh ix bndy `addDim` (i-sz)+ Constant e -> Left e+ | otherwise = bound sh ix bndy `addDim` i+ where+ Right ix `addDim` i = Right (ix, i)+ Left e `addDim` _ = Left e iter (sh, sz) f c r = iter' 0 where iter' i | i >= sz = r@@ -76,9 +106,13 @@ in ((low, 0), (high, sz - 1)) + shapeToList (sh,sz) = sz : shapeToList sh+ listToShape [] = error "Data.Array.Accelerate.Array: empty list when converting to Ix"+ listToShape (x:xs) = (listToShape xs,x) + -- |Slice representation--- -+-- -- |Class of slice representations (which are nested pairs) --
Data/Array/Accelerate/Array/Sugar.hs view
@@ -23,7 +23,7 @@ liftToElem, liftToElem2, sinkFromElem, sinkFromElem2, -- * Array shapes- DIM0, DIM1, DIM2, DIM3, DIM4, DIM5,+ DIM0, DIM1, DIM2, DIM3, DIM4, DIM5, DIM6, DIM7, DIM8, DIM9, -- * Array indexing and slicing ShapeBase, Shape, Ix(..), All(..), SliceIx(..), convertSliceIndex,@@ -94,6 +94,11 @@ type instance ElemRepr (a, b, c) = (ElemRepr (a, b), ElemRepr' c) type instance ElemRepr (a, b, c, d) = (ElemRepr (a, b, c), ElemRepr' d) type instance ElemRepr (a, b, c, d, e) = (ElemRepr (a, b, c, d), ElemRepr' e)+type instance ElemRepr (a, b, c, d, e, f) = (ElemRepr (a, b, c, d, e), ElemRepr' f)+type instance ElemRepr (a, b, c, d, e, f, g) = (ElemRepr (a, b, c, d, e, f), ElemRepr' g)+type instance ElemRepr (a, b, c, d, e, f, g, h) = (ElemRepr (a, b, c, d, e, f, g), ElemRepr' h)+type instance ElemRepr (a, b, c, d, e, f, g, h, i) + = (ElemRepr (a, b, c, d, e, f, g, h), ElemRepr' i) -- To avoid overly nested pairs, we use a flattened representation at the -- leaves.@@ -132,6 +137,11 @@ type instance ElemRepr' (a, b, c) = (ElemRepr (a, b), ElemRepr' c) type instance ElemRepr' (a, b, c, d) = (ElemRepr (a, b, c), ElemRepr' d) type instance ElemRepr' (a, b, c, d, e) = (ElemRepr (a, b, c, d), ElemRepr' e)+type instance ElemRepr' (a, b, c, d, e, f) = (ElemRepr (a, b, c, d, e), ElemRepr' f)+type instance ElemRepr' (a, b, c, d, e, f, g) = (ElemRepr (a, b, c, d, e, f), ElemRepr' g)+type instance ElemRepr' (a, b, c, d, e, f, g, h) = (ElemRepr (a, b, c, d, e, f, g), ElemRepr' h)+type instance ElemRepr' (a, b, c, d, e, f, g, h, i) + = (ElemRepr (a, b, c, d, e, f, g, h), ElemRepr' i) -- Array elements (tuples of scalars)@@ -471,8 +481,67 @@ fromElem' (a, b, c, d, e) = (fromElem (a, b, c, d), fromElem' e) toElem' (abcd, e) = let (a, b, c, d) = toElem abcd in (a, b, c, d, toElem' e) +instance (Elem a, Elem b, Elem c, Elem d, Elem e, Elem f) => Elem (a, b, c, d, e, f) where+ elemType (_::(a, b, c, d, e, f)) + = PairTuple (elemType (undefined :: (a, b, c, d, e))) + (elemType' (undefined :: f))+ fromElem (a, b, c, d, e, f) = (fromElem (a, b, c, d, e), fromElem' f)+ toElem (abcde, f) = let (a, b, c, d, e) = toElem abcde in (a, b, c, d, e, toElem' f)++ elemType' (_::(a, b, c, d, e, f)) + = PairTuple (elemType (undefined :: (a, b, c, d, e))) + (elemType' (undefined :: f))+ fromElem' (a, b, c, d, e, f) = (fromElem (a, b, c, d, e), fromElem' f)+ toElem' (abcde, f) = let (a, b, c, d, e) = toElem abcde in (a, b, c, d, e, toElem' f)++instance (Elem a, Elem b, Elem c, Elem d, Elem e, Elem f, Elem g) + => Elem (a, b, c, d, e, f, g) where+ elemType (_::(a, b, c, d, e, f, g)) + = PairTuple (elemType (undefined :: (a, b, c, d, e, f))) + (elemType' (undefined :: g))+ fromElem (a, b, c, d, e, f, g) = (fromElem (a, b, c, d, e, f), fromElem' g)+ toElem (abcdef, g) = let (a, b, c, d, e, f) = toElem abcdef in (a, b, c, d, e, f, toElem' g)++ elemType' (_::(a, b, c, d, e, f, g)) + = PairTuple (elemType (undefined :: (a, b, c, d, e, f))) + (elemType' (undefined :: g))+ fromElem' (a, b, c, d, e, f, g) = (fromElem (a, b, c, d, e, f), fromElem' g)+ toElem' (abcdef, g) = let (a, b, c, d, e, f) = toElem abcdef in (a, b, c, d, e, f, toElem' g)++instance (Elem a, Elem b, Elem c, Elem d, Elem e, Elem f, Elem g, Elem h) + => Elem (a, b, c, d, e, f, g, h) where+ elemType (_::(a, b, c, d, e, f, g, h)) + = PairTuple (elemType (undefined :: (a, b, c, d, e, f, g))) + (elemType' (undefined :: h))+ fromElem (a, b, c, d, e, f, g, h) = (fromElem (a, b, c, d, e, f, g), fromElem' h)+ toElem (abcdefg, h) = let (a, b, c, d, e, f, g) = toElem abcdefg + in (a, b, c, d, e, f, g, toElem' h)++ elemType' (_::(a, b, c, d, e, f, g, h)) + = PairTuple (elemType (undefined :: (a, b, c, d, e, f, g))) + (elemType' (undefined :: h))+ fromElem' (a, b, c, d, e, f, g, h) = (fromElem (a, b, c, d, e, f, g), fromElem' h)+ toElem' (abcdefg, h) = let (a, b, c, d, e, f, g) = toElem abcdefg + in (a, b, c, d, e, f, g, toElem' h)++instance (Elem a, Elem b, Elem c, Elem d, Elem e, Elem f, Elem g, Elem h, Elem i) + => Elem (a, b, c, d, e, f, g, h, i) where+ elemType (_::(a, b, c, d, e, f, g, h, i)) + = PairTuple (elemType (undefined :: (a, b, c, d, e, f, g, h))) + (elemType' (undefined :: i))+ fromElem (a, b, c, d, e, f, g, h, i) = (fromElem (a, b, c, d, e, f, g, h), fromElem' i)+ toElem (abcdefgh, i) = let (a, b, c, d, e, f, g, h) = toElem abcdefgh+ in (a, b, c, d, e, f, g, h, toElem' i)++ elemType' (_::(a, b, c, d, e, f, g, h, i)) + = PairTuple (elemType (undefined :: (a, b, c, d, e, f, g, h))) + (elemType' (undefined :: i))+ fromElem' (a, b, c, d, e, f, g, h, i) = (fromElem (a, b, c, d, e, f, g, h), fromElem' i)+ toElem' (abcdefgh, i) = let (a, b, c, d, e, f, g, h) = toElem abcdefgh+ in (a, b, c, d, e, f, g, h, toElem' i)+ -- |Convenience functions--- -+-- singletonScalarType :: IsScalar a => a -> TupleType ((), a) singletonScalarType _ = PairTuple UnitTuple (SingleTuple scalarType)@@ -548,6 +617,10 @@ type DIM3 = (Int, Int, Int) type DIM4 = (Int, Int, Int, Int) type DIM5 = (Int, Int, Int, Int, Int)+type DIM6 = (Int, Int, Int, Int, Int, Int)+type DIM7 = (Int, Int, Int, Int, Int, Int, Int)+type DIM8 = (Int, Int, Int, Int, Int, Int, Int, Int)+type DIM9 = (Int, Int, Int, Int, Int, Int, Int, Int, Int) -- Shape constraints and indexing -- @@ -569,6 +642,17 @@ => Shape (a, b, c, d) instance (ShapeBase a, ShapeBase b, ShapeBase c, ShapeBase d, ShapeBase e) => Shape (a, b, c, d, e)+instance (ShapeBase a, ShapeBase b, ShapeBase c, ShapeBase d, ShapeBase e, ShapeBase f) + => Shape (a, b, c, d, e, f)+instance (ShapeBase a, ShapeBase b, ShapeBase c, ShapeBase d, ShapeBase e, ShapeBase f, + ShapeBase g) + => Shape (a, b, c, d, e, f, g)+instance (ShapeBase a, ShapeBase b, ShapeBase c, ShapeBase d, ShapeBase e, ShapeBase f, + ShapeBase g, ShapeBase h) + => Shape (a, b, c, d, e, f, g, h)+instance (ShapeBase a, ShapeBase b, ShapeBase c, ShapeBase d, ShapeBase e, ShapeBase f, + ShapeBase g, ShapeBase h, ShapeBase i) + => Shape (a, b, c, d, e, f, g, h, i) type family FromShapeBase shb :: * type instance FromShapeBase Int = Int@@ -585,6 +669,19 @@ type instance FromShapeRepr ((((((), a), b), c), d), e) = (FromShapeBase a, FromShapeBase b, FromShapeBase c, FromShapeBase d, FromShapeBase e)+type instance FromShapeRepr (((((((), a), b), c), d), e), f) + = (FromShapeBase a, FromShapeBase b, FromShapeBase c, FromShapeBase d, + FromShapeBase e, FromShapeBase f)+type instance FromShapeRepr ((((((((), a), b), c), d), e), f), g) + = (FromShapeBase a, FromShapeBase b, FromShapeBase c, FromShapeBase d, + FromShapeBase e, FromShapeBase f, FromShapeBase g)+type instance FromShapeRepr (((((((((), a), b), c), d), e), f), g), h) + = (FromShapeBase a, FromShapeBase b, FromShapeBase c, FromShapeBase d, + FromShapeBase e, FromShapeBase f, FromShapeBase g, FromShapeBase h)+type instance FromShapeRepr ((((((((((), a), b), c), d), e), f), g), h), i) + = (FromShapeBase a, FromShapeBase b, FromShapeBase c, FromShapeBase d, + FromShapeBase e, FromShapeBase f, FromShapeBase g, FromShapeBase h,+ FromShapeBase i) -- |Shapes and indices of multi-dimensional arrays --@@ -596,7 +693,7 @@ -- Total number of elements in an array of the given /shape/ size :: ix -> Int - -- |Magic value identifing elements ignored in 'permute'+ -- |Magic value identifying elements ignored in 'permute' ignore :: ix -- |Map a multi-dimensional index into one in a linear, row-major @@ -604,6 +701,9 @@ -- argument is the index) index :: ix -> ix -> Int + -- |Apply a boundary condition to an index+ bound :: ix -> ix -> Boundary a -> Either a ix+ -- |Iterate through the entire shape, applying the function; third argument -- combines results and fourth is returned in case of an empty iteration -- space; the index space is traversed in row-major order@@ -615,13 +715,22 @@ -- |Convert a /shape/ into a minpoint-maxpoint index shapeToRange :: ix -> (ix, ix) - dim = Repr.dim . fromElem- size = Repr.size . fromElem- ignore = toElem Repr.ignore- index sh ix = Repr.index (fromElem sh) (fromElem ix)- + -- |Convert a shape to a list of dimensions+ shapeToList :: ix -> [Int]++ -- |Convert a list of dimensions into a shape+ listToShape :: [Int] -> ix++ dim = Repr.dim . fromElem+ size = Repr.size . fromElem+ ignore = toElem Repr.ignore+ index sh ix = Repr.index (fromElem sh) (fromElem ix)+ bound sh ix bndy = case Repr.bound (fromElem sh) (fromElem ix) bndy of+ Left v -> Left v+ Right ix' -> Right $ toElem ix'+ iter sh f c r = Repr.iter (fromElem sh) (f . toElem) c r- + rangeToShape (low, high) = toElem (Repr.rangeToShape (fromElem low, fromElem high)) shapeToRange ix@@ -629,12 +738,20 @@ in (toElem low, toElem high) + shapeToList = Repr.shapeToList . fromElem+ listToShape = toElem . Repr.listToShape+ instance Ix () instance Ix (Int) instance Ix (Int, Int) instance Ix (Int, Int, Int) instance Ix (Int, Int, Int, Int) instance Ix (Int, Int, Int, Int, Int)+instance Ix (Int, Int, Int, Int, Int, Int)+instance Ix (Int, Int, Int, Int, Int, Int, Int)+instance Ix (Int, Int, Int, Int, Int, Int, Int, Int)+instance Ix (Int, Int, Int, Int, Int, Int, Int, Int, Int)+ -- |Slices -aka generalised indices- as n-tuples and mappings of slice -- indicies to slices, co-slices, and slice dimensions
Data/Array/Accelerate/CUDA.hs view
@@ -1,3 +1,4 @@+{-# LANGUAGE CPP #-} -- | -- Module : Data.Array.Accelerate.CUDA -- Copyright : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee, Trevor L. McDonell@@ -17,22 +18,31 @@ ) where +import Prelude hiding (catch)+import Control.Exception++import Foreign.CUDA.Driver.Error+ import Data.Array.Accelerate.AST import Data.Array.Accelerate.CUDA.State import Data.Array.Accelerate.CUDA.Compile import Data.Array.Accelerate.CUDA.Execute import Data.Array.Accelerate.CUDA.Array.Device-import qualified Data.Array.Accelerate.Smart as Sugar+import qualified Data.Array.Accelerate.CUDA.Smart as Sugar +#include "accelerate.h" + -- Accelerate: CUDA -- ~~~~~~~~~~~~~~~~ -- | Compiles and run a complete embedded array program using the CUDA backend -- run :: Arrays a => Sugar.Acc a -> IO a-run acc = evalCUDA- $ execute (Sugar.convertAcc acc) >>= collect+run acc =+ evalCUDA (execute (Sugar.convertAcc acc) >>= collect)+ `catch`+ \e -> INTERNAL_ERROR(error) "unhandled" (show (e :: CUDAException)) execute :: Arrays a => Acc a -> CIO a execute acc = compileAcc acc >> executeAcc acc
Data/Array/Accelerate/CUDA/Analysis/Hash.hs view
@@ -1,4 +1,4 @@-{-# LANGUAGE GADTs #-}+{-# LANGUAGE CPP, GADTs #-} -- | -- Module : Data.Array.Accelerate.CUDA.Analysis.Hash -- Copyright : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee, Trevor L. McDonell@@ -13,36 +13,47 @@ where import Data.Char+import Language.C+import Control.Monad.State+import Text.PrettyPrint+ import Data.Array.Accelerate.AST import Data.Array.Accelerate.Type import Data.Array.Accelerate.Pretty () import Data.Array.Accelerate.Analysis.Type-import Data.Array.Accelerate.CUDA.State+import Data.Array.Accelerate.CUDA.CodeGen +#include "accelerate.h" + -- | -- Generate a unique key for each kernel computation (not extensively tested...) ---accToKey :: OpenAcc aenv a -> Key-accToKey (Unit e) = chr 2 : showTy (expType e) ++ show e-accToKey (Reshape _ a) = chr 7 : showTy (accType a)-accToKey (Replicate _ e a) = chr 17 : showTy (accType a) ++ show e-accToKey (Index _ a e) = chr 29 : showTy (accType a) ++ show e-accToKey (Map f a) = chr 41 : showTy (accType a) ++ show f-accToKey (ZipWith f x y) = chr 53 : showTy (accType x) ++ showTy (accType y) ++ show f-accToKey (Fold f e a) = chr 67 : showTy (accType a) ++ show f ++ show e-accToKey (FoldSeg f e _ a) = chr 79 : showTy (accType a) ++ show f ++ show e-accToKey (Scanl f e a) = chr 97 : showTy (accType a) ++ show f ++ show e-accToKey (Scanr f e a) = chr 107 : showTy (accType a) ++ show f ++ show e-accToKey (Permute c e p a) = chr 127 : showTy (accType a) ++ show c ++ show e ++ show p-accToKey (Backpermute p _ a) = chr 139 : showTy (accType a) ++ show p-accToKey _ =- error "we should never get here"+accToKey :: OpenAcc aenv a -> String+accToKey (Replicate _ e a) = chr 17 : showTy (accType a) ++ showExp e+accToKey (Index _ a e) = chr 29 : showTy (accType a) ++ showExp e+accToKey (Map f a) = chr 41 : showTy (accType a) ++ showFun f+accToKey (ZipWith f x y) = chr 53 : showTy (accType x) ++ showTy (accType y) ++ showFun f+accToKey (Fold f e a) = chr 67 : showTy (accType a) ++ showFun f ++ showExp e+accToKey (FoldSeg f e _ a) = chr 79 : showTy (accType a) ++ showFun f ++ showExp e+accToKey (Scanl f e a) = chr 97 : showTy (accType a) ++ showFun f ++ showExp e+accToKey (Scanr f e a) = chr 107 : showTy (accType a) ++ showFun f ++ showExp e+accToKey (Permute c _ p a) = chr 127 : showTy (accType a) ++ showFun c ++ showFun p+accToKey (Backpermute _ p a) = chr 139 : showTy (accType a) ++ showFun p+accToKey x =+ INTERNAL_ERROR(error) "accToKey"+ (unlines ["incomplete patterns for key generation", render . nest 2 $ text (show x)]) showTy :: TupleType a -> String showTy UnitTuple = [] showTy (SingleTuple ty) = show ty showTy (PairTuple a b) = showTy a ++ showTy b++showFun :: OpenFun env aenv a -> String+showFun f = render . hcat . map pretty $ evalState (codeGenFun f) []++showExp :: OpenExp env aenv a -> String+showExp e = render . hcat . map pretty $ evalState (codeGenExp e) [] {- -- hash function from the dragon book pp437; assumes 7 bit characters and needs
Data/Array/Accelerate/CUDA/Analysis/Launch.hs view
@@ -14,12 +14,13 @@ import Control.Monad.IO.Class +import Data.Int import Data.Array.Accelerate.AST import Data.Array.Accelerate.Analysis.Type- import Data.Array.Accelerate.CUDA.State import qualified Foreign.CUDA.Analysis as CUDA import qualified Foreign.CUDA.Driver as CUDA+import qualified Foreign.Storable as F -- |@@ -35,16 +36,19 @@ -- TLM: this could probably be stored in the KernelEntry -- launchConfig :: OpenAcc aenv a -> Int -> CUDA.Fun -> CIO (Int, Int, Integer)+launchConfig acc@(Fold _ _ _) n _ =+ return (256, gridSize undefined acc n 256, toInteger (sharedMem undefined acc 256))+ launchConfig acc n fn = do regs <- liftIO $ CUDA.requires fn CUDA.NumRegs stat <- liftIO $ CUDA.requires fn CUDA.SharedSizeBytes -- static memory only prop <- getM deviceProps - let dyn = sharedMem acc+ let dyn = sharedMem prop acc (cta, occ) = CUDA.optimalBlockSize prop (const regs) ((stat+) . dyn) mbk = CUDA.multiProcessorCount prop * CUDA.activeThreadBlocks occ - return (cta, mbk `min` gridSize acc n cta, toInteger (dyn cta))+ return (cta, mbk `min` gridSize prop acc n cta, toInteger (dyn cta)) -- |@@ -52,8 +56,11 @@ -- given array expression. This should understand things like #elements per -- thread for the various kernels. ---gridSize :: OpenAcc aenv a -> Int -> Int -> Int-gridSize acc size cta =+-- foldSeg: 'size' is the number of segments, require one warp per segment+--+gridSize :: CUDA.DeviceProperties -> OpenAcc aenv a -> Int -> Int -> Int+gridSize p (FoldSeg _ _ _ _) size cta = ((size * CUDA.warpSize p) + cta - 1) `div` cta+gridSize _ acc size cta = let between arr n = (n+arr-1) `div` n in 1 `max` ((cta - 1 + (size `between` elementsPerThread acc)) `div` cta) @@ -66,9 +73,15 @@ -- memory usage as a function of thread block size. This can be used by the -- occupancy calculator to optimise kernel launch shape. ---sharedMem :: OpenAcc aenv a -> Int -> Int-sharedMem (Fold _ x _) t = sizeOf (expType x) * t-sharedMem (Scanl _ x _) t = sizeOf (expType x) * t-sharedMem (Scanr _ x _) t = sizeOf (expType x) * t-sharedMem _ _ = 0+sharedMem :: CUDA.DeviceProperties -> OpenAcc aenv a -> Int -> Int+sharedMem _ (Fold _ x _) blockDim = sizeOf (expType x) * blockDim+sharedMem _ (Scanl _ x _) blockDim = sizeOf (expType x) * blockDim+sharedMem _ (Scanr _ x _) blockDim = sizeOf (expType x) * blockDim+sharedMem p (FoldSeg _ x _ _) blockDim =+ let warp = CUDA.warpSize p+ in+ (blockDim `div` warp * 2) * F.sizeOf (undefined :: Int32) ++ (blockDim + warp `div` 2) * sizeOf (expType x)++sharedMem _ _ _ = 0
Data/Array/Accelerate/CUDA/Array/Data.hs view
@@ -15,7 +15,7 @@ ArrayElem(..) ) where -import Prelude hiding (id, (.), mod)+import Prelude hiding (id, (.)) import Control.Category import Foreign.Ptr@@ -28,7 +28,7 @@ import Control.Monad import Control.Applicative import Control.Monad.IO.Class-import qualified Data.IntMap as IM+import qualified Data.HashTable as HT import Data.Array.Accelerate.CUDA.State import qualified Data.Array.Accelerate.Array.Data as Acc@@ -36,7 +36,9 @@ import qualified Foreign.CUDA.Driver.Stream as CUDA import qualified Foreign.CUDA.Driver.Texture as CUDA +#include "accelerate.h" + -- Instances -- ~~~~~~~~~ @@ -203,14 +205,16 @@ -- Allocate a new device array to accompany the given host-side Accelerate array ---mallocArray' :: (Acc.ArrayPtrs e ~ Ptr a, Storable a, Acc.ArrayElem e) => Acc.ArrayData e -> Int -> CIO ()+mallocArray' :: forall a e. (Acc.ArrayPtrs e ~ Ptr a, Storable a, Acc.ArrayElem e) => Acc.ArrayData e -> Int -> CIO () mallocArray' ad n = do- exists <- IM.member key <$> getM memoryEntry- unless exists $ insertArray ad =<< liftIO (CUDA.mallocArray n)+ table <- getM memoryTable+ exists <- isJust <$> liftIO (HT.lookup table key)+ unless exists . liftIO $ insertArray ad table =<< CUDA.mallocArray n where- insertArray :: (Acc.ArrayPtrs e ~ Ptr a, Acc.ArrayElem e) => Acc.ArrayData e -> CUDA.DevicePtr a -> CIO ()- insertArray _ = modM memoryEntry . IM.insert key . MemoryEntry 0 . CUDA.devPtrToWordPtr- key = arrayToKey ad+ insertArray :: (Acc.ArrayPtrs e ~ Ptr a, Acc.ArrayElem e) => Acc.ArrayData e -> MemTable -> CUDA.DevicePtr a -> IO ()+ insertArray _ t = HT.insert t key . MemoryEntry 0 bytes . CUDA.devPtrToWordPtr+ bytes = fromIntegral $ n * sizeOf (undefined :: a)+ key = arrayToKey ad -- Array indexing@@ -239,14 +243,14 @@ peekArray' :: (Acc.ArrayPtrs e ~ Ptr a, Storable a, Acc.ArrayElem e) => Acc.ArrayData e -> Int -> CIO () peekArray' ad n = let dst = Acc.ptrsOfArrayData ad- src = CUDA.wordPtrToDevPtr . get arena+ src = CUDA.wordPtrToDevPtr . getL arena in lookupArray ad >>= \me -> liftIO $ CUDA.peekArray n (src me) dst peekArrayAsync' :: (Acc.ArrayPtrs e ~ Ptr a, Storable a, Acc.ArrayElem e) => Acc.ArrayData e -> Int -> Maybe CUDA.Stream -> CIO () peekArrayAsync' ad n st = let dst = CUDA.HostPtr . Acc.ptrsOfArrayData- src = CUDA.wordPtrToDevPtr . get arena+ src = CUDA.wordPtrToDevPtr . getL arena in lookupArray ad >>= \me -> liftIO $ CUDA.peekArrayAsync n (src me) (dst ad) st @@ -258,44 +262,49 @@ pokeArray' :: (Acc.ArrayPtrs e ~ Ptr a, Storable a, Acc.ArrayElem e) => Acc.ArrayData e -> Int -> CIO () pokeArray' ad n = let src = Acc.ptrsOfArrayData ad- dst = CUDA.wordPtrToDevPtr . get arena+ dst = CUDA.wordPtrToDevPtr . getL arena in do- me <- mod refcount (+1) <$> lookupArray ad- when (get refcount me <= 1) . liftIO $ CUDA.pokeArray n src (dst me)- modM memoryEntry (IM.insert (arrayToKey ad) me)+ tab <- getM memoryTable+ me <- modL refcount (+1) <$> lookupArray ad+ when (getL refcount me <= 1) . liftIO $ CUDA.pokeArray n src (dst me)+ liftIO $ HT.insert tab (arrayToKey ad) me pokeArrayAsync' :: (Acc.ArrayPtrs e ~ Ptr a, Storable a, Acc.ArrayElem e) => Acc.ArrayData e -> Int -> Maybe CUDA.Stream -> CIO () pokeArrayAsync' ad n st = let src = CUDA.HostPtr . Acc.ptrsOfArrayData- dst = CUDA.wordPtrToDevPtr . get arena+ dst = CUDA.wordPtrToDevPtr . getL arena in do- me <- mod refcount (+1) <$> lookupArray ad- when (get refcount me <= 1) . liftIO $ CUDA.pokeArrayAsync n (src ad) (dst me) st- modM memoryEntry (IM.insert (arrayToKey ad) me)+ tab <- getM memoryTable+ me <- modL refcount (+1) <$> lookupArray ad+ when (getL refcount me <= 1) . liftIO $ CUDA.pokeArrayAsync n (src ad) (dst me) st+ liftIO $ HT.insert tab (arrayToKey ad) me -- Release a device array, when its reference counter drops to zero -- freeArray' :: (Acc.ArrayPtrs e ~ Ptr a, Acc.ArrayElem e) => Acc.ArrayData e -> CIO () freeArray' ad = do- me <- mod refcount (subtract 1) <$> lookupArray ad- if get refcount me > 0- then modM memoryEntry (IM.insert (arrayToKey ad) me)- else do liftIO . CUDA.free . CUDA.wordPtrToDevPtr $ get arena me- modM memoryEntry (IM.delete (arrayToKey ad))+ tab <- getM memoryTable+ me <- modL refcount (subtract 1) <$> lookupArray ad+ liftIO $ if getL refcount me > 0+ then HT.insert tab (arrayToKey ad) me+ else do CUDA.free . CUDA.wordPtrToDevPtr $ getL arena me+ HT.delete tab (arrayToKey ad) -- Increase the reference count of an array. You may wish to call this right -- before another method that will implicitly attempt to release the array. -- touchArray' :: (Acc.ArrayPtrs e ~ Ptr a, Acc.ArrayElem e) => Acc.ArrayData e -> CIO ()-touchArray' ad = modM memoryEntry =<< IM.insert (arrayToKey ad) . mod refcount (+1) <$> lookupArray ad+touchArray' ad = do+ tab <- getM memoryTable+ liftIO . HT.insert tab (arrayToKey ad) . modL refcount (+1) =<< lookupArray ad -- Return the device pointers wrapped in a list of function parameters -- devicePtrs' :: (Acc.ArrayPtrs e ~ Ptr a, Acc.ArrayElem e) => Acc.ArrayData e -> CIO [CUDA.FunParam]-devicePtrs' ad = (: []) . CUDA.VArg . CUDA.wordPtrToDevPtr . get arena <$> lookupArray ad+devicePtrs' ad = (: []) . CUDA.VArg . CUDA.wordPtrToDevPtr . getL arena <$> lookupArray ad -- Retrieve texture references from the module (beginning with the given seed),@@ -316,21 +325,26 @@ -- Generate a memory map key from the given ArrayData ---arrayToKey :: (Acc.ArrayPtrs e ~ Ptr a, Acc.ArrayElem e) => Acc.ArrayData e -> IM.Key-arrayToKey = fromIntegral . ptrToWordPtr . Acc.ptrsOfArrayData+arrayToKey :: (Acc.ArrayPtrs e ~ Ptr a, Acc.ArrayElem e) => Acc.ArrayData e -> WordPtr+arrayToKey = ptrToWordPtr . Acc.ptrsOfArrayData {-# INLINE arrayToKey #-} -- Retrieve the device memory entry from the state structure associated with a -- particular Accelerate array. -- lookupArray :: (Acc.ArrayPtrs e ~ Ptr a, Acc.ArrayElem e) => Acc.ArrayData e -> CIO MemoryEntry-lookupArray ad = fromMaybe (error "ArrayElem: internal error") . IM.lookup (arrayToKey ad) <$> getM memoryEntry {-# INLINE lookupArray #-}+lookupArray ad = do+ t <- getM memoryTable+ x <- liftIO $ HT.lookup t (arrayToKey ad)+ case x of+ Just e -> return e+ _ -> INTERNAL_ERROR(error) "lookupArray" "lost device memory reference" -- Return the device pointer associated with a host-side Accelerate array -- getArray :: (Acc.ArrayPtrs e ~ Ptr a, Acc.ArrayElem e) => Acc.ArrayData e -> CIO (CUDA.DevicePtr a)-getArray ad = CUDA.wordPtrToDevPtr . get arena <$> lookupArray ad+getArray ad = CUDA.wordPtrToDevPtr . getL arena <$> lookupArray ad {-# INLINE getArray #-} -- Array tuple extraction
Data/Array/Accelerate/CUDA/CodeGen.hs view
@@ -1,4 +1,4 @@-{-# LANGUAGE GADTs, PatternGuards, ExistentialQuantification, ScopedTypeVariables #-}+{-# LANGUAGE CPP, GADTs, PatternGuards, ScopedTypeVariables #-} -- | -- Module : Data.Array.Accelerate.CUDA.CodeGen -- Copyright : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee, Trevor L. McDonell@@ -9,7 +9,11 @@ -- Portability : non-portable (GHC extensions) -- -module Data.Array.Accelerate.CUDA.CodeGen (codeGenAcc, CUTranslSkel)+module Data.Array.Accelerate.CUDA.CodeGen+ (+ CUTranslSkel,+ codeGenAcc, codeGenFun, codeGenExp+ ) where import Prelude hiding (mod)@@ -18,11 +22,13 @@ import Language.C import Control.Applicative import Control.Monad.State+import Text.PrettyPrint import Data.Array.Accelerate.Type import Data.Array.Accelerate.Tuple import Data.Array.Accelerate.Pretty () import Data.Array.Accelerate.Analysis.Type+import Data.Array.Accelerate.Array.Representation import qualified Data.Array.Accelerate.AST as AST import qualified Data.Array.Accelerate.Array.Sugar as Sugar import qualified Foreign.Storable as F@@ -31,7 +37,10 @@ import Data.Array.Accelerate.CUDA.CodeGen.Util import Data.Array.Accelerate.CUDA.CodeGen.Skeleton +#include "accelerate.h" +type CG a = State [CExtDecl] a+ -- Array expressions -- ~~~~~~~~~~~~~~~~~ @@ -52,24 +61,38 @@ -- -- FRAGILE. ---type CG a = State [CExtDecl] a-+-- TLM 2010-07-16:+-- Make AST.Acc an instance of Traversable, and use mapAccumL to collect the+-- types of free array variables ??+-- codeGenAcc' :: AST.OpenAcc aenv a -> CG CUTranslSkel-codeGenAcc' op@(AST.Map f1 a1) = mkMap (codeGenAccType op) (codeGenAccType a1) <$> codeGenFun f1-codeGenAcc' op@(AST.ZipWith f1 a1 a2) = mkZipWith (codeGenAccType op) (codeGenAccType a1) (codeGenAccType a2) <$> codeGenFun f1-codeGenAcc' (AST.Replicate _ e1 a1) = mkReplicate (codeGenAccType a1) <$> codeGenExp e1-codeGenAcc' (AST.Index _ a1 e1) = mkIndex (codeGenAccType a1) <$> codeGenExp e1-codeGenAcc' (AST.Fold f1 e1 _) = mkFold (codeGenExpType e1) <$> codeGenExp e1 <*> codeGenFun f1-codeGenAcc' (AST.FoldSeg f1 e1 _ _) = mkFoldSeg (codeGenExpType e1) <$> codeGenExp e1 <*> codeGenFun f1-codeGenAcc' (AST.Scanl f1 e1 _) = mkScanl (codeGenExpType e1) <$> codeGenExp e1 <*> codeGenFun f1-codeGenAcc' (AST.Scanr f1 e1 _) = mkScanr (codeGenExpType e1) <$> codeGenExp e1 <*> codeGenFun f1-codeGenAcc' (AST.Permute f1 _ f2 a1) = mkPermute (codeGenAccType a1) <$> codeGenFun f1 <*> codeGenFun f2-codeGenAcc' (AST.Backpermute _ f1 a1) = mkBackpermute (codeGenAccType a1) <$> codeGenFun f1+codeGenAcc' op@(AST.Replicate sl e1 a1) = codeGenReplicate sl e1 a1 op <* codeGenExp e1+codeGenAcc' op@(AST.Index sl a1 e1) = codeGenIndex sl a1 op e1 <* codeGenExp e1+codeGenAcc' (AST.Fold f1 e1 _) = mkFold (codeGenExpType e1) <$> codeGenExp e1 <*> codeGenFun f1+codeGenAcc' (AST.FoldSeg f1 e1 _ s) = mkFoldSeg (codeGenExpType e1) (codeGenAccType s) <$> codeGenExp e1 <*> codeGenFun f1+codeGenAcc' (AST.Scanl f1 e1 _) = mkScanl (codeGenExpType e1) <$> codeGenExp e1 <*> codeGenFun f1+codeGenAcc' (AST.Scanr f1 e1 _) = mkScanr (codeGenExpType e1) <$> codeGenExp e1 <*> codeGenFun f1+codeGenAcc' op@(AST.Map f1 a1) = mkMap (codeGenAccType op) (codeGenAccType a1) <$> codeGenFun f1+codeGenAcc' op@(AST.ZipWith f1 a1 a0)+ = mkZipWith (codeGenAccType op) (codeGenShapeType op)+ (codeGenAccType a1) (codeGenShapeType a1)+ (codeGenAccType a0) (codeGenShapeType a0)+ <$> codeGenFun f1 -codeGenAcc' _ =- error "codeGenAcc: internal error"+codeGenAcc' op@(AST.Permute f1 _ f2 a1)+ = mkPermute (codeGenAccType a1) (codeGenShapeType op) (codeGenShapeType a1)+ <$> codeGenFun f1+ <*> codeGenFun f2 +codeGenAcc' op@(AST.Backpermute _ f1 a1)+ = mkBackpermute (codeGenAccType a1) (codeGenShapeType op) (codeGenShapeType a1)+ <$> codeGenFun f1 +codeGenAcc' x =+ INTERNAL_ERROR(error) "codeGenAcc"+ (unlines ["unsupported array primitive", render . nest 2 $ text (show x)])++ -- Embedded expressions -- ~~~~~~~~~~~~~~~~~~~~ @@ -93,12 +116,20 @@ codeGenExp (AST.PrimConst c) = return . unit $ codeGenPrimConst c codeGenExp (AST.PrimApp f arg) = unit . codeGenPrim f <$> codeGenExp arg codeGenExp (AST.Const c) = return $ codeGenConst (Sugar.elemType (undefined::t)) c+codeGenExp (AST.Tuple t) = codeGenTup t+codeGenExp prj@(AST.Prj idx e)+ = reverse+ . take (length $ codeGenTupleType (expType prj))+ . drop (prjToInt idx (expType e))+ . reverse+ <$> codeGenExp e+ codeGenExp (AST.Var i) = let var = CVar (internalIdent ('x' : show (idxToInt i))) internalNode in case codeGenTupleType (Sugar.elemType (undefined::t)) of [_] -> return [var]- cps -> return . reverse . take (length cps) . flip map ['a'..] $- \c -> CMember var (internalIdent [c]) False internalNode+ cps -> return . reverse . take (length cps) . flip map (enumFrom 0 :: [Int]) $+ \c -> CMember var (internalIdent ('a':show c)) False internalNode codeGenExp (AST.Cond p e1 e2) = do [a] <- codeGenExp p@@ -106,11 +137,6 @@ [c] <- codeGenExp e2 return [CCond a (Just b) c internalNode] -codeGenExp (AST.Tuple t) = codeGenTup t-codeGenExp (AST.Prj idx e) = do- [var] <- codeGenExp e- return [CMember var (internalIdent [enumFrom 'a' !! prjToInt idx]) False internalNode]- codeGenExp (AST.IndexScalar a1 e1) = do n <- length <$> get [i] <- codeGenExp e1@@ -132,20 +158,81 @@ codeGenTup NilTup = return [] codeGenTup (t `SnocTup` e) = (++) <$> codeGenTup t <*> codeGenExp e - -- Convert a typed de Brujin index to the corresponding integer -- idxToInt :: AST.Idx env t -> Int idxToInt AST.ZeroIdx = 0 idxToInt (AST.SuccIdx idx) = 1 + idxToInt idx --- Convert a tuple index into the corresponding integer+-- Convert a tuple index into the corresponding integer. Since the internal+-- representation is flat, be sure to walk over all sub components when indexing+-- past nested tuples. ---prjToInt :: TupleIdx t e -> Int-prjToInt ZeroTupIdx = 0-prjToInt (SuccTupIdx idx) = 1 + prjToInt idx+prjToInt :: TupleIdx t e -> TupleType a -> Int+prjToInt ZeroTupIdx _ = 0+prjToInt (SuccTupIdx i) (b `PairTuple` a) = length (codeGenTupleType a) + prjToInt i b+prjToInt _ _ =+ INTERNAL_ERROR(error) "prjToInt" "inconsistent valuation" +-- multidimensional array slice and index+--+codeGenIndex+ :: SliceIndex (Sugar.ElemRepr slix)+ (Sugar.ElemRepr sl)+ co+ (Sugar.ElemRepr dim)+ -> AST.OpenAcc aenv (Sugar.Array dim e)+ -> AST.OpenAcc aenv (Sugar.Array sl e)+ -> AST.Exp aenv slix+ -> CG CUTranslSkel+codeGenIndex sl acc acc' slix =+ return . mkIndex ty dimSl dimCo dimIn0 $ restrict sl (length dimCo-1,length dimSl-1)+ where+ ty = codeGenAccType acc+ dimCo = codeGenExpType slix+ dimSl = codeGenShapeType acc'+ dimIn0 = codeGenShapeType acc++ restrict :: SliceIndex slix sl co dim -> (Int,Int) -> [CExpr]+ restrict (SliceNil) _ = []+ restrict (SliceAll sliceIdx) (m,n) = mkPrj dimSl "sl" n : restrict sliceIdx (m,n-1)+ restrict (SliceFixed sliceIdx) (m,n) = mkPrj dimCo "co" m : restrict sliceIdx (m-1,n)+++-- multidimensional replicate+--+codeGenReplicate+ :: SliceIndex (Sugar.ElemRepr slix)+ (Sugar.ElemRepr sl)+ co+ (Sugar.ElemRepr dim)+ -> AST.Exp aenv slix+ -> AST.OpenAcc aenv (Sugar.Array sl e)+ -> AST.OpenAcc aenv (Sugar.Array dim e)+ -> CG CUTranslSkel+codeGenReplicate sl _slix acc acc' =+ return . mkReplicate ty dimSl dimOut . post $ extend sl (length dimOut-1)+ where+ ty = codeGenAccType acc+ dimSl = codeGenShapeType acc+ dimOut = codeGenShapeType acc'++ extend :: SliceIndex slix sl co dim -> Int -> [CExpr]+ extend (SliceNil) _ = []+ extend (SliceAll sliceIdx) n = mkPrj dimOut "dim" n : extend sliceIdx (n-1)+ extend (SliceFixed sliceIdx) n = extend sliceIdx (n-1)++ post [] = [CConst (CIntConst (cInteger 0) internalNode)]+ post xs = xs+++mkPrj :: [CType] -> String -> Int -> CExpr+mkPrj ty var c+ | length ty <= 1 = CVar (internalIdent var) internalNode+ | otherwise = CMember (CVar (internalIdent var) internalNode) (internalIdent ('a':show c)) False internalNode++ -- Types -- ~~~~~ @@ -157,7 +244,10 @@ codeGenExpType :: AST.OpenExp aenv env t -> [CType] codeGenExpType = codeGenTupleType . expType +codeGenShapeType :: AST.OpenAcc aenv (Sugar.Array dim e) -> [CType]+codeGenShapeType = codeGenTupleType . accShapeType + -- Implementation -- codeGenTupleType :: TupleType a -> [CType]@@ -295,18 +385,18 @@ codeGenPrim (AST.PrimNeg _) [a] = CUnary CMinOp a internalNode codeGenPrim (AST.PrimAbs ty) [a] = codeGenAbs ty a codeGenPrim (AST.PrimSig ty) [a] = codeGenSig ty a-codeGenPrim (AST.PrimQuot ty) [a,b] = codeGenQuot ty a b+codeGenPrim (AST.PrimQuot _) [a,b] = CBinary CDivOp a b internalNode codeGenPrim (AST.PrimRem _) [a,b] = CBinary CRmdOp a b internalNode-codeGenPrim (AST.PrimIDiv _) [a,b] = CBinary CDivOp a b internalNode-codeGenPrim (AST.PrimMod ty) [a,b] = codeGenMod ty a b+codeGenPrim (AST.PrimIDiv _) [a,b] = CCall (CVar (internalIdent "idiv") internalNode) [a,b] internalNode+codeGenPrim (AST.PrimMod _) [a,b] = CCall (CVar (internalIdent "mod") internalNode) [a,b] internalNode codeGenPrim (AST.PrimBAnd _) [a,b] = CBinary CAndOp a b internalNode codeGenPrim (AST.PrimBOr _) [a,b] = CBinary COrOp a b internalNode codeGenPrim (AST.PrimBXor _) [a,b] = CBinary CXorOp a b internalNode codeGenPrim (AST.PrimBNot _) [a] = CUnary CCompOp a internalNode codeGenPrim (AST.PrimBShiftL _) [a,b] = CBinary CShlOp a b internalNode codeGenPrim (AST.PrimBShiftR _) [a,b] = CBinary CShrOp a b internalNode-codeGenPrim (AST.PrimBRotateL _) [a,b] = codeGenBRotateL a b-codeGenPrim (AST.PrimBRotateR _) [a,b] = codeGenBRotateR a b+codeGenPrim (AST.PrimBRotateL _) [a,b] = CCall (CVar (internalIdent "rotateL") internalNode) [a,b] internalNode+codeGenPrim (AST.PrimBRotateR _) [a,b] = CCall (CVar (internalIdent "rotateR") internalNode) [a,b] internalNode codeGenPrim (AST.PrimFDiv _) [a,b] = CBinary CDivOp a b internalNode codeGenPrim (AST.PrimRecip ty) [a] = codeGenRecip ty a codeGenPrim (AST.PrimSin ty) [a] = ccall (FloatingNumType ty) "sin" [a]@@ -323,6 +413,7 @@ codeGenPrim (AST.PrimLog ty) [a] = ccall (FloatingNumType ty) "log" [a] codeGenPrim (AST.PrimFPow ty) [a,b] = ccall (FloatingNumType ty) "pow" [a,b] codeGenPrim (AST.PrimLogBase ty) [a,b] = codeGenLogBase ty a b+codeGenPrim (AST.PrimAtan2 ty) [a,b] = ccall (FloatingNumType ty) "atan2" [a,b] codeGenPrim (AST.PrimLt _) [a,b] = CBinary CLeOp a b internalNode codeGenPrim (AST.PrimGt _) [a,b] = CBinary CGrOp a b internalNode codeGenPrim (AST.PrimLtEq _) [a,b] = CBinary CLeqOp a b internalNode@@ -343,7 +434,7 @@ -- If the argument lists are not the correct length codeGenPrim _ _ =- error "Data.Array.Accelerate.CUDA: inconsistent valuation"+ INTERNAL_ERROR(error) "codeGenPrim" "inconsistent valuation" -- Implementation@@ -353,14 +444,27 @@ codeGenConst (SingleTuple ty) c = [codeGenScalar ty c] codeGenConst (PairTuple ty1 ty0) (cs,c) = codeGenConst ty1 cs ++ codeGenConst ty0 c -+-- FIXME:+-- Language-c isn't pretty printing float constants with a trailing 'f', so as+-- per the C spec nvcc considers them to be double constants. This causes+-- warnings on pre-1.3 series devices, and unnecessary runtime conversion and+-- register pressure on later hardware. Work around this with an explicit type+-- cast. This is quite ugly and should be fixed, but appears to work for now.+-- codeGenScalar :: ScalarType a -> a -> CExpr codeGenScalar (NumScalarType (IntegralNumType ty)) | IntegralDict <- integralDict ty = CConst . flip CIntConst internalNode . cInteger . fromIntegral-codeGenScalar (NumScalarType (FloatingNumType ty))- | FloatingDict <- floatingDict ty+codeGenScalar (NumScalarType (FloatingNumType (TypeFloat _)))+ = flip (CCast (CDecl [CTypeSpec (CFloatType internalNode)] [] internalNode)) internalNode+ . CConst . flip CFloatConst internalNode . cFloat . fromRational . toRational+codeGenScalar (NumScalarType (FloatingNumType (TypeDouble _))) = CConst . flip CFloatConst internalNode . cFloat . fromRational . toRational+codeGenScalar (NumScalarType (FloatingNumType (TypeCFloat _)))+ = flip (CCast (CDecl [CTypeSpec (CFloatType internalNode)] [] internalNode)) internalNode+ . CConst . flip CFloatConst internalNode . cFloat . fromRational . toRational+codeGenScalar (NumScalarType (FloatingNumType (TypeCDouble _)))+ = CConst . flip CFloatConst internalNode . cFloat . fromRational . toRational codeGenScalar (NonNumScalarType (TypeBool _)) = fromBool codeGenScalar (NonNumScalarType (TypeChar _)) = CConst . flip CCharConst internalNode . cChar@@ -410,35 +514,6 @@ (Just (codeGenScalar (NumScalarType ty) 1)) (codeGenScalar (NumScalarType ty) 0) internalNode--codeGenQuot :: IntegralType a -> CExpr -> CExpr -> CExpr-codeGenQuot = error "Data.Array.Accelerate.CUDA.CodeGen: PrimQuot"--codeGenMod :: IntegralType a -> CExpr -> CExpr -> CExpr-codeGenMod = error "Data.Array.Accelerate.CUDA.CodeGen: PrimMod"---- TLM 2010-06-29:--- It would be nice we could use something like Language.C.Parser.execParser--- to suck in the C code directly, instead of storing this long and messy--- abstract syntax. The problem lies in injecting our `x' and `i' expressions.------ T rotl(T x, int i)--- {--- return (i &= 8 * sizeof(x) - 1) == 0 ? x : x << i | x >> 8 * sizeof(x) - i;--- }----codeGenBRotateL :: CExpr -> CExpr -> CExpr-codeGenBRotateL x i =- CCond (CBinary CEqOp (CAssign CAndAssOp i (CBinary CSubOp (CBinary CMulOp (CConst (CIntConst (cInteger 8) internalNode)) (CSizeofExpr x internalNode) internalNode) (CConst (CIntConst (cInteger 1) internalNode)) internalNode) internalNode) (CConst (CIntConst (cInteger 0) internalNode)) internalNode) (Just x) (CBinary COrOp (CBinary CShlOp x i internalNode) (CBinary CShrOp x (CBinary CSubOp (CBinary CMulOp (CConst (CIntConst (cInteger 8) internalNode)) (CSizeofExpr x internalNode) internalNode) i internalNode) internalNode) internalNode) internalNode---- T rotr(T x, int i)--- {--- return (i &= 8 * sizeof(x) - 1) == 0 ? x : x >> i | x << 8 * sizeof(x) - i;--- }----codeGenBRotateR :: CExpr -> CExpr -> CExpr-codeGenBRotateR x i =- CCond (CBinary CEqOp (CAssign CAndAssOp i (CBinary CSubOp (CBinary CMulOp (CConst (CIntConst (cInteger 8) internalNode)) (CSizeofExpr x internalNode) internalNode) (CConst (CIntConst (cInteger 1) internalNode)) internalNode) internalNode) (CConst (CIntConst (cInteger 0) internalNode)) internalNode) (Just x) (CBinary COrOp (CBinary CShrOp x i internalNode) (CBinary CShlOp x (CBinary CSubOp (CBinary CMulOp (CConst (CIntConst (cInteger 8) internalNode)) (CSizeofExpr x internalNode) internalNode) i internalNode) internalNode) internalNode) internalNode codeGenRecip :: FloatingType a -> CExpr -> CExpr codeGenRecip ty x | FloatingDict <- floatingDict ty
Data/Array/Accelerate/CUDA/CodeGen/Skeleton.hs view
@@ -18,6 +18,7 @@ where import Language.C+import System.FilePath import Data.Array.Accelerate.CUDA.CodeGen.Data import Data.Array.Accelerate.CUDA.CodeGen.Util import Data.Array.Accelerate.CUDA.CodeGen.Tuple@@ -34,10 +35,18 @@ ( mkTupleTypeAsc 2 ty ++ [ mkIdentity identity , mkApply 2 apply ])- (mkNodeInfo (initPos "fold.cu") (Name 0))+ (mkNodeInfo (initPos skel) (Name 0)) -mkFoldSeg :: [CType] -> [CExpr] -> [CExpr] -> CUTranslSkel-mkFoldSeg _ty _identity _apply = undefined+mkFoldSeg :: [CType] -> [CType] -> [CExpr] -> [CExpr] -> CUTranslSkel+mkFoldSeg ty int identity apply = CUTranslSkel code skel+ where+ skel = "fold_segmented.inl"+ code = CTranslUnit+ ( mkTupleTypeAsc 2 ty +++ [ mkTypedef "Int" False (head int)+ , mkIdentity identity+ , mkApply 2 apply ])+ (mkNodeInfo (initPos skel) (Name 0)) --------------------------------------------------------------------------------@@ -52,19 +61,22 @@ ( mkTupleType Nothing tyOut ++ mkTupleType (Just 0) tyIn0 ++ [ mkApply 1 apply ])- (mkNodeInfo (initPos "map.cu") (Name 0))+ (mkNodeInfo (initPos skel) (Name 0)) -mkZipWith :: [CType] -> [CType] -> [CType] -> [CExpr] -> CUTranslSkel-mkZipWith tyOut tyIn1 tyIn0 apply = CUTranslSkel code skel+mkZipWith :: [CType] -> [CType] -> [CType] -> [CType] -> [CType] -> [CType] -> [CExpr] -> CUTranslSkel+mkZipWith tyOut shOut tyIn1 shIn1 tyIn0 shIn0 apply = CUTranslSkel code skel where skel = "zipWith.inl" code = CTranslUnit ( mkTupleType Nothing tyOut ++ mkTupleType (Just 1) tyIn1 ++ mkTupleType (Just 0) tyIn0 ++- [ mkApply 2 apply ])- (mkNodeInfo (initPos "zipWith.cu") (Name 0))+ [ mkApply 2 apply+ , mkDim "DimOut" shOut+ , mkDim "DimIn1" shIn1+ , mkDim "DimIn0" shIn0 ])+ (mkNodeInfo (initPos skel) (Name 0)) --------------------------------------------------------------------------------@@ -75,15 +87,15 @@ mkScan isBackward ty identity apply = CUTranslSkel code skel where- skel | length ty == 1 = "thrust/scan_safe.inl" -- TODO: use fast scan for primitive types- | otherwise = "thrust/scan_safe.inl"+ skel | length ty == 1 = "thrust" </> "scan_safe.inl" -- TODO: use fast scan for primitive types+ | otherwise = "thrust" </> "scan_safe.inl" code = CTranslUnit ( mkTupleTypeAsc 2 ty ++ [ mkIdentity identity , mkApply 2 apply , mkFlag "reverse" (fromBool isBackward) ])- (mkNodeInfo (initPos "scan.cu") (Name 0))+ (mkNodeInfo (initPos (takeFileName skel)) (Name 0)) mkScanl :: [CType] -> [CExpr] -> [CExpr] -> CUTranslSkel@@ -107,44 +119,55 @@ -- Permutation -------------------------------------------------------------------------------- -mkPermute :: [CType] -> [CExpr] -> [CExpr] -> CUTranslSkel-mkPermute ty combine index = CUTranslSkel code skel+mkPermute :: [CType] -> [CType] -> [CType] -> [CExpr] -> [CExpr] -> CUTranslSkel+mkPermute ty dimOut dimIn0 combinefn indexfn = CUTranslSkel code skel where skel = "permute.inl" code = CTranslUnit ( mkTupleTypeAsc 2 ty ++- [ mkIgnore (-1)- , mkIndexFun index- , mkApply 2 combine ])- (mkNodeInfo (initPos "permute.cu") (Name 0))---- TLM 2010-07-10:--- should generate from Sugar.Ix (or Representation.Ix) notion of ignore----mkIgnore :: Integer -> CExtDecl-mkIgnore n =- CDeclExt- (CDecl [CTypeQual (CConstQual internalNode),CTypeSpec (CTypeDef (internalIdent "Ix") internalNode)]- [(Just (CDeclr (Just (internalIdent "ignore")) [] Nothing [] internalNode),Just (CInitExpr (CCast (CDecl [CTypeSpec (CTypeDef (internalIdent "Ix") internalNode)] [] internalNode) (CConst (CIntConst (cInteger n) internalNode)) internalNode) internalNode),Nothing)]- internalNode)+ [ mkDim "DimOut" dimOut+ , mkDim "DimIn0" dimIn0+ , mkProject indexfn+ , mkApply 2 combinefn ])+ (mkNodeInfo (initPos skel) (Name 0)) -mkBackpermute :: [CType] -> [CExpr] -> CUTranslSkel-mkBackpermute ty index = CUTranslSkel code skel+mkBackpermute :: [CType] -> [CType] -> [CType] -> [CExpr] -> CUTranslSkel+mkBackpermute ty dimOut dimIn0 index = CUTranslSkel code skel where skel = "backpermute.inl" code = CTranslUnit ( mkTupleTypeAsc 1 ty ++- [ mkIndexFun index ])- (mkNodeInfo (initPos "backpermute.cu") (Name 0))+ [ mkDim "DimOut" dimOut+ , mkDim "DimIn0" dimIn0+ , mkProject index ])+ (mkNodeInfo (initPos skel) (Name 0)) -------------------------------------------------------------------------------- -- Multidimensional Index and Replicate -------------------------------------------------------------------------------- -mkIndex :: [CType] -> [CExpr] -> CUTranslSkel-mkIndex _ty _slix = undefined+mkIndex :: [CType] -> [CType] -> [CType] -> [CType] -> [CExpr] -> CUTranslSkel+mkIndex ty dimSl dimCo dimIn0 slix = CUTranslSkel code skel+ where+ skel = "slice.inl"+ code = CTranslUnit+ ( mkTupleTypeAsc 1 ty +++ [ mkDim "Slice" dimSl+ , mkDim "CoSlice" dimCo+ , mkDim "SliceDim" dimIn0+ , mkSliceIndex slix ])+ (mkNodeInfo (initPos skel) (Name 0)) -mkReplicate :: [CType] -> [CExpr] -> CUTranslSkel-mkReplicate _ty _slix = undefined++mkReplicate :: [CType] -> [CType] -> [CType] -> [CExpr] -> CUTranslSkel+mkReplicate ty dimSl dimOut slix = CUTranslSkel code skel+ where+ skel = "replicate.inl"+ code = CTranslUnit+ ( mkTupleTypeAsc 1 ty +++ [ mkDim "Slice" dimSl+ , mkDim "SliceDim" dimOut+ , mkSliceReplicate slix ])+ (mkNodeInfo (initPos skel) (Name 0))
Data/Array/Accelerate/CUDA/CodeGen/Tuple.hs view
@@ -66,8 +66,8 @@ arrIn = internalIdent ("d_in" ++ show prj) initList | n <= 1 = CInitExpr (CIndex (CVar arrIn internalNode) (CVar (internalIdent "idx") internalNode) internalNode) internalNode- | otherwise = flip CInitList internalNode . reverse . take n . flip map (enumFrom 'a') $ \v ->- ([], CInitExpr (CIndex (CMember (CVar arrIn internalNode) (internalIdent [v]) False internalNode) (CVar (internalIdent "idx") internalNode) internalNode) internalNode)+ | otherwise = flip CInitList internalNode . reverse . take n . flip map (enumFrom 0 :: [Int]) $ \v ->+ ([], CInitExpr (CIndex (CMember (CVar arrIn internalNode) (internalIdent ('a':show v)) False internalNode) (CVar (internalIdent "idx") internalNode) internalNode) internalNode) mkSet :: Int -> CExtDecl@@ -82,6 +82,6 @@ where assignList | n <= 1 = [CBlockStmt (CExpr (Just (CAssign CAssignOp (CIndex (CVar (internalIdent "d_out") internalNode) (CVar (internalIdent "idx") internalNode) internalNode) (CVar (internalIdent "val") internalNode) internalNode)) internalNode)]- | otherwise = reverse . take n . flip map (enumFrom 'a') $ \v ->- CBlockStmt (CExpr (Just (CAssign CAssignOp (CIndex (CMember (CVar (internalIdent "d_out") internalNode) (internalIdent [v]) False internalNode) (CVar (internalIdent "idx") internalNode) internalNode) (CMember (CVar (internalIdent "val") internalNode) (internalIdent [v]) False internalNode) internalNode)) internalNode)+ | otherwise = reverse . take n . flip map (enumFrom 0 :: [Int]) $ \v ->+ CBlockStmt (CExpr (Just (CAssign CAssignOp (CIndex (CMember (CVar (internalIdent "d_out") internalNode) (internalIdent ('a':show v)) False internalNode) (CVar (internalIdent "idx") internalNode) internalNode) (CMember (CVar (internalIdent "val") internalNode) (internalIdent ('a':show v)) False internalNode) internalNode)) internalNode)
Data/Array/Accelerate/CUDA/CodeGen/Util.hs view
@@ -15,20 +15,59 @@ import Data.Array.Accelerate.CUDA.CodeGen.Data +--------------------------------------------------------------------------------+-- Common device functions+--------------------------------------------------------------------------------++mkIdentity :: [CExpr] -> CExtDecl+mkIdentity = mkDeviceFun "identity" (typename "TyOut") []++mkApply :: Int -> [CExpr] -> CExtDecl+mkApply argc+ = mkDeviceFun "apply" (typename "TyOut")+ $ map (\n -> (typename ("TyIn"++show n), 'x':show n)) [argc-1,argc-2..0]++mkProject :: [CExpr] -> CExtDecl+mkProject = mkDeviceFun "project" (typename "DimOut") [(typename "DimIn0","x0")]++mkSliceIndex :: [CExpr] -> CExtDecl+mkSliceIndex =+ mkDeviceFun "sliceIndex" (typename "SliceDim") [(typename "Slice","sl"), (typename "CoSlice","co")]++mkSliceReplicate :: [CExpr] -> CExtDecl+mkSliceReplicate =+ mkDeviceFun "sliceIndex" (typename "Slice") [(typename "SliceDim","dim")]+++--------------------------------------------------------------------------------+-- Helper functions+--------------------------------------------------------------------------------++typename :: String -> CType+typename var = [CTypeDef (internalIdent var) internalNode]+ fromBool :: Bool -> CExpr fromBool True = CConst $ CIntConst (cInteger 1) internalNode fromBool False = CConst $ CIntConst (cInteger 0) internalNode +mkDim :: String -> [CType] -> CExtDecl+mkDim name ty =+ mkTypedef name False [CTypeDef (internalIdent ("DIM" ++ show (length ty))) internalNode] --- typedef ty (*?) var;--- mkTypedef :: String -> Bool -> CType -> CExtDecl mkTypedef var ptr ty =- CDeclExt (CDecl (CStorageSpec (CTypedef internalNode) : map CTypeSpec ty) [(Just (CDeclr (Just (internalIdent var)) [CPtrDeclr [] internalNode | ptr] Nothing [] internalNode), Nothing, Nothing)] internalNode)+ CDeclExt $ CDecl+ (CStorageSpec (CTypedef internalNode) : map CTypeSpec ty)+ [(Just (CDeclr (Just (internalIdent var)) [CPtrDeclr [] internalNode | ptr] Nothing [] internalNode), Nothing, Nothing)]+ internalNode +mkInitList :: [CExpr] -> CInit+mkInitList [x] = CInitExpr x internalNode+mkInitList xs = CInitList (map (\e -> ([],CInitExpr e internalNode)) xs) internalNode + -- typedef struct {--- ... ty1 (*?) b; ty0 (*?) a;+-- ... ty1 (*?) a1; ty0 (*?) a0; -- } var; -- -- NOTE: The Accelerate language uses snoc based tuple projection, so the last@@ -36,10 +75,15 @@ -- mkStruct :: String -> Bool -> [CType] -> CExtDecl mkStruct name ptr types =- CDeclExt (CDecl [ CStorageSpec (CTypedef internalNode) , CTypeSpec (CSUType (CStruct CStructTag Nothing (Just (zipWith field names types)) [] internalNode) internalNode)] [(Just (CDeclr (Just (internalIdent name)) [] Nothing [] internalNode),Nothing,Nothing)] internalNode)+ CDeclExt $ CDecl+ [CStorageSpec (CTypedef internalNode) , CTypeSpec (CSUType (CStruct CStructTag Nothing (Just (zipWith field names types)) [] internalNode) internalNode)]+ [(Just (CDeclr (Just (internalIdent name)) [] Nothing [] internalNode),Nothing,Nothing)]+ internalNode where- names = reverse . take (length types) $ enumFrom 'a'- field v ty = CDecl (map CTypeSpec ty) [(Just (CDeclr (Just (internalIdent [v])) [CPtrDeclr [] internalNode | ptr] Nothing [] internalNode),Nothing,Nothing)] internalNode+ names = reverse . take (length types) $ (enumFrom 0 :: [Int])+ field v ty = CDecl (map CTypeSpec ty)+ [(Just (CDeclr (Just (internalIdent ('a':show v))) [CPtrDeclr [] internalNode | ptr] Nothing [] internalNode), Nothing, Nothing)]+ internalNode -- typedef struct __attribute__((aligned(n * sizeof(ty)))) {@@ -48,55 +92,32 @@ -- mkTyVector :: String -> Int -> CType -> CExtDecl mkTyVector var n ty =- CDeclExt (CDecl [ CStorageSpec (CTypedef internalNode) , CTypeSpec (CSUType (CStruct CStructTag Nothing (Just [CDecl (map CTypeSpec ty) fields internalNode]) [CAttr (internalIdent "aligned") [CBinary CMulOp (CConst (CIntConst (cInteger (toInteger n)) internalNode)) (CSizeofType (CDecl (map CTypeSpec ty) [] internalNode) internalNode) internalNode] internalNode] internalNode) internalNode)] [ (Just (CDeclr (Just (internalIdent var)) [] Nothing [] internalNode), Nothing, Nothing)] internalNode)+ CDeclExt $ CDecl+ [CStorageSpec (CTypedef internalNode), CTypeSpec (CSUType (CStruct CStructTag Nothing (Just [CDecl (map CTypeSpec ty) fields internalNode]) [CAttr (internalIdent "aligned") [CBinary CMulOp (CConst (CIntConst (cInteger (toInteger n)) internalNode)) (CSizeofType (CDecl (map CTypeSpec ty) [] internalNode) internalNode) internalNode] internalNode] internalNode) internalNode)]+ [(Just (CDeclr (Just (internalIdent var)) [] Nothing [] internalNode), Nothing, Nothing)]+ internalNode where fields = take n . flip map "xyzw" $ \f -> (Just (CDeclr (Just (internalIdent [f])) [] Nothing [] internalNode), Nothing, Nothing) --- static inline __attribute__((device)) TyOut identity()--- {--- return expr;--- }----mkIdentity :: [CExpr] -> CExtDecl-mkIdentity expr =- CFDefExt (CFunDef [CStorageSpec (CStatic internalNode),CTypeQual (CInlineQual internalNode),CTypeQual (CAttrQual (CAttr (builtinIdent "device") [] internalNode)),CTypeSpec (CTypeDef (internalIdent "TyOut") internalNode)] (CDeclr (Just (internalIdent "identity")) [CFunDeclr (Right ([],False)) [] internalNode] Nothing [] internalNode) [] (CCompound [] [CBlockDecl (CDecl [CTypeSpec (CTypeDef (internalIdent "TyOut") internalNode)] [(Just (CDeclr (Just (internalIdent "x")) [] Nothing [] internalNode),Just (mkInitList expr),Nothing)] internalNode),CBlockStmt (CReturn (Just (CVar (internalIdent "x") internalNode)) internalNode)] internalNode) internalNode)----- static inline __attribute__((device)) TyOut--- apply(..., const TyIn1 x1, const TyIn0 x0)--- {--- TyOut x; x.a = x0; x.b = x1; ...--- return x;--- }----mkApply :: Int -> [CExpr] -> CExtDecl-mkApply argc expr =- CFDefExt (CFunDef [CStorageSpec (CStatic internalNode),CTypeQual (CInlineQual internalNode),CTypeQual (CAttrQual (CAttr (builtinIdent "device") [] internalNode)),CTypeSpec (CTypeDef (internalIdent "TyOut") internalNode)] (CDeclr (Just (internalIdent "apply")) [CFunDeclr (Right (argv,False)) [] internalNode] Nothing [] internalNode) [] (CCompound [] [CBlockDecl (CDecl [CTypeSpec (CTypeDef (internalIdent "TyOut") internalNode)] [(Just (CDeclr (Just (internalIdent "x")) [] Nothing [] internalNode),Just (mkInitList expr),Nothing)] internalNode),CBlockStmt (CReturn (Just (CVar (internalIdent "x") internalNode)) internalNode)] internalNode) internalNode)- where- argv = reverse . take argc . flip map (enumFrom 0 :: [Int]) $ \x ->- let ty = "TyIn" ++ show x- var = 'x' : show x- in- CDecl [CTypeQual (CConstQual internalNode), CTypeSpec (CTypeDef (internalIdent ty) internalNode)]- [(Just (CDeclr (Just (internalIdent var)) [] Nothing [] internalNode), Nothing, Nothing)]- internalNode---- Not strictly necessary, just attention to detail----mkInitList :: [CExpr] -> CInit-mkInitList [x] = CInitExpr x internalNode-mkInitList xs = CInitList (map (\e -> ([],CInitExpr e internalNode)) xs) internalNode----- static inline __attribute__((device)) Ix--- project(const Ix x0)+-- static inline __attribute__((device)) tyout name(args) -- {--- return expr;+-- tyout r = { expr };+-- return r; -- } ---mkIndexFun :: [CExpr] -> CExtDecl-mkIndexFun [expr] = CFDefExt (CFunDef [CStorageSpec (CStatic internalNode),CTypeQual (CInlineQual internalNode),CTypeQual (CAttrQual (CAttr (builtinIdent "device") [] internalNode)),CTypeSpec (CTypeDef (internalIdent "Ix") internalNode)] (CDeclr (Just (internalIdent "project")) [CFunDeclr (Right ([CDecl [CTypeQual (CConstQual internalNode),CTypeSpec (CTypeDef (internalIdent "Ix") internalNode)] [(Just (CDeclr (Just (internalIdent "x0")) [] Nothing [] internalNode),Nothing,Nothing)] internalNode],False)) [] internalNode] Nothing [] internalNode) [] (CCompound [] [CBlockStmt (CReturn (Just expr) internalNode)] internalNode) internalNode)-mkIndexFun _ = error "mkIndexFun: internal error"+mkDeviceFun :: String -> CType -> [(CType,String)] -> [CExpr] -> CExtDecl+mkDeviceFun name tyout args expr =+ CFDefExt $ CFunDef+ ([CStorageSpec (CStatic internalNode), CTypeQual (CInlineQual internalNode), CTypeQual (CAttrQual (CAttr (builtinIdent "device") [] internalNode))] ++ map CTypeSpec tyout)+ (CDeclr (Just (internalIdent name)) [CFunDeclr (Right (argv,False)) [] internalNode] Nothing [] internalNode)+ []+ (CCompound [] [CBlockDecl (CDecl (map CTypeSpec tyout) [(Just (CDeclr (Just (internalIdent "r")) [] Nothing [] internalNode), Just (mkInitList expr), Nothing)] internalNode), CBlockStmt (CReturn (Just (CVar (internalIdent "r") internalNode)) internalNode)] internalNode)+ internalNode+ where+ argv = flip map args $ \(ty,var) ->+ CDecl (CTypeQual (CConstQual internalNode) : map CTypeSpec ty)+ [(Just (CDeclr (Just (internalIdent var)) [] Nothing [] internalNode), Nothing, Nothing)]+ internalNode
Data/Array/Accelerate/CUDA/Compile.hs view
@@ -22,7 +22,7 @@ import Control.Applicative hiding (Const) import Language.C import Text.PrettyPrint-import qualified Data.Map as M+import qualified Data.HashTable as HT import System.Directory import System.FilePath@@ -49,9 +49,9 @@ compileAcc :: OpenAcc aenv a -> CIO () compileAcc (Use (Array sh ad)) = let n = size sh- in do- mallocArray ad n- pokeArrayAsync ad n Nothing+ in when (n > 0) $ do+ mallocArray ad n+ pokeArrayAsync ad n Nothing compileAcc (Let a1 a2) = compileAcc a1 >> compileAcc a2 compileAcc (Let2 a1 a2) = compileAcc a1 >> compileAcc a2@@ -69,6 +69,10 @@ compileAcc op@(Backpermute e1 f1 a1) = compileExp e1 >> compileFun f1 >> compileAcc a1 >> compile op compileAcc op@(Permute f1 a1 f2 a2) = compileFun f1 >> compileAcc a1 >> compileFun f2 >> compileAcc a2 >> compile op compileAcc op@(FoldSeg f1 e1 a1 a2) = compileFun f1 >> compileExp e1 >> compileAcc a1 >> compileAcc a2 >> compile op+compileAcc (Stencil _ _ _) = + error "D.A.Accelerate.CUDA.Compile: the CUDA backend does not support 'stencil' operations yet"+compileAcc (Stencil2 _ _ _ _ _) = + error "D.A.Accelerate.CUDA.Compile: the CUDA backend does not support 'stencil2' operations yet" -- | Lift array expressions out of closed functions for code generation@@ -100,7 +104,8 @@ compile :: OpenAcc aenv a -> CIO () compile acc = do let key = accToKey acc- compiled <- M.member key <$> getM kernelEntry+ table <- getM kernelTable+ compiled <- isJust <$> liftIO (HT.lookup table key) unless compiled $ do nvcc <- fromMaybe (error "nvcc: command not found") <$> liftIO (findExecutable "nvcc") dir <- getM outputDir@@ -110,7 +115,7 @@ writeCode cufile (codeGenAcc acc) forkProcess $ executeFile nvcc False flags Nothing - modM kernelEntry $ M.insert key (KernelEntry cufile (Left pid))+ liftIO $ HT.insert table key (KernelEntry cufile (Left pid)) -- Write the generated code to file
Data/Array/Accelerate/CUDA/Execute.hs view
@@ -1,4 +1,4 @@-{-# LANGUAGE GADTs, TupleSections #-}+{-# LANGUAGE CPP, GADTs, TupleSections #-} -- | -- Module : Data.Array.Accelerate.CUDA.Execute -- Copyright : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee, Trevor L. McDonell@@ -12,24 +12,25 @@ module Data.Array.Accelerate.CUDA.Execute (executeAcc) where -import Prelude hiding (id, (.), mod, sum)+import Prelude hiding (id, (.), sum) import Control.Category import Data.Maybe import Control.Monad import Control.Monad.Trans import Control.Applicative hiding (Const)-import qualified Data.Map as M+import qualified Data.HashTable as HT import System.FilePath import System.Posix.Process import System.Exit (ExitCode(..)) import System.Posix.Types (ProcessID)+import Text.PrettyPrint import Data.Array.Accelerate.AST import Data.Array.Accelerate.Tuple import Data.Array.Accelerate.Array.Data-import Data.Array.Accelerate.Array.Representation+import Data.Array.Accelerate.Array.Representation hiding (sliceIndex) import Data.Array.Accelerate.Array.Sugar (Array(..)) import qualified Data.Array.Accelerate.Array.Sugar as Sugar import qualified Data.Array.Accelerate.Interpreter as I@@ -41,7 +42,9 @@ import qualified Foreign.CUDA.Driver as CUDA +#include "accelerate.h" + -- Expression evaluation -- ~~~~~~~~~~~~~~~~~~~~~ @@ -96,7 +99,7 @@ -- the same kernel will be able to extract the loaded kernel directly. -- executeOpenAcc :: OpenAcc aenv a -> Val aenv -> CIO a-executeOpenAcc (Use xs) _env = return xs+executeOpenAcc (Use xs) _env = return xs executeOpenAcc (Avar ix) env = return $ prj ix env executeOpenAcc (Let x y) env = do ax <- executeOpenAcc x env@@ -119,11 +122,12 @@ executeOpenAcc acc env = do- krn <- fromMaybe (error "code generation failed") . M.lookup key <$> getM kernelEntry- mdl <- either' (get kernelStatus krn) return $ \pid -> do+ tab <- getM kernelTable+ krn <- fromMaybe (error "code generation failed") <$> liftIO (HT.lookup tab key)+ mdl <- either' (getL kernelStatus krn) return $ \pid -> do liftIO (waitFor pid)- mdl <- liftIO $ CUDA.loadFile (get kernelName krn `replaceExtension` ".cubin")- modM kernelEntry $ M.insert key (set kernelStatus (Right mdl) krn)+ mdl <- liftIO $ CUDA.loadFile (getL kernelName krn `replaceExtension` ".cubin")+ liftIO $ HT.insert tab key (setL kernelStatus (Right mdl) krn) return mdl -- determine dispatch pattern, extract parameters, allocate storage, run@@ -151,6 +155,7 @@ liftExp :: OpenExp env aenv a -> Val aenv -> CIO [FVar] liftExp (Tuple t) aenv = liftTup t aenv+liftExp (Prj _ e) aenv = liftExp e aenv liftExp (PrimApp _ e) aenv = liftExp e aenv liftExp (Cond e1 e2 e3) aenv = concat <$> sequence [liftExp e1 aenv, liftExp e2 aenv, liftExp e3 aenv] liftExp (IndexScalar a e) aenv = (:) . FArr <$> executeOpenAcc a aenv <*> liftExp e aenv@@ -167,7 +172,7 @@ bind :: CUDA.Module -> [FVar] -> CIO [CUDA.FunParam] bind mdl var = let tex n (FArr (Array sh ad)) = textureRefs ad mdl (size sh) n- in concat <$> zipWithM tex [0..] var+ in foldM (\texs farr -> (texs ++) <$> tex (length texs) farr) [] var release :: [FVar] -> CIO () release = mapM_ (\(FArr (Array _ ad)) -> freeArray ad)@@ -219,7 +224,7 @@ f_var <- liftFun f env t_var <- bind mdl f_var - launch acc n fn (d_out ++ d_in1 ++ d_in0 ++ t_var ++ [CUDA.IArg n])+ launch acc n fn (d_out ++ d_in1 ++ d_in0 ++ t_var ++ convertIx sh' ++ convertIx sh1 ++ convertIx sh0) freeArray in0 freeArray in1 release f_var@@ -234,7 +239,7 @@ mallocArray out grid d_out <- devicePtrs out d_in0 <- devicePtrs in0- f_arr <- liftFun f env+ f_arr <- (++) <$> liftExp x env <*> liftFun f env t_var <- bind mdl f_arr launch' (cta,grid,smem) fn (d_out ++ d_in0 ++ t_var ++ [CUDA.IArg (size sh)])@@ -243,15 +248,36 @@ if grid > 1 then dispatch (Fold f x (Use res)) env mdl else return (Array (Sugar.fromElem ()) out) +dispatch acc@(FoldSeg f x ad sd) env mdl = do+ fn <- liftIO $ CUDA.getFun mdl "fold_segmented"+ (Array sh in0) <- executeOpenAcc ad env+ (Array sh' seg) <- executeOpenAcc sd env+ (cta,grid,smem) <- launchConfig acc (size sh') fn+ let res@(Array _ out) = newArray (size sh')++ mallocArray out (size sh')+ d_out <- devicePtrs out+ d_in0 <- devicePtrs in0+ d_seg <- devicePtrs seg+ f_arr <- (++) <$> liftExp x env <*> liftFun f env+ t_var <- bind mdl f_arr++ launch' (cta,grid,smem) fn (d_out ++ d_in0 ++ d_seg ++ t_var ++ map (CUDA.IArg . size) [sh', sh])+ freeArray in0+ freeArray seg+ release f_arr+ return res++ dispatch acc@(Scanl _ _ _) env mdl = dispatchScan acc env mdl dispatch acc@(Scanr _ _ _) env mdl = dispatchScan acc env mdl dispatch acc@(Permute f1 df f2 ad) env mdl = do fn <- liftIO $ CUDA.getFun mdl "permute" (Array sh def) <- executeOpenAcc df env- (Array sh' in0) <- executeOpenAcc ad env+ (Array sh0 in0) <- executeOpenAcc ad env let res@(Array _ out) = newArray (Sugar.toElem sh)- n = size sh'+ n = size sh0 mallocArray out n copyArray def out n@@ -260,16 +286,16 @@ f_arr <- (++) <$> liftFun f1 env <*> liftFun f2 env t_var <- bind mdl f_arr - launch acc n fn (d_out ++ d_in0 ++ t_var ++ [CUDA.IArg n])+ launch acc n fn (d_out ++ d_in0 ++ t_var ++ convertIx sh ++ convertIx sh0) freeArray def freeArray in0 release f_arr return res dispatch acc@(Backpermute e f ad) env mdl = do- fn <- liftIO $ CUDA.getFun mdl "backpermute"- sh <- executeExp e env- (Array _ in0) <- executeOpenAcc ad env+ fn <- liftIO $ CUDA.getFun mdl "backpermute"+ sh <- executeExp e env+ (Array sh0 in0) <- executeOpenAcc ad env let res@(Array sh' out) = newArray sh n = size sh' @@ -279,15 +305,67 @@ f_arr <- liftFun f env t_var <- bind mdl f_arr - launch acc n fn (d_out ++ d_in0 ++ t_var ++ [CUDA.IArg n])+ launch acc n fn (d_out ++ d_in0 ++ t_var ++ convertIx sh' ++ convertIx sh0) freeArray in0 release f_arr return res -dispatch _ _ _ =- error "Data.Array.Accelerate.CUDA: dispatch: internal error"+dispatch acc@(Replicate sliceIndex e ad) env mdl = do+ fn <- liftIO $ CUDA.getFun mdl "replicate"+ slix <- executeExp e env+ (Array sh0 in0) <- executeOpenAcc ad env + let block = extend sliceIndex (Sugar.fromElem slix) sh0+ res@(Array sh' out) = newArray (Sugar.toElem block)+ n = size sh' + extend :: SliceIndex slix sl co dim -> slix -> sl -> dim+ extend (SliceNil) () () = ()+ extend (SliceAll sliceIdx) (slx,()) (sl,sz) = (extend sliceIdx slx sl, sz)+ extend (SliceFixed sliceIdx) (slx,sz) sl = (extend sliceIdx slx sl, sz)++ mallocArray out n+ d_out <- devicePtrs out+ d_in0 <- devicePtrs in0+ f_arr <- liftExp e env+ t_var <- bind mdl f_arr++ launch acc n fn (d_out ++ d_in0 ++ t_var ++ convertIx sh0 ++ convertIx sh')+ freeArray in0+ release f_arr+ return res++dispatch acc@(Index sliceIndex ad e) env mdl = do+ fn <- liftIO $ CUDA.getFun mdl "slice"+ slix <- executeExp e env+ (Array sh0 in0) <- executeOpenAcc ad env+ let slice = restrict sliceIndex (Sugar.fromElem slix) sh0+ slix' = convertSliceIndex sliceIndex (Sugar.fromElem slix)+ res@(Array sh' out) = newArray (Sugar.toElem slice)+ n = size sh'++ restrict :: SliceIndex slix sl co dim -> slix -> dim -> sl+ restrict (SliceNil) () () = ()+ restrict (SliceAll sliceIdx) (slx,()) (sh,sz) = (restrict sliceIdx slx sh, sz)+ restrict (SliceFixed sliceIdx) (slx,i) (sh,sz)+ = BOUNDS_CHECK(checkIndex) "slice" i sz $ restrict sliceIdx slx sh++ mallocArray out n+ d_out <- devicePtrs out+ d_in0 <- devicePtrs in0+ f_arr <- liftExp e env+ t_var <- bind mdl f_arr++ launch acc n fn (d_out ++ d_in0 ++ t_var ++ convertIx sh' ++ slix' ++ convertIx sh0)+ freeArray in0+ release f_arr+ return res++dispatch x _ _ =+ INTERNAL_ERROR(error) "dispatch"+ (unlines ["unsupported array primitive", render . nest 2 $ text (show x)])++ -- Unified dispatch handler for left/right scan. -- -- TLM 2010-07-02:@@ -303,7 +381,7 @@ -- dispatchScan :: OpenAcc aenv a -> Val aenv -> CUDA.Module -> CIO a dispatchScan (Scanr f x ad) env mdl = dispatchScan (Scanl f x ad) env mdl-dispatchScan acc@(Scanl f _ ad) env mdl = do+dispatchScan acc@(Scanl f x ad) env mdl = do fscan <- liftIO $ CUDA.getFun mdl "inclusive_scan" fadd <- liftIO $ CUDA.getFun mdl "exclusive_update" (Array sh in0) <- executeOpenAcc ad env@@ -312,13 +390,12 @@ a_sum@(Array _ sum) = newArray () a_bks@(Array _ bks) = newArray grid n = size sh- units = (n + cta - 1) `div` cta- interval = cta * ((units + grid - 1) `div` grid)+ interval = (n + grid - 1) `div` grid - unify :: Array dim e -> Array dim' e -> CIO ()+ unify :: Array dim e -> Array dim e -> CIO () unify _ _ = return () - unify a_sum a_bks -- TLM: *cough*+ unify a_out a_bks -- TLM: *cough* mallocArray out n mallocArray sum 1@@ -327,7 +404,7 @@ d_in0 <- devicePtrs in0 d_bks <- devicePtrs bks d_sum <- devicePtrs sum- f_arr <- liftFun f env+ f_arr <- (++) <$> liftExp x env <*> liftFun f env t_var <- bind mdl f_arr launch' (cta,grid,smem) fscan (d_out ++ d_in0 ++ d_bks ++ t_var ++ map CUDA.IArg [n,interval])@@ -336,10 +413,11 @@ freeArray in0 freeArray bks+ release f_arr return (a_out, a_sum) dispatchScan _ _ _ =- error "Data.Array.Accelerate.CUDA: internal error"+ error "we can never get here" -- Initiate the device computation. The first version selects launch parameters@@ -369,6 +447,22 @@ -- FIXME: small arrays are relocated by the GC ad = fst . runArrayData $ (,undefined) <$> newArrayData (1024 `max` Sugar.size sh) +-- Extract shape dimensions as a list of function parameters. Not that this will+-- convert to the base integer width of the device, namely, 32-bits. Singleton+-- dimensions are considered to be of unit size.+--+convertIx :: Ix dim => dim -> [CUDA.FunParam]+convertIx = post . map CUDA.IArg . shapeToList+ where post [] = [CUDA.IArg 1]+ post xs = xs++-- Convert a slice specification into storable index projection components.+-- Note: implicit conversion Int -> Int32+--+convertSliceIndex :: SliceIndex slix sl co dim -> slix -> [CUDA.FunParam]+convertSliceIndex (SliceNil) () = []+convertSliceIndex (SliceAll sliceIdx) (s,()) = convertSliceIndex sliceIdx s+convertSliceIndex (SliceFixed sliceIdx) (s,i) = CUDA.IArg i : convertSliceIndex sliceIdx s -- | Wait for the compilation process to finish --
+ Data/Array/Accelerate/CUDA/Smart.hs view
@@ -0,0 +1,77 @@+{-# LANGUAGE GADTs #-}+-- |+-- Module : Data.Array.Accelerate.CUDA.Smart+-- Copyright : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee, Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Manuel M T Chakravarty <chak@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-partable (GHC extensions)+--++module Data.Array.Accelerate.CUDA.Smart (convertAcc, Sugar.Acc) where++import Prelude hiding (exp)++import Data.Array.Accelerate.AST+import qualified Data.Array.Accelerate.Smart as Sugar+import qualified Data.Array.Accelerate.Array.Sugar as Sugar+++-- Convert a closed array expression from HOAS to de Bruijn form AST+--+convertAcc :: Sugar.Acc a -> Acc a+convertAcc = refineAcc . Sugar.convertAcc+++-- Inject extra computational steps specific to the CUDA backend implementation.+-- TODO: Implement array fusion+--+refineAcc :: Acc a -> Acc a+refineAcc = refineOpenAcc Empty++refineOpenAcc :: Val aenv -> OpenAcc aenv a -> OpenAcc aenv a+refineOpenAcc aenv (Let a1 a2)+ = Let (refineOpenAcc aenv a1)+ (refineOpenAcc (aenv `Push` undefined) a2)++refineOpenAcc aenv (Let2 a1 a2)+ = Let2 (refineOpenAcc aenv a1)+ (refineOpenAcc (aenv `Push` undefined `Push` undefined) a2)++refineOpenAcc aenv (Unit e) = Unit (refineExp aenv e)+refineOpenAcc aenv (Reshape e a) = Reshape (refineExp aenv e) (refineOpenAcc aenv a)+refineOpenAcc aenv (Replicate slix e a) = Replicate slix (refineExp aenv e) (refineOpenAcc aenv a)+refineOpenAcc aenv (Index slix a e) = Index slix (refineOpenAcc aenv a) (refineExp aenv e)+refineOpenAcc aenv (Map f a) = Map (refineFun aenv f) (refineOpenAcc aenv a)+refineOpenAcc aenv (ZipWith f a1 a2) = ZipWith (refineFun aenv f) (refineOpenAcc aenv a1) (refineOpenAcc aenv a2)+refineOpenAcc aenv (Fold f e a) = Fold (refineFun aenv f) (refineExp aenv e) (refineOpenAcc aenv a)+refineOpenAcc aenv (Scanl f e a) = Scanl (refineFun aenv f) (refineExp aenv e) (refineOpenAcc aenv a)+refineOpenAcc aenv (Scanr f e a) = Scanr (refineFun aenv f) (refineExp aenv e) (refineOpenAcc aenv a)+refineOpenAcc aenv (Permute f1 a1 f2 a2) = Permute (refineFun aenv f1) (refineOpenAcc aenv a1) (refineFun aenv f2) (refineOpenAcc aenv a2)+refineOpenAcc aenv (Backpermute e f a) = Backpermute (refineExp aenv e) (refineFun aenv f) (refineOpenAcc aenv a)+refineOpenAcc aenv (FoldSeg f e a s) = FoldSeg (refineFun aenv f) (refineExp aenv e) (refineOpenAcc aenv a) offsets+ where+ offsets = Let2 scanop (Avar (SuccIdx ZeroIdx))+ scanop = Scanl (Sugar.convertFun2 undefined Sugar.mkAdd)+ (Const (Sugar.fromElem (0::Int)))+ (refineOpenAcc aenv s)++refineOpenAcc _ acc = acc++-- Functions ??+--+refineFun :: Val aenv -> Fun aenv t -> Fun aenv t+refineFun _ = id+++-- Expressions+--+refineExp :: Val aenv -> Exp aenv a -> Exp aenv a+refineExp = refineOpenExp Empty++refineOpenExp :: Val env -> Val aenv -> OpenExp env aenv a -> OpenExp env aenv a+refineOpenExp env aenv (IndexScalar a e) = IndexScalar (refineOpenAcc aenv a) (refineOpenExp env aenv e)+refineOpenExp _ aenv (Shape a) = Shape (refineOpenAcc aenv a)+refineOpenExp _ _ exp = exp+
Data/Array/Accelerate/CUDA/State.hs view
@@ -8,24 +8,26 @@ -- Stability : experimental -- Portability : non-partable (GHC extensions) ----- This module defines a state monad which keeps track of the code generator--- state, including memory transfers and external compilation processes.+-- This module defines a state monad token which keeps track of the code+-- generator state, including memory transfers and external compilation+-- processes. -- module Data.Array.Accelerate.CUDA.State (- evalCUDA, runCUDA, CIO, unique, outputDir, deviceProps, memoryEntry, kernelEntry,- KernelEntry(KernelEntry), kernelName, kernelStatus, Key,- MemoryEntry(MemoryEntry), refcount, arena,+ evalCUDA, runCUDA, CIO, unique, outputDir, deviceProps, memoryTable, kernelTable,+ AccTable, KernelEntry(KernelEntry), kernelName, kernelStatus,+ MemTable, MemoryEntry(MemoryEntry), refcount, memsize, arena, freshVar, module Data.Record.Label ) where -import Prelude hiding (id, (.), mod)+import Prelude hiding (id, (.)) import Control.Category +import Data.Int import Data.Maybe import Data.Binary import Data.Record.Label@@ -33,10 +35,8 @@ import Control.Applicative import Control.Exception import Control.Monad.State (StateT(..), liftM)-import Data.Map (Map)-import Data.IntMap (IntMap)-import qualified Data.Map as M-import qualified Data.IntMap as IM+import Data.HashTable (HashTable)+import qualified Data.HashTable as HT import System.Directory import System.FilePath@@ -47,6 +47,51 @@ import qualified Foreign.CUDA.Driver as CUDA +-- Types+-- ~~~~~++type AccTable = HashTable String KernelEntry+type MemTable = HashTable WordPtr MemoryEntry++-- | The state token for accelerated CUDA array operations+--+type CIO = StateT CUDAState IO+data CUDAState = CUDAState+ {+ _unique :: Int,+ _outputDir :: FilePath,+ _deviceProps :: CUDA.DeviceProperties,+ _memoryTable :: MemTable,+ _kernelTable :: AccTable+ }++-- |+-- Associate an array expression with an external compilation tool (nvcc) or the+-- loaded function module+--+data KernelEntry = KernelEntry+ {+ _kernelName :: String,+ _kernelStatus :: Either ProcessID CUDA.Module+ }++-- |+-- Reference tracking for device memory allocations. Associates the products of+-- an `Array dim e' with data stored on the graphics device. Facilitates reuse+-- and delayed allocation at the cost of explicit release.+--+-- This maps to a single concrete array. Arrays of pairs, for example, which are+-- represented internally as pairs of arrays, will generate two entries.+--+data MemoryEntry = MemoryEntry+ {+ _refcount :: Int,+ _memsize :: Int64,+ _arena :: WordPtr+ }+++ -- The CUDA State Monad -- ~~~~~~~~~~~~~~~~~~~~ @@ -64,8 +109,8 @@ -- Store the kernel module map to file to the given directory ---save :: FilePath -> Map Key KernelEntry -> IO ()-save f m = encodeFile f . map (second _kernelName) . filter compiled $ M.toAscList m+save :: FilePath -> AccTable -> IO ()+save f m = encodeFile f . map (second _kernelName) . filter compiled =<< HT.toList m where compiled (_,KernelEntry _ (Right _)) = True compiled _ = False@@ -73,15 +118,18 @@ -- Read the kernel index map file (if it exists), loading modules into the -- current context ---load :: FilePath -> IO (Map Key KernelEntry)+load :: FilePath -> IO (AccTable, Int) load f = do x <- doesFileExist f- if x then M.fromDistinctAscList <$> (mapM reload =<< decodeFile f)- else return M.empty+ e <- if x then mapM reload =<< decodeFile f+ else return []++ (,length e) <$> HT.fromList HT.hashString e where reload (k,n) = (k,) . KernelEntry n . Right <$> CUDA.loadFile (n `replaceExtension` ".cubin") + -- | -- Evaluate a CUDA array computation under a newly initialised environment, -- discarding the final state.@@ -94,12 +142,13 @@ bracket (initialise Nothing) finalise $ \(dev,_ctx) -> do props <- CUDA.props dev dir <- getOutputDir- dict <- load (dir </> ".dict")- (a,s) <- runStateT acc (CUDAState (M.size dict) dir props IM.empty dict)+ tab <- HT.new (==) fromIntegral+ (k,n) <- load (dir </> "_index")+ (a,s) <- runStateT acc (CUDAState n dir props tab k) -- -- TLM 2010-06-05: assert all memory has been released ?? -- does CUDA.destroy release memory ??- save (dir </> ".dict") (_kernelEntry s)+ save (dir </> "_index") (_kernelTable s) return (a,s) where@@ -111,60 +160,7 @@ return (dev, ctx) --- | The state token for accelerated CUDA array operations----type CIO = StateT CUDAState IO-data CUDAState = CUDAState- {- _unique :: Int,- _outputDir :: FilePath,- _deviceProps :: CUDA.DeviceProperties,- _memoryEntry :: IntMap MemoryEntry,- _kernelEntry :: Map Key KernelEntry- }---- |--- Associate an array expression with an external compilation tool (nvcc) or the--- loaded function module----type Key = String-data KernelEntry = KernelEntry- {- _kernelName :: String,- _kernelStatus :: Either ProcessID CUDA.Module- }----- |--- Reference tracking for device memory allocations. Associates the products of--- an `Array dim e' with data stored on the graphics device. Facilitates reuse--- and delayed allocation at the cost of explicit release.------ This maps to a single concrete array. Arrays of pairs, for example, which are--- represented internally as pairs of arrays, will generate two entries.----data MemoryEntry = MemoryEntry- {- _refcount :: Int,- _arena :: WordPtr- }- $(mkLabels [''CUDAState, ''MemoryEntry, ''KernelEntry])----- The derived labels (documentation goes here)----unique :: CUDAState :-> Int-outputDir :: CUDAState :-> FilePath-deviceProps :: CUDAState :-> CUDA.DeviceProperties-memoryEntry :: CUDAState :-> IntMap MemoryEntry-kernelEntry :: CUDAState :-> Map Key KernelEntry--refcount :: MemoryEntry :-> Int-arena :: MemoryEntry :-> WordPtr--kernelName :: KernelEntry :-> String-kernelStatus :: KernelEntry :-> Either ProcessID CUDA.Module -- | A unique name supply
+ Data/Array/Accelerate/Internal/Check.hs view
@@ -0,0 +1,107 @@+{-# LANGUAGE CPP #-}+-- |+-- Module : Data.Array.Accelerate.Internal.Check+-- Copyright : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee, Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Manuel M T Chakravarty <chak@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--+-- Bounds checking infrastructure+--+-- Sstolen from the Vector package by Roman Leshchinskiy+-- <http://hackage.haskell.org/package/vector>+--++module Data.Array.Accelerate.Internal.Check+ (+ Checks(..), doChecks,+ error, check, assert, checkIndex, checkLength, checkSlice+ ) where++import Prelude hiding( error )+import qualified Prelude as P++data Checks = Bounds | Unsafe | Internal deriving( Eq )++doBoundsChecks :: Bool+#ifdef ACCELERATE_BOUNDS_CHECKS+doBoundsChecks = True+#else+doBoundsChecks = False+#endif++doUnsafeChecks :: Bool+#ifdef ACCELERATE_UNSAFE_CHECKS+doUnsafeChecks = True+#else+doUnsafeChecks = False+#endif++doInternalChecks :: Bool+#ifdef ACCELERATE_INTERNAL_CHECKS+doInternalChecks = True+#else+doInternalChecks = False+#endif+++doChecks :: Checks -> Bool+{-# INLINE doChecks #-}+doChecks Bounds = doBoundsChecks+doChecks Unsafe = doUnsafeChecks+doChecks Internal = doInternalChecks++error :: String -> Int -> Checks -> String -> String -> a+error file line kind loc msg+ = P.error . unlines $+ (if kind == Internal+ then (["*** Internal error in package accelerate ***"+ ,"*** Please submit a bug report at http://trac.haskell.org/accelerate"]++)+ else id)+ [ file ++ ":" ++ show line ++ " (" ++ loc ++ "): " ++ msg ]++check :: String -> Int -> Checks -> String -> String -> Bool -> a -> a+{-# INLINE check #-}+check file line kind loc msg cond x+ | not (doChecks kind) || cond = x+ | otherwise = error file line kind loc msg++assert_msg :: String+assert_msg = "assertion failure"++assert :: String -> Int -> Checks -> String -> Bool -> a -> a+{-# INLINE assert #-}+assert file line kind loc = check file line kind loc assert_msg++checkIndex_msg :: Int -> Int -> String+{-# NOINLINE checkIndex_msg #-}+checkIndex_msg i n = "index out of bounds " ++ show (i,n)++checkIndex :: String -> Int -> Checks -> String -> Int -> Int -> a -> a+{-# INLINE checkIndex #-}+checkIndex file line kind loc i n x+ = check file line kind loc (checkIndex_msg i n) (i >= 0 && i<n) x+++checkLength_msg :: Int -> String+{-# NOINLINE checkLength_msg #-}+checkLength_msg n = "negative length " ++ show n++checkLength :: String -> Int -> Checks -> String -> Int -> a -> a+{-# INLINE checkLength #-}+checkLength file line kind loc n x+ = check file line kind loc (checkLength_msg n) (n >= 0) x+++checkSlice_msg :: Int -> Int -> Int -> String+{-# NOINLINE checkSlice_msg #-}+checkSlice_msg i m n = "invalid slice " ++ show (i,m,n)++checkSlice :: String -> Int -> Checks -> String -> Int -> Int -> Int -> a -> a+{-# INLINE checkSlice #-}+checkSlice file line kind loc i m n x+ = check file line kind loc (checkSlice_msg i m n)+ (i >= 0 && m >= 0 && i+m <= n) x+
Data/Array/Accelerate/Interpreter.hs view
@@ -13,6 +13,15 @@ -- This interpreter is meant to be a reference implementation of the semantics -- of the embedded array language. The emphasis is on defining the semantics -- clearly, not on performance.+--+-- /Surface types versus representation types/+--+-- As a general rule, we perform all computations on representation types and we store all data+-- as values of representation types. To guarantee the type safety of the interpreter, this+-- currently implies a lot of conversions between surface and representation types. Optimising+-- the code by eliminating back and forth conversions is fine, but only where it doesn't+-- negatively affects clarity — after all, the main purpose of the interpreter is to serve as an+-- executable specification. module Data.Array.Accelerate.Interpreter ( @@ -114,6 +123,12 @@ evalOpenAcc (Backpermute e p acc) aenv = backpermuteOp (evalExp e aenv) (evalFun p aenv) (evalOpenAcc acc aenv) +evalOpenAcc (Stencil sten bndy acc) aenv+ = stencilOp (evalFun sten aenv) bndy (evalOpenAcc acc aenv)++evalOpenAcc (Stencil2 sten bndy1 acc1 bndy2 acc2) aenv+ = stencil2Op (evalFun sten aenv) bndy1 (evalOpenAcc acc1 aenv) bndy2 (evalOpenAcc acc2 aenv)+ -- Evaluate a closed array expressions -- evalAcc :: Delayable a => Acc a -> Delayed a@@ -323,7 +338,50 @@ backpermuteOp sh' p (DelayedArray _sh rf) = DelayedArray (Sugar.fromElem sh') (rf . Sugar.sinkFromElem p) +stencilOp :: forall dim e e' stencil. (Sugar.Elem e, Sugar.Elem e', Stencil dim e stencil)+ => (stencil -> e')+ -> Boundary (Sugar.ElemRepr e)+ -> Delayed (Array dim e)+ -> Delayed (Array dim e')+stencilOp sten bndy (DelayedArray sh rf)+ = DelayedArray sh rf'+ where+ rf' = Sugar.sinkFromElem (sten . stencilAccess rfBounded) + -- add a boundary to the source array as specified by the boundary condition+ rfBounded :: dim -> e+ rfBounded ix = Sugar.toElem $ case Sugar.bound (Sugar.toElem sh) ix bndy of+ Left v -> v+ Right ix' -> rf (Sugar.fromElem ix')++stencil2Op :: forall dim e1 e2 e' stencil1 stencil2. + (Sugar.Elem e1, Sugar.Elem e2, Sugar.Elem e', + Stencil dim e1 stencil1, Stencil dim e2 stencil2)+ => (stencil1 -> stencil2 -> e')+ -> Boundary (Sugar.ElemRepr e1)+ -> Delayed (Array dim e1)+ -> Boundary (Sugar.ElemRepr e2)+ -> Delayed (Array dim e2)+ -> Delayed (Array dim e')+stencil2Op sten bndy1 (DelayedArray sh1 rf1) bndy2 (DelayedArray sh2 rf2)+ = DelayedArray (sh1 `intersect` sh2) rf'+ where+ rf' = Sugar.sinkFromElem (\ix -> sten (stencilAccess rf1Bounded ix)+ (stencilAccess rf2Bounded ix))++ -- add a boundary to the source arrays as specified by the boundary conditions+ + rf1Bounded :: dim -> e1+ rf1Bounded ix = Sugar.toElem $ case Sugar.bound (Sugar.toElem sh1) ix bndy1 of+ Left v -> v+ Right ix' -> rf1 (Sugar.fromElem ix')++ rf2Bounded :: dim -> e2+ rf2Bounded ix = Sugar.toElem $ case Sugar.bound (Sugar.toElem sh2) ix bndy2 of+ Left v -> v+ Right ix' -> rf2 (Sugar.fromElem ix')++ -- Expression evaluation -- --------------------- @@ -431,6 +489,7 @@ evalPrim (PrimLog ty) = evalLog ty evalPrim (PrimFPow ty) = evalFPow ty evalPrim (PrimLogBase ty) = evalLogBase ty+evalPrim (PrimAtan2 ty) = evalAtan2 ty evalPrim (PrimLt ty) = evalLt ty evalPrim (PrimGt ty) = evalGt ty evalPrim (PrimLtEq ty) = evalLtEq ty@@ -562,6 +621,9 @@ evalLogBase :: FloatingType a -> ((a, a) -> a) evalLogBase ty | FloatingDict <- floatingDict ty = uncurry logBase++evalAtan2 :: FloatingType a -> ((a, a) -> a)+evalAtan2 ty | FloatingDict <- floatingDict ty = uncurry atan2 -- Methods of Num
Data/Array/Accelerate/Language.hs view
@@ -1,14 +1,14 @@ {-# LANGUAGE FlexibleContexts, TypeFamilies, RankNTypes, ScopedTypeVariables #-}-{-# LANGUAGE FlexibleInstances #-}+{-# LANGUAGE FlexibleInstances, MultiParamTypeClasses #-} {-# OPTIONS_GHC -fno-warn-missing-methods #-} --- |Embedded array processing language: user-visible language------ Copyright (c) 2009 Manuel M T Chakravarty, Gabriele Keller, Sean Lee------ License: BSD3+-- Module : Data.Array.Accelerate.Language+-- Copyright : [2009..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee+-- License : BSD3 ------ Description ---------------------------------------------------------------+-- Maintainer : Manuel M T Chakravarty <chak@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions) -- -- We use the dictionary view of overloaded operations (such as arithmetic and -- bit manipulation) to reify such expressions. With non-overloaded@@ -21,7 +21,16 @@ -- ** Array and scalar expressions Acc, Exp, -- re-exporting from 'Smart'+ + -- ** Stencil specification+ Boundary(..), Stencil, -- re-exporting from 'Smart' + -- ** Common stencil types+ Stencil3, Stencil5, Stencil7, Stencil9,+ Stencil3x3, Stencil5x3, Stencil3x5, Stencil5x5,+ Stencil3x3x3, Stencil5x3x3, Stencil3x5x3, Stencil3x3x5, Stencil5x5x3, Stencil5x3x5,+ Stencil3x5x5, Stencil5x5x5,+ -- ** Scalar introduction constant, -- re-exporting from 'Smart' @@ -33,7 +42,7 @@ -- ** Collective array operations slice, replicate, zip, unzip, map, zipWith, scanl, scanr, fold, foldSeg,- permute, backpermute,+ permute, backpermute, stencil, stencil2, -- ** Tuple construction and destruction Tuple(..), fst, snd, curry, uncurry,@@ -113,7 +122,7 @@ -- yields a three dimensional array, where 'arr' is replicated twice across the -- first and three times across the third dimension. ---replicate :: forall slix e. (SliceIx slix, Elem e) +replicate :: (SliceIx slix, Elem e) => Exp slix -> Acc (Array (Slice slix) e) -> Acc (Array (SliceDim slix) e)@@ -123,7 +132,7 @@ -- argument). The result is a new array (possibly a singleton) containing -- all dimensions in their entirety. ---slice :: forall slix e. (SliceIx slix, Elem e) +slice :: (SliceIx slix, Elem e) => Acc (Array (SliceDim slix) e) -> Exp slix -> Acc (Array (Slice slix) e)@@ -154,7 +163,8 @@ -> Acc (Array dim b) map = Map --- |Apply the given binary function elementwise to the two arrays.+-- |Apply the given binary function elementwise to the two arrays. The extent of the resulting+-- array is the intersection of the extents of the two source arrays. -- zipWith :: (Ix dim, Elem a, Elem b, Elem c) => (Exp a -> Exp b -> Exp c) @@ -163,7 +173,7 @@ -> Acc (Array dim c) zipWith = ZipWith --- |Prescan of a vector. The type \'a\' together with the binary function+-- |Prescan of a vector. The type 'a' together with the binary function -- (first argument) and value (second argument) must form a monoid; i.e., the -- function must be /associative/ and the value must be its /neutral element/. --@@ -224,6 +234,7 @@ permute = Permute -- |Backward permutation +-- backpermute :: (Ix dim, Ix dim', Elem a) => Exp dim' -- ^shape of the result array -> (Exp dim' -> Exp dim) -- ^permutation@@ -232,6 +243,64 @@ backpermute = Backpermute +-- Common stencil types+--++-- DIM1 stencil type+type Stencil3 a = (Exp a, Exp a, Exp a)+type Stencil5 a = (Exp a, Exp a, Exp a, Exp a, Exp a)+type Stencil7 a = (Exp a, Exp a, Exp a, Exp a, Exp a, Exp a, Exp a)+type Stencil9 a = (Exp a, Exp a, Exp a, Exp a, Exp a, Exp a, Exp a, Exp a, Exp a)++-- DIM2 stencil type+type Stencil3x3 a = (Stencil3 a, Stencil3 a, Stencil3 a)+type Stencil5x3 a = (Stencil5 a, Stencil5 a, Stencil5 a)+type Stencil3x5 a = (Stencil3 a, Stencil3 a, Stencil3 a, Stencil3 a, Stencil3 a)+type Stencil5x5 a = (Stencil5 a, Stencil5 a, Stencil5 a, Stencil5 a, Stencil5 a)++-- DIM3 stencil type+type Stencil3x3x3 a = (Stencil3x3 a, Stencil3x3 a, Stencil3x3 a)+type Stencil5x3x3 a = (Stencil5x3 a, Stencil5x3 a, Stencil5x3 a)+type Stencil3x5x3 a = (Stencil3x5 a, Stencil3x5 a, Stencil3x5 a)+type Stencil3x3x5 a = (Stencil3x3 a, Stencil3x3 a, Stencil3x3 a, Stencil3x3 a, Stencil3x3 a)+type Stencil5x5x3 a = (Stencil5x5 a, Stencil5x5 a, Stencil5x5 a)+type Stencil5x3x5 a = (Stencil5x3 a, Stencil5x3 a, Stencil5x3 a, Stencil5x3 a, Stencil5x3 a)+type Stencil3x5x5 a = (Stencil3x5 a, Stencil3x5 a, Stencil3x5 a, Stencil3x5 a, Stencil3x5 a)+type Stencil5x5x5 a = (Stencil5x5 a, Stencil5x5 a, Stencil5x5 a, Stencil5x5 a, Stencil5x5 a)++-- |Map a stencil over an array. In contrast to 'map', the domain of a stencil function is an+-- entire /neighbourhood/ of each array element. Neighbourhoods are sub-arrays centred around a+-- focal point. They are not necessarily rectangular, but they are symmetric in each dimension+-- and have an extent of at least three in each dimensions — due to the symmetry requirement, the+-- extent is necessarily odd. The focal point is the array position that is determined by the+-- stencil.+--+-- For those array positions where the neighbourhood extends past the boundaries of the source+-- array, a boundary condition determines the contents of the out-of-bounds neighbourhood+-- positions.+--+stencil :: (Ix dim, Elem a, Elem b, Stencil dim a stencil)+ => (stencil -> Exp b) -- ^stencil function+ -> Boundary a -- ^boundary condition+ -> Acc (Array dim a) -- ^source array+ -> Acc (Array dim b) -- ^destination array+stencil = Stencil++-- |Map a binary stencil of an array. The extent of the resulting array is the intersection of+-- the extents of the two source arrays.+--+stencil2 :: (Ix dim, Elem a, Elem b, Elem c, + Stencil dim a stencil1, + Stencil dim b stencil2)+ => (stencil1 -> stencil2 -> Exp c) -- ^binary stencil function+ -> Boundary a -- ^boundary condition #1+ -> Acc (Array dim a) -- ^source array #1+ -> Boundary b -- ^boundary condition #2+ -> Acc (Array dim b) -- ^source array #2+ -> Acc (Array dim c) -- ^destination array+stencil2 = Stencil2++ -- Tuples -- ------ @@ -411,7 +480,8 @@ instance (Elem t, IsFloating t) => RealFrac (Exp t) -- FIXME: add ops -instance (Elem t, IsFloating t) => RealFloat (Exp t)+instance (Elem t, IsFloating t) => RealFloat (Exp t) where+ atan2 = mkAtan2 -- FIXME: add ops
Data/Array/Accelerate/Pretty.hs view
@@ -23,6 +23,7 @@ import Data.Array.Accelerate.Array.Sugar import Data.Array.Accelerate.Tuple import Data.Array.Accelerate.AST+import Data.Array.Accelerate.Type -- |Show instances@@ -48,7 +49,7 @@ = text "let a" <> int lvl <+> char '=' <+> prettyAcc lvl acc1 <+> text " in " <+> prettyAcc (lvl + 1) acc2 prettyAcc lvl (Let2 acc1 acc2) - = text "let (a" <> int lvl <> text ", a" <> int (lvl + 1) <+> char '=' <+>+ = text "let (a" <> int lvl <> text ", a" <> int (lvl + 1) <> char ')' <+> char '=' <+> prettyAcc lvl acc1 <+> text " in " <+> prettyAcc (lvl + 2) acc2 prettyAcc _ (Avar idx) = text $ "a" ++ show (idxToInt idx)@@ -86,6 +87,25 @@ parens (prettyFun lvl p), prettyAccParens lvl acc] +prettyAcc lvl (Stencil sten bndy acc)+ = prettyArrOp "stencil" [parens (prettyFun lvl sten),+ prettyBoundary acc bndy,+ prettyAccParens lvl acc]++prettyAcc lvl (Stencil2 sten bndy1 acc1 bndy2 acc2)+ = prettyArrOp "stencil2" [parens (prettyFun lvl sten),+ prettyBoundary acc1 bndy1,+ prettyAccParens lvl acc1,+ prettyBoundary acc2 bndy2,+ prettyAccParens lvl acc2]++prettyBoundary :: forall aenv dim e. Elem e + => {-dummy-}OpenAcc aenv (Array dim e) -> Boundary (ElemRepr e) -> Doc+prettyBoundary _ Clamp = text "Clamp"+prettyBoundary _ Mirror = text "Mirror"+prettyBoundary _ Wrap = text "Wrap"+prettyBoundary _ (Constant e) = parens $ text "Wrap" <+> text (show (toElem e :: e))+ prettyArrOp :: String -> [Doc] -> Doc prettyArrOp name docs = hang (text name) 2 $ sep docs @@ -195,6 +215,7 @@ prettyPrim (PrimLog _) = text "log" prettyPrim (PrimFPow _) = text "(**)" prettyPrim (PrimLogBase _) = text "logBase"+prettyPrim (PrimAtan2 _) = text "atan2" prettyPrim (PrimLt _) = text "(<*)" prettyPrim (PrimGt _) = text "(>*)" prettyPrim (PrimLtEq _) = text "(<=*)"
Data/Array/Accelerate/Smart.hs view
@@ -1,26 +1,27 @@ {-# LANGUAGE GADTs, TypeFamilies, ScopedTypeVariables, FlexibleContexts #-}-{-# LANGUAGE FlexibleInstances #-}+{-# LANGUAGE FlexibleInstances, MultiParamTypeClasses, TypeSynonymInstances #-} --- |Embedded array processing language: smart expression constructors------ Copyright (c) [2008..2009] Manuel M T Chakravarty, Gabriele Keller, Sean Lee------ License: BSD3+-- Module : Data.Array.Accelerate.Smart+-- Copyright : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee+-- License : BSD3 ------ Description ---------------------------------------------------------------+-- Maintainer : Manuel M T Chakravarty <chak@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions) ----- This modules defines the AST of the user-visible embedded language using--- more convenient higher-order abstract syntax (instead of de Bruijn--- indices). Moreover, it defines smart constructors to construct programs.+-- This modules defines the AST of the user-visible embedded language using+-- more convenient higher-order abstract syntax (instead of de Bruijn+-- indices). Moreover, it defines smart constructors to construct programs. module Data.Array.Accelerate.Smart ( -- * HOAS AST- Acc(..), Exp(..), + Acc(..), Exp(..), Boundary(..), Stencil(..), -- * HOAS -> de Bruijn conversion- convertAcc, convertClosedExp,- + convertAcc,+ convertExp, convertFun1, convertFun2,+ -- * Smart constructors for unpairing unpair, @@ -37,6 +38,7 @@ mkAsinh, mkAcosh, mkAtanh, mkExpFloating, mkSqrt, mkLog, mkFPow, mkLogBase,+ mkAtan2, -- * Smart constructors for primitive functions mkAdd, mkSub, mkMul, mkNeg, mkAbs, mkSig, mkQuot, mkRem, mkIDiv, mkMod,@@ -56,7 +58,7 @@ import Data.Array.Accelerate.Array.Sugar import Data.Array.Accelerate.Tuple hiding (Tuple) import qualified Data.Array.Accelerate.Tuple as Tuple-import Data.Array.Accelerate.AST hiding (OpenAcc(..), Acc, OpenExp(..), Exp)+import Data.Array.Accelerate.AST hiding (OpenAcc(..), Acc, Stencil, OpenExp(..), Exp) import qualified Data.Array.Accelerate.AST as AST import Data.Array.Accelerate.Pretty () @@ -68,10 +70,10 @@ -- data Acc a where - FstArray :: (Elem e1, Elem e2)+ FstArray :: (Ix dim1, Elem e1, Elem e2) => Acc (Array dim1 e1, Array dim2 e2) -> Acc (Array dim1 e1)- SndArray :: (Elem e1, Elem e2)+ SndArray :: (Ix dim2, Elem e1, Elem e2) => Acc (Array dim1 e1, Array dim2 e2) -> Acc (Array dim2 e2) @@ -132,6 +134,19 @@ -> (Exp dim' -> Exp dim) -> Acc (Array dim e) -> Acc (Array dim' e)+ Stencil :: (Ix dim, Elem a, Elem b, Stencil dim a stencil)+ => (stencil -> Exp b)+ -> Boundary a+ -> Acc (Array dim a)+ -> Acc (Array dim b)+ Stencil2 :: (Ix dim, Elem a, Elem b, Elem c, + Stencil dim a stencil1, Stencil dim b stencil2)+ => (stencil1 -> stencil2 -> Exp c)+ -> Boundary a+ -> Acc (Array dim a)+ -> Boundary b+ -> Acc (Array dim b)+ -> Acc (Array dim c) -- |Conversion from HOAS to de Bruijn computation AST@@ -178,7 +193,25 @@ = AST.Backpermute (convertExp alyt newDim) (convertFun1 alyt perm) (convertOpenAcc alyt acc)+convertOpenAcc alyt (Stencil stencil boundary acc) + = AST.Stencil (convertStencilFun acc alyt stencil) + (convertBoundary boundary) + (convertOpenAcc alyt acc)+convertOpenAcc alyt (Stencil2 stencil bndy1 acc1 bndy2 acc2) + = AST.Stencil2 (convertStencilFun2 acc1 acc2 alyt stencil) + (convertBoundary bndy1) + (convertOpenAcc alyt acc1)+ (convertBoundary bndy2) + (convertOpenAcc alyt acc2) +-- |Convert a boundary condition+--+convertBoundary :: Elem e => Boundary e -> Boundary (ElemRepr e)+convertBoundary Clamp = Clamp+convertBoundary Mirror = Mirror+convertBoundary Wrap = Wrap+convertBoundary (Constant e) = Constant (fromElem e)+ -- |Convert a closed array expression -- convertAcc :: Acc a -> AST.Acc a@@ -316,7 +349,52 @@ (ZeroIdx :: Idx (((), ElemRepr a), ElemRepr b) (ElemRepr b)) openF = convertOpenExp lyt alyt (f a b) +-- Convert a unary stencil function+--+convertStencilFun :: forall dim a stencil b aenv. (Elem a, Stencil dim a stencil)+ => Acc (Array dim a) -- just passed to fix the type variables+ -> Layout aenv aenv + -> (stencil -> Exp b)+ -> AST.Fun aenv (StencilRepr dim stencil -> b)+convertStencilFun _ alyt stencilFun = Lam (Body openStencilFun)+ where+ stencil = Tag 0 :: Exp (StencilRepr dim stencil)+ lyt = EmptyLayout + `PushLayout` + (ZeroIdx :: Idx ((), ElemRepr (StencilRepr dim stencil)) + (ElemRepr (StencilRepr dim stencil)))+ openStencilFun = convertOpenExp lyt alyt $+ stencilFun (stencilPrj (undefined::dim) (undefined::a) stencil) +-- Convert a binary stencil function+--+convertStencilFun2 :: forall dim a b stencil1 stencil2 c aenv. + (Elem a, Stencil dim a stencil1,+ Elem b, Stencil dim b stencil2)+ => Acc (Array dim a) -- just passed to fix the type variables+ -> Acc (Array dim b) -- just passed to fix the type variables+ -> Layout aenv aenv + -> (stencil1 -> stencil2 -> Exp c)+ -> AST.Fun aenv (StencilRepr dim stencil1 ->+ StencilRepr dim stencil2 -> c)+convertStencilFun2 _ _ alyt stencilFun = Lam (Lam (Body openStencilFun))+ where+ stencil1 = Tag 1 :: Exp (StencilRepr dim stencil1)+ stencil2 = Tag 0 :: Exp (StencilRepr dim stencil2)+ lyt = EmptyLayout + `PushLayout` + (SuccIdx ZeroIdx :: Idx (((), ElemRepr (StencilRepr dim stencil1)),+ ElemRepr (StencilRepr dim stencil2)) + (ElemRepr (StencilRepr dim stencil1)))+ `PushLayout` + (ZeroIdx :: Idx (((), ElemRepr (StencilRepr dim stencil1)),+ ElemRepr (StencilRepr dim stencil2)) + (ElemRepr (StencilRepr dim stencil2)))+ openStencilFun = convertOpenExp lyt alyt $+ stencilFun (stencilPrj (undefined::dim) (undefined::a) stencil1)+ (stencilPrj (undefined::dim) (undefined::b) stencil2)++ -- Pretty printing -- @@ -349,9 +427,186 @@ slix = undefined :: slix --- |Smart constructors to construct HOAS AST expressions--- -----------------------------------------------------+-- |Smart constructors for stencil reification+-- ------------------------------------------- +-- Stencil reification+--+-- In the AST representation, we turn the stencil type from nested tuples of Accelerate expressions+-- into an Accelerate expression whose type is a tuple nested in the same manner. This enables us+-- to represent the stencil function as a unary function (which also only needs one de Bruijn+-- index). The various positions in the stencil are accessed via tuple indices (i.e., projections).++class (Elem (StencilRepr dim stencil), AST.Stencil dim a (StencilRepr dim stencil)) + => Stencil dim a stencil where+ type StencilRepr dim stencil :: *+ stencilPrj :: dim{-dummy-} -> a{-dummy-} -> Exp (StencilRepr dim stencil) -> stencil+ +-- DIM1+instance Elem a => Stencil DIM1 a (Exp a, Exp a, Exp a) where+ type StencilRepr DIM1 (Exp a, Exp a, Exp a) = (a, a, a)+ stencilPrj _ _ s = (Prj tix2 s, Prj tix1 s, Prj tix0 s)+instance Elem a => Stencil DIM1 a (Exp a, Exp a, Exp a, Exp a, Exp a) where+ type StencilRepr DIM1 (Exp a, Exp a, Exp a, Exp a, Exp a) = (a, a, a, a, a)+ stencilPrj _ _ s = (Prj tix4 s, Prj tix3 s, Prj tix2 s, Prj tix1 s, Prj tix0 s)+instance Elem a => Stencil DIM1 a (Exp a, Exp a, Exp a, Exp a, Exp a, Exp a, Exp a) where+ type StencilRepr DIM1 (Exp a, Exp a, Exp a, Exp a, Exp a, Exp a, Exp a) = (a, a, a, a, a, a, a)+ stencilPrj _ _ s = (Prj tix6 s, Prj tix5 s, Prj tix4 s, Prj tix3 s, Prj tix2 s, Prj tix1 s, + Prj tix0 s)+instance Elem a + => Stencil DIM1 a (Exp a, Exp a, Exp a, Exp a, Exp a, Exp a, Exp a, Exp a, Exp a) where+ type StencilRepr DIM1 (Exp a, Exp a, Exp a, Exp a, Exp a, Exp a, Exp a, Exp a, Exp a) + = (a, a, a, a, a, a, a, a, a)+ stencilPrj _ _ s = (Prj tix8 s, Prj tix7 s, Prj tix6 s, Prj tix5 s, Prj tix4 s, Prj tix3 s,+ Prj tix2 s, Prj tix1 s, Prj tix0 s)++-- DIM2+instance (Stencil DIM1 a row2, + Stencil DIM1 a row1,+ Stencil DIM1 a row0) => Stencil DIM2 a (row2, row1, row0) where+ type StencilRepr DIM2 (row2, row1, row0) + = (StencilRepr DIM1 row2, StencilRepr DIM1 row1, StencilRepr DIM1 row0)+ stencilPrj _ a s = (stencilPrj (undefined::DIM1) a (Prj tix2 s), + stencilPrj (undefined::DIM1) a (Prj tix1 s), + stencilPrj (undefined::DIM1) a (Prj tix0 s))+instance (Stencil DIM1 a row1,+ Stencil DIM1 a row2,+ Stencil DIM1 a row3,+ Stencil DIM1 a row4,+ Stencil DIM1 a row5) => Stencil DIM2 a (row1, row2, row3, row4, row5) where+ type StencilRepr DIM2 (row1, row2, row3, row4, row5) + = (StencilRepr DIM1 row1, StencilRepr DIM1 row2, StencilRepr DIM1 row3, StencilRepr DIM1 row4,+ StencilRepr DIM1 row5)+ stencilPrj _ a s = (stencilPrj (undefined::DIM1) a (Prj tix4 s), + stencilPrj (undefined::DIM1) a (Prj tix3 s), + stencilPrj (undefined::DIM1) a (Prj tix2 s), + stencilPrj (undefined::DIM1) a (Prj tix1 s), + stencilPrj (undefined::DIM1) a (Prj tix0 s))+instance (Stencil DIM1 a row1,+ Stencil DIM1 a row2,+ Stencil DIM1 a row3,+ Stencil DIM1 a row4,+ Stencil DIM1 a row5,+ Stencil DIM1 a row6,+ Stencil DIM1 a row7) => Stencil DIM2 a (row1, row2, row3, row4, row5, row6, row7) where+ type StencilRepr DIM2 (row1, row2, row3, row4, row5, row6, row7) + = (StencilRepr DIM1 row1, StencilRepr DIM1 row2, StencilRepr DIM1 row3, StencilRepr DIM1 row4,+ StencilRepr DIM1 row5, StencilRepr DIM1 row6, StencilRepr DIM1 row7)+ stencilPrj _ a s = (stencilPrj (undefined::DIM1) a (Prj tix6 s), + stencilPrj (undefined::DIM1) a (Prj tix5 s), + stencilPrj (undefined::DIM1) a (Prj tix4 s), + stencilPrj (undefined::DIM1) a (Prj tix3 s), + stencilPrj (undefined::DIM1) a (Prj tix2 s), + stencilPrj (undefined::DIM1) a (Prj tix1 s), + stencilPrj (undefined::DIM1) a (Prj tix0 s))+instance (Stencil DIM1 a row1,+ Stencil DIM1 a row2,+ Stencil DIM1 a row3,+ Stencil DIM1 a row4,+ Stencil DIM1 a row5,+ Stencil DIM1 a row6,+ Stencil DIM1 a row7,+ Stencil DIM1 a row8,+ Stencil DIM1 a row9) + => Stencil DIM2 a (row1, row2, row3, row4, row5, row6, row7, row8, row9) where+ type StencilRepr DIM2 (row1, row2, row3, row4, row5, row6, row7, row8, row9) + = (StencilRepr DIM1 row1, StencilRepr DIM1 row2, StencilRepr DIM1 row3, StencilRepr DIM1 row4,+ StencilRepr DIM1 row5, StencilRepr DIM1 row6, StencilRepr DIM1 row7, StencilRepr DIM1 row8,+ StencilRepr DIM1 row9)+ stencilPrj _ a s = (stencilPrj (undefined::DIM1) a (Prj tix8 s), + stencilPrj (undefined::DIM1) a (Prj tix7 s), + stencilPrj (undefined::DIM1) a (Prj tix6 s), + stencilPrj (undefined::DIM1) a (Prj tix5 s), + stencilPrj (undefined::DIM1) a (Prj tix4 s), + stencilPrj (undefined::DIM1) a (Prj tix3 s), + stencilPrj (undefined::DIM1) a (Prj tix2 s), + stencilPrj (undefined::DIM1) a (Prj tix1 s), + stencilPrj (undefined::DIM1) a (Prj tix0 s))++-- DIM3+instance (Stencil DIM2 a row1, + Stencil DIM2 a row2,+ Stencil DIM2 a row3) => Stencil DIM3 a (row1, row2, row3) where+ type StencilRepr DIM3 (row1, row2, row3) + = (StencilRepr DIM2 row1, StencilRepr DIM2 row2, StencilRepr DIM2 row3)+ stencilPrj _ a s = (stencilPrj (undefined::DIM2) a (Prj tix2 s), + stencilPrj (undefined::DIM2) a (Prj tix1 s), + stencilPrj (undefined::DIM2) a (Prj tix0 s))+instance (Stencil DIM2 a row1,+ Stencil DIM2 a row2,+ Stencil DIM2 a row3,+ Stencil DIM2 a row4,+ Stencil DIM2 a row5) => Stencil DIM3 a (row1, row2, row3, row4, row5) where+ type StencilRepr DIM3 (row1, row2, row3, row4, row5) + = (StencilRepr DIM2 row1, StencilRepr DIM2 row2, StencilRepr DIM2 row3, StencilRepr DIM2 row4,+ StencilRepr DIM2 row5)+ stencilPrj _ a s = (stencilPrj (undefined::DIM2) a (Prj tix4 s), + stencilPrj (undefined::DIM2) a (Prj tix3 s), + stencilPrj (undefined::DIM2) a (Prj tix2 s), + stencilPrj (undefined::DIM2) a (Prj tix1 s), + stencilPrj (undefined::DIM2) a (Prj tix0 s))+instance (Stencil DIM2 a row1,+ Stencil DIM2 a row2,+ Stencil DIM2 a row3,+ Stencil DIM2 a row4,+ Stencil DIM2 a row5,+ Stencil DIM2 a row6,+ Stencil DIM2 a row7) => Stencil DIM3 a (row1, row2, row3, row4, row5, row6, row7) where+ type StencilRepr DIM3 (row1, row2, row3, row4, row5, row6, row7) + = (StencilRepr DIM2 row1, StencilRepr DIM2 row2, StencilRepr DIM2 row3, StencilRepr DIM2 row4,+ StencilRepr DIM2 row5, StencilRepr DIM2 row6, StencilRepr DIM2 row7)+ stencilPrj _ a s = (stencilPrj (undefined::DIM2) a (Prj tix6 s), + stencilPrj (undefined::DIM2) a (Prj tix5 s), + stencilPrj (undefined::DIM2) a (Prj tix4 s), + stencilPrj (undefined::DIM2) a (Prj tix3 s), + stencilPrj (undefined::DIM2) a (Prj tix2 s), + stencilPrj (undefined::DIM2) a (Prj tix1 s), + stencilPrj (undefined::DIM2) a (Prj tix0 s))+instance (Stencil DIM2 a row1,+ Stencil DIM2 a row2,+ Stencil DIM2 a row3,+ Stencil DIM2 a row4,+ Stencil DIM2 a row5,+ Stencil DIM2 a row6,+ Stencil DIM2 a row7,+ Stencil DIM2 a row8,+ Stencil DIM2 a row9) + => Stencil DIM3 a (row1, row2, row3, row4, row5, row6, row7, row8, row9) where+ type StencilRepr DIM3 (row1, row2, row3, row4, row5, row6, row7, row8, row9) + = (StencilRepr DIM2 row1, StencilRepr DIM2 row2, StencilRepr DIM2 row3, StencilRepr DIM2 row4,+ StencilRepr DIM2 row5, StencilRepr DIM2 row6, StencilRepr DIM2 row7, StencilRepr DIM2 row8,+ StencilRepr DIM2 row9)+ stencilPrj _ a s = (stencilPrj (undefined::DIM2) a (Prj tix8 s), + stencilPrj (undefined::DIM2) a (Prj tix7 s), + stencilPrj (undefined::DIM2) a (Prj tix6 s), + stencilPrj (undefined::DIM2) a (Prj tix5 s), + stencilPrj (undefined::DIM2) a (Prj tix4 s), + stencilPrj (undefined::DIM2) a (Prj tix3 s), + stencilPrj (undefined::DIM2) a (Prj tix2 s), + stencilPrj (undefined::DIM2) a (Prj tix1 s), + stencilPrj (undefined::DIM2) a (Prj tix0 s))++-- Auxilliary tuple index constants+--+tix0 :: Elem s => TupleIdx (t, s) s+tix0 = ZeroTupIdx+tix1 :: Elem s => TupleIdx ((t, s), s1) s+tix1 = SuccTupIdx tix0+tix2 :: Elem s => TupleIdx (((t, s), s1), s2) s+tix2 = SuccTupIdx tix1+tix3 :: Elem s => TupleIdx ((((t, s), s1), s2), s3) s+tix3 = SuccTupIdx tix2+tix4 :: Elem s => TupleIdx (((((t, s), s1), s2), s3), s4) s+tix4 = SuccTupIdx tix3+tix5 :: Elem s => TupleIdx ((((((t, s), s1), s2), s3), s4), s5) s+tix5 = SuccTupIdx tix4+tix6 :: Elem s => TupleIdx (((((((t, s), s1), s2), s3), s4), s5), s6) s+tix6 = SuccTupIdx tix5+tix7 :: Elem s => TupleIdx ((((((((t, s), s1), s2), s3), s4), s5), s6), s7) s+tix7 = SuccTupIdx tix6+tix8 :: Elem s => TupleIdx (((((((((t, s), s1), s2), s3), s4), s5), s6), s7), s8) s+tix8 = SuccTupIdx tix7+ -- Pushes the 'Acc' constructor through a pair -- unpair :: (Ix dim1, Ix dim2, Elem e1, Elem e2)@@ -424,6 +679,9 @@ mkPi :: (Elem r, IsFloating r) => Exp r mkPi = PrimConst (PrimPi floatingType) +-- Operators from Floating+--+ mkSin :: (Elem t, IsFloating t) => Exp t -> Exp t mkSin x = PrimSin floatingType `PrimApp` x @@ -534,6 +792,9 @@ mkRecip :: (Elem t, IsFloating t) => Exp t -> Exp t mkRecip x = PrimRecip floatingType `PrimApp` x++mkAtan2 :: (Elem t, IsFloating t) => Exp t -> Exp t -> Exp t+mkAtan2 x y = PrimAtan2 floatingType `PrimApp` tup2 (x, y) -- FIXME: add operations from Floating, RealFrac & RealFloat
+ Data/Array/Accelerate/Test.hs view
@@ -0,0 +1,18 @@+-- |+-- Module : Data.Array.Accelerate.Test+-- Copyright : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee, Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Manuel M T Chakravarty <chak@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.Test+ (+ module Data.Array.Accelerate.Test.QuickCheck+ )+ where++import Data.Array.Accelerate.Test.QuickCheck+
+ Data/Array/Accelerate/Test/QuickCheck.hs view
@@ -0,0 +1,18 @@+-- |+-- Module : Data.Array.Accelerate.Test.QuickCheck+-- Copyright : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee, Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Manuel M T Chakravarty <chak@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.Test.QuickCheck+ (+ module Data.Array.Accelerate.Test.QuickCheck.Arbitrary+ )+ where++import Data.Array.Accelerate.Test.QuickCheck.Arbitrary+
+ Data/Array/Accelerate/Test/QuickCheck/Arbitrary.hs view
@@ -0,0 +1,236 @@+{-# LANGUAGE NoMonomorphismRestriction, ScopedTypeVariables #-}+{-# LANGUAGE FlexibleContexts, FlexibleInstances, TypeSynonymInstances #-}+{-# OPTIONS_GHC -fno-warn-orphans #-}+-- |+-- Module : Data.Array.Accelerate.Test.QuickCheck.Arbitrary+-- Copyright : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee, Trevor L. McDonell+-- License : BSD3+--+-- Maintainer : Manuel M T Chakravarty <chak@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions)+--++module Data.Array.Accelerate.Test.QuickCheck.Arbitrary+ (+ -- Instances of Arbitrary+ arbitraryIntegralExp, arbitraryIntegralVector,+ arbitraryFloatingExp, arbitraryFloatingVector+ )+ where++import Data.Array.Accelerate+import Data.Array.Accelerate.Smart++import Control.Applicative hiding (Const)+import Control.Monad+import Data.List+import Foreign.Storable+import Test.QuickCheck+++-- Primitive Types+-- ---------------+instance Arbitrary Int8 where arbitrary = arbitrarySizedIntegral+instance Arbitrary Int16 where arbitrary = arbitrarySizedIntegral+instance Arbitrary Int32 where arbitrary = arbitrarySizedIntegral+instance Arbitrary Int64 where arbitrary = arbitrarySizedIntegral++instance Arbitrary Word where arbitrary = arbitrarySizedIntegral+instance Arbitrary Word8 where arbitrary = arbitrarySizedIntegral+instance Arbitrary Word16 where arbitrary = arbitrarySizedIntegral+instance Arbitrary Word32 where arbitrary = arbitrarySizedIntegral+instance Arbitrary Word64 where arbitrary = arbitrarySizedIntegral+++-- Arrays+-- ------+instance (Elem e, Arbitrary e) => Arbitrary (Array DIM0 e) where+ arbitrary = fromList () <$> vector 1++instance (Elem e, Arbitrary e) => Arbitrary (Array DIM1 e) where+ arbitrary = do+ n <- suchThat arbitrary (>= 0)+ fromList n <$> vector n++instance (Elem e, Arbitrary e) => Arbitrary (Array DIM2 e) where+ arbitrary = do+ (n, [a,b]) <- clamp <$> vectorOf 2 (suchThat arbitrary (> 0))+ fromList (a,b) <$> vector n++instance (Elem e, Arbitrary e) => Arbitrary (Array DIM3 e) where+ arbitrary = do+ (n, [a,b,c]) <- clamp <$> vectorOf 3 (suchThat arbitrary (> 0))+ fromList (a,b,c) <$> vector n++instance (Elem e, Arbitrary e) => Arbitrary (Array DIM4 e) where+ arbitrary = do+ (n, [a,b,c,d]) <- clamp <$> vectorOf 4 (suchThat arbitrary (> 0))+ fromList (a,b,c,d) <$> vector n++instance (Elem e, Arbitrary e) => Arbitrary (Array DIM5 e) where+ arbitrary = do+ (n, [a,b,c,d,e]) <- clamp <$> vectorOf 5 (suchThat arbitrary (> 0))+ fromList (a,b,c,d,e) <$> vector n++-- make sure not to create an index too large that we get integer overflow when+-- converting it to linear form+--+clamp :: [Int] -> (Int, [Int])+clamp = mapAccumL k 1+ where k a x = let n = Prelude.max 1 (x `mod` (maxBound `div` a))+ in (n*a, n)+++-- Expressions+-- -----------+instance Arbitrary (Acc (Vector Int)) where arbitrary = sized arbitraryIntegralVector+instance Arbitrary (Acc (Vector Int8)) where arbitrary = sized arbitraryIntegralVector+instance Arbitrary (Acc (Vector Int16)) where arbitrary = sized arbitraryIntegralVector+instance Arbitrary (Acc (Vector Int32)) where arbitrary = sized arbitraryIntegralVector+instance Arbitrary (Acc (Vector Int64)) where arbitrary = sized arbitraryIntegralVector++instance Arbitrary (Acc (Vector Word)) where arbitrary = sized arbitraryIntegralVector+instance Arbitrary (Acc (Vector Word8)) where arbitrary = sized arbitraryIntegralVector+instance Arbitrary (Acc (Vector Word16)) where arbitrary = sized arbitraryIntegralVector+instance Arbitrary (Acc (Vector Word32)) where arbitrary = sized arbitraryIntegralVector+instance Arbitrary (Acc (Vector Word64)) where arbitrary = sized arbitraryIntegralVector++instance Arbitrary (Acc (Vector Float)) where arbitrary = sized arbitraryFloatingVector+instance Arbitrary (Acc (Vector Double)) where arbitrary = sized arbitraryFloatingVector+++-- ugly ugly duplication...+--++arbitraryIntegralExp :: (IsIntegral e, Elem e, Arbitrary e) => Int -> Gen (Exp e)+arbitraryIntegralExp 0 = liftM Const arbitrary+arbitraryIntegralExp n =+ let m = n `div` 4+ in oneof [ do vec <- arbitraryIntegralVector m+ idx <- abs <$> arbitrary+ let idx' = shape vec ==* 0 ? (0, constant idx `mod` shape vec)+ return (IndexScalar vec idx')+ , liftM (!constant ()) (Fold <$> associative <*> arbitraryIntegralExp m <*> arbitraryIntegralVector m)+ ]++arbitraryIntegralVector :: (IsNum e, IsIntegral e, Elem e, Arbitrary e) => Int -> Gen (Acc (Vector e))+arbitraryIntegralVector 0 = Use <$> arbitrary+arbitraryIntegralVector n =+ let m = n `div` 4+ in oneof [ Map <$> unaryIntegral <*> arbitraryIntegralVector m+ , ZipWith <$> binaryIntegral <*> arbitraryIntegralVector m <*> arbitraryIntegralVector m+ , liftM FstArray (Scanl <$> associative <*> arbitraryIntegralExp m <*> arbitraryIntegralVector m)+ , liftM FstArray (Scanr <$> associative <*> arbitraryIntegralExp m <*> arbitraryIntegralVector m)+ ]+++arbitraryFloatingExp :: (IsFloating e, Elem e, Arbitrary e) => Int -> Gen (Exp e)+arbitraryFloatingExp 0 = liftM Const arbitrary+arbitraryFloatingExp n =+ let m = n `div` 4+ in oneof [ do vec <- arbitraryFloatingVector m+ idx <- abs <$> arbitrary+ let idx' = shape vec ==* 0 ? (0, constant idx `mod` shape vec)+ return (IndexScalar vec idx')+ , liftM (!constant ()) (Fold <$> associative <*> arbitraryFloatingExp m <*> arbitraryFloatingVector m)+ ]++arbitraryFloatingVector :: (IsNum e, IsFloating e, Elem e, Arbitrary e) => Int -> Gen (Acc (Vector e))+arbitraryFloatingVector 0 = Use <$> arbitrary+arbitraryFloatingVector n =+ let m = n `div` 4+ in oneof [ Map <$> unaryFloating <*> arbitraryFloatingVector m+ , ZipWith <$> binaryFloating <*> arbitraryFloatingVector m <*> arbitraryFloatingVector m+ , liftM FstArray (Scanl <$> associative <*> arbitraryFloatingExp m <*> arbitraryFloatingVector m)+ , liftM FstArray (Scanr <$> associative <*> arbitraryFloatingExp m <*> arbitraryFloatingVector m)+ ]+++-- Functions+-- ---------++wordSize :: Int+wordSize = 8 * sizeOf (undefined :: Int)+++unaryNum :: (Elem t, IsNum t, Arbitrary t) => Gen (Exp t -> Exp t)+unaryNum = oneof+ [ return mkNeg+ , return mkAbs+ , return mkSig+ ]++binaryNum :: (Elem t, IsNum t) => Gen (Exp t -> Exp t -> Exp t)+binaryNum = oneof+ [ return mkAdd+ , return mkSub+ , return mkMul+ ]++unaryIntegral :: (Elem t, IsIntegral t, Arbitrary t) => Gen (Exp t -> Exp t)+unaryIntegral = sized $ \n -> oneof+ [ return mkBNot+ , flip shift . constant <$> choose (-wordSize,wordSize)+ , flip rotate . constant <$> choose (-wordSize,wordSize)+ , flip mkBShiftL . constant <$> choose (0,wordSize)+ , flip mkBShiftR . constant <$> choose (0,wordSize)+ , flip mkBRotateL . constant <$> choose (0,wordSize)+ , flip mkBRotateR . constant <$> choose (0,wordSize)+ , flip setBit . constant <$> choose (0,wordSize)+ , flip clearBit . constant <$> choose (0,wordSize)+ , flip complementBit . constant <$> choose (0,wordSize)+ , unaryNum+ , binaryNum <*> arbitraryIntegralExp (n `div` 2)+ , binaryIntegral <*> arbitraryIntegralExp (n `div` 2)+ ]++binaryIntegral :: (Elem t, IsIntegral t) => Gen (Exp t -> Exp t -> Exp t)+binaryIntegral = oneof+ [ return mkQuot+ , return mkRem+ , return mkIDiv+ , return mkMod+ , return mkBAnd+ , return mkBOr+ , return mkBXor+ , return mkMax+ , return mkMin+ , binaryNum+ ]++unaryFloating :: (Elem t, IsFloating t, Arbitrary t) => Gen (Exp t -> Exp t)+unaryFloating = sized $ \n -> oneof+ [ return mkSin+ , return mkCos+ , return mkTan+-- , return mkAsin -- can't be sure inputs will be in the valid range+-- , return mkAcos+ , return mkAtan+ , return mkAsinh+-- , return mkAcosh+-- , return mkAtanh+-- , return mkExpFloating+-- , return mkSqrt+-- , return mkLog+ , unaryNum+ , binaryFloating `ap` arbitraryFloatingExp (n `div` 2)+ ]++binaryFloating :: (Elem t, IsFloating t) => Gen (Exp t -> Exp t -> Exp t)+binaryFloating = oneof+ [ return mkFPow+ , return mkLogBase+ , return mkMax+ , return mkMin+ , binaryNum+ ]++associative :: (Elem t, IsNum t) => Gen (Exp t -> Exp t -> Exp t)+associative = oneof+ [ return mkMax+ , return mkMin+ , return mkAdd+ , return mkMul+ ]+
Data/Array/Accelerate/Tuple.hs view
@@ -1,17 +1,17 @@ {-# LANGUAGE GADTs, TypeFamilies, FlexibleInstances #-} --- |Embedded array processing language: tuple representation------ Copyright (c) 2009 Manuel M T Chakravarty, Gabriele Keller, Sean Lee------ License: BSD3+-- Module : Data.Array.Accelerate.Tuple+-- Copyright : [2009..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee+-- License : BSD3 ------ Description ---------------------------------------------------------------+-- Maintainer : Manuel M T Chakravarty <chak@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions) ----- Our representation of tuples are heterogenous snoc lists, which are typed --- by type lists, where '()' and '(,)' are type-level nil and snoc,--- respectively. The components may only be drawn from types that can be--- used as array elements.+-- Our representation of tuples are heterogenous snoc lists, which are typed +-- by type lists, where '()' and '(,)' are type-level nil and snoc,+-- respectively. The components may only be drawn from types that can be+-- used as array elements. module Data.Array.Accelerate.Tuple ( @@ -67,3 +67,26 @@ type TupleRepr (a, b, c, d, e) = (TupleRepr (a, b, c, d), e) fromTuple (x, y, z, v, w) = ((((((), x), y), z), v), w) toTuple ((((((), x), y), z), v), w) = (x, y, z, v, w)++instance IsTuple (a, b, c, d, e, f) where+ type TupleRepr (a, b, c, d, e, f) = (TupleRepr (a, b, c, d, e), f)+ fromTuple (x, y, z, v, w, r) = (((((((), x), y), z), v), w), r)+ toTuple (((((((), x), y), z), v), w), r) = (x, y, z, v, w, r)++instance IsTuple (a, b, c, d, e, f, g) where+ type TupleRepr (a, b, c, d, e, f, g) = (TupleRepr (a, b, c, d, e, f), g)+ fromTuple (x, y, z, v, w, r, s) = ((((((((), x), y), z), v), w), r), s)+ toTuple ((((((((), x), y), z), v), w), r), s) = (x, y, z, v, w, r, s)++instance IsTuple (a, b, c, d, e, f, g, h) where+ type TupleRepr (a, b, c, d, e, f, g, h) = (TupleRepr (a, b, c, d, e, f, g), h)+ fromTuple (x, y, z, v, w, r, s, t) = (((((((((), x), y), z), v), w), r), s), t)+ toTuple (((((((((), x), y), z), v), w), r), s), t) = (x, y, z, v, w, r, s, t)++instance IsTuple (a, b, c, d, e, f, g, h, i) where+ type TupleRepr (a, b, c, d, e, f, g, h, i) = (TupleRepr (a, b, c, d, e, f, g, h), i)+ fromTuple (x, y, z, v, w, r, s, t, u) + = ((((((((((), x), y), z), v), w), r), s), t), u)+ toTuple ((((((((((), x), y), z), v), w), r), s), t), u) + = (x, y, z, v, w, r, s, t, u)+
Data/Array/Accelerate/Type.hs view
@@ -1,15 +1,15 @@ {-# LANGUAGE GADTs, TypeFamilies, FlexibleInstances #-} --- |Embedded array processing language: data types------ Copyright (c) [2008..2009] Manuel M T Chakravarty, Gabriele Keller, Sean Lee+-- Module : Data.Array.Accelerate.Type+-- Copyright : [2008..2010] Manuel M T Chakravarty, Gabriele Keller, Sean Lee+-- License : BSD3 ----- License: BSD3+-- Maintainer : Manuel M T Chakravarty <chak@cse.unsw.edu.au>+-- Stability : experimental+-- Portability : non-portable (GHC extensions) ------ Description ---------------------------------------------------------------+-- /Scalar types supported in array computations/ ----- Scalar types supported in array computations--- ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ -- Integral types: Int, Int8, Int16, Int32, Int64, Word, Word8, Word16, Word32, -- Word64, CShort, CUShort, CInt, CUInt, CLong, CULong, CLLong, CULLong --@@ -17,8 +17,8 @@ -- -- Non-numeric types: Bool, Char, CChar, CSChar, CUChar ----- `Int' has the same bitwidth as in plain Haskell computations, and `Float'--- and `Double' represent IEEE single and double precision floating point+-- 'Int' has the same bitwidth as in plain Haskell computations, and 'Float'+-- and 'Double' represent IEEE single and double precision floating point -- numbers, respectively. module Data.Array.Accelerate.Type (@@ -40,6 +40,42 @@ -- in the future, CHalf +-- Extend Typeable support for 8- and 9-tuple+-- ------------------------------------------++class Typeable8 t where+ typeOf8 :: t a b c d e f g h -> TypeRep++instance Typeable8 (,,,,,,,) where+ typeOf8 _ = mkTyCon "(,,,,,,,)" `mkTyConApp` []++typeOf7Default :: (Typeable8 t, Typeable a) => t a b c d e f g h -> TypeRep+typeOf7Default x = typeOf7 x `mkAppTy` typeOf (argType x)+ where+ argType :: t a b c d e f g h -> a+ argType = undefined++instance (Typeable8 s, Typeable a)+ => Typeable7 (s a) where+ typeOf7 = typeOf7Default+ +class Typeable9 t where+ typeOf9 :: t a b c d e f g h i -> TypeRep++instance Typeable9 (,,,,,,,,) where+ typeOf9 _ = mkTyCon "(,,,,,,,,)" `mkTyConApp` []++typeOf8Default :: (Typeable9 t, Typeable a) => t a b c d e f g h i -> TypeRep+typeOf8Default x = typeOf8 x `mkAppTy` typeOf (argType x)+ where+ argType :: t a b c d e f g h i -> a+ argType = undefined++instance (Typeable9 s, Typeable a)+ => Typeable8 (s a) where+ typeOf8 = typeOf8Default++ -- Scalar types -- ------------ @@ -544,6 +580,17 @@ SingleTuple :: ScalarType a -> TupleType a PairTuple :: TupleType a -> TupleType b -> TupleType (a, b) ++-- Stencil support+-- ---------------++-- |Boundary condition specification for stencil operations.+--+data Boundary a = Clamp -- ^clamp coordinates to the extent of the array+ | Mirror -- ^mirror coordinates beyond the array extent+ | Wrap -- ^wrap coordinates around on each dimension+ | Constant a -- ^use a constant value for outlying coordinates + deriving Show {- -- Vector GPU data types
LICENSE view
@@ -1,5 +1,5 @@-Copyright (c) [2007..2009] Manuel M T Chakravarty, Gabriele Keller & Sean Lee,-University of New South Wales. All rights reserved.+Copyright (c) [2007..2009] Manuel M T Chakravarty, Gabriele Keller, Sean Lee &+Trevor L. McDonell, University of New South Wales. All rights reserved. Redistribution and use in source and binary forms, with or without modification, are permitted provided that the following conditions are met:
accelerate.cabal view
@@ -1,7 +1,7 @@ Name: accelerate-Version: 0.7.1.0+Version: 0.8.0.0 Cabal-version: >= 1.6-Tested-with: GHC >= 6.12.1+Tested-with: GHC >= 6.12.3 Build-type: Simple Synopsis: An embedded language for accelerated array processing@@ -15,12 +15,15 @@ CUDA-capable NVIDIA GPUs. . To use the CUDA backend, you need to have CUDA version 3.x- installed. The CUDA backend still misses some features of- the full language; in particular, the array operations- 'replicate', 'slice', and 'foldSeg' are not yet supported.+ installed. The CUDA backend currently doesn't support 'Char'+ and 'Bool' arrays. .- Known bugs in this version: http://trac.haskell.org/accelerate/query?status=new&status=assigned&status=reopened&status=closed&version=0.7.1.0&order=priority+ Known bugs in this version:+ <http://trac.haskell.org/accelerate/query?status=new&status=assigned&status=reopened&status=closed&version=0.8.0.0&order=priority> .+ * New in 0.8.0.0: 'replicate', 'slice' and 'foldSeg' supported in the + CUDA backend; frontend and interpreter support for 'stencil'; bug fixes+ . * New in 0.7.1.0: the CUDA backend and a number of scalar functions License: BSD3 License-file: LICENSE@@ -42,21 +45,26 @@ accelerate_cuda_utils.h backpermute.inl fold.inl+ fold_segmented.inl map.inl permute.inl+ replicate.inl+ slice.inl zipWith.inl thrust/scan_safe.inl Extra-source-files: INSTALL- examples/simple/DotP.hs- examples/simple/Filter.hs- examples/simple/Main.hs+ include/accelerate.h+ examples/simple/src/DotP.hs+ examples/simple/src/Filter.hs+ examples/simple/src/Main.hs examples/simple/Makefile- examples/simple/SAXPY.hs- examples/simple/SMVM.hs- examples/simple/Square.hs- examples/simple/Sum.hs- examples/simple/Time.hs+ examples/simple/src/Random.hs+ examples/simple/src/SAXPY.hs+ examples/simple/src/SMVM.hs+ examples/simple/src/Stencil.hs+ examples/simple/src/Square.hs+ examples/simple/src/Sum.hs examples/rasterize/RasterizeAcc.hs examples/rasterize/rasterize-test1.txt examples/rasterize/rasterize-test2.txt@@ -65,19 +73,38 @@ examples/rasterize/rasterize.hs Flag llvm- Description: enable the LLVM backend (sequential)+ Description: Enable the LLVM backend (sequential) Default: False Flag cuda- Description: enable the CUDA parallel backend for NVIDIA GPUs+ Description: Enable the CUDA parallel backend for NVIDIA GPUs Default: True +Flag test-suite+ Description: Export extra test modules+ Default: False++Flag bounds-checks+ Description: Enable bounds checking+ Default: True++Flag unsafe-checks+ Description: Enable bounds checking in unsafe operations+ Default: False++Flag internal-checks+ Description: Enable internal consistency checks+ Default: False+ Library Build-depends: array, base == 4.*, ghc-prim, haskell98, pretty++ Include-Dirs: include+ If flag(llvm) Build-depends: llvm >= 0.6.8 @@ -87,13 +114,16 @@ containers, cuda >= 0.2 && < 0.3, directory,- fclabels,+ fclabels >= 0.9 && < 1.0, filepath, language-c >= 0.3 && < 0.4, monads-fd, transformers >= 0.2 && < 0.3, unix + if flag(test-suite)+ Build-depends: QuickCheck == 2.*+ Exposed-modules: Data.Array.Accelerate Data.Array.Accelerate.Interpreter -- If flag(llvm)@@ -101,8 +131,14 @@ If flag(cuda) Exposed-modules: Data.Array.Accelerate.CUDA- - Other-modules: Data.Array.Accelerate.Array.Data++ If flag(test-suite)+ Exposed-modules: Data.Array.Accelerate.Test+ Data.Array.Accelerate.Test.QuickCheck+ Data.Array.Accelerate.Test.QuickCheck.Arbitrary++ Other-modules: Data.Array.Accelerate.Internal.Check+ Data.Array.Accelerate.Array.Data Data.Array.Accelerate.Array.Delayed Data.Array.Accelerate.Array.Representation Data.Array.Accelerate.Array.Sugar@@ -131,11 +167,21 @@ Data.Array.Accelerate.CUDA.CodeGen Data.Array.Accelerate.CUDA.Compile Data.Array.Accelerate.CUDA.Execute+ Data.Array.Accelerate.CUDA.Smart Data.Array.Accelerate.CUDA.State + if flag(bounds-checks)+ cpp-options: -DACCELERATE_BOUNDS_CHECKS++ if flag(unsafe-checks)+ cpp-options: -DACCELERATE_UNSAFE_CHECKS++ if flag(internal-checks)+ cpp-options: -DACCELERATE_INTERNAL_CHECKS+ Ghc-options: -O2 -Wall -fno-warn-orphans -fno-warn-name-shadowing- Extensions: FlexibleContexts, FlexibleInstances, - ExistentialQuantification, GADTs, TypeFamilies, + Extensions: FlexibleContexts, FlexibleInstances, TypeSynonymInstances,+ ExistentialQuantification, GADTs, TypeFamilies, MultiParamTypeClasses, ScopedTypeVariables, DeriveDataTypeable, BangPatterns, PatternGuards, TypeOperators, RankNTypes
cubits/accelerate_cuda_extras.h view
@@ -135,72 +135,209 @@ * Shapes * -------------------------------------------------------------------------- */ -typedef int Ix;-typedef Ix DIM1;-typedef struct { Ix a,b; } DIM2;-typedef struct { Ix a,b,c; } DIM3;-typedef struct { Ix a,b,c,d; } DIM4;-typedef struct { Ix a,b,c,d,e; } DIM5;+typedef int32_t Ix;+typedef Ix DIM0;+typedef Ix DIM1;+typedef struct { Ix a1,a0; } DIM2;+typedef struct { Ix a2,a1,a0; } DIM3;+typedef struct { Ix a3,a2,a1,a0; } DIM4;+typedef struct { Ix a4,a3,a2,a1,a0; } DIM5; #ifdef __cplusplus +/*+ * Number of dimensions of a shape+ */ static __inline__ __device__ int dim(DIM1 sh) { return 1; } static __inline__ __device__ int dim(DIM2 sh) { return 2; } static __inline__ __device__ int dim(DIM3 sh) { return 3; } static __inline__ __device__ int dim(DIM4 sh) { return 4; } static __inline__ __device__ int dim(DIM5 sh) { return 5; } +/*+ * Yield the total number of elements in a shape+ */ static __inline__ __device__ int size(DIM1 sh) { return sh; }-static __inline__ __device__ int size(DIM2 sh) { return sh.a * sh.b; }-static __inline__ __device__ int size(DIM3 sh) { return sh.a * sh.b * sh.c; }-static __inline__ __device__ int size(DIM4 sh) { return sh.a * sh.b * sh.c * sh.d; }-static __inline__ __device__ int size(DIM5 sh) { return sh.a * sh.b * sh.c * sh.d * sh.e; }+static __inline__ __device__ int size(DIM2 sh) { return sh.a0 * sh.a1; }+static __inline__ __device__ int size(DIM3 sh) { return sh.a0 * sh.a1 * sh.a2; }+static __inline__ __device__ int size(DIM4 sh) { return sh.a0 * sh.a1 * sh.a2 * sh.a3; }+static __inline__ __device__ int size(DIM5 sh) { return sh.a0 * sh.a1 * sh.a2 * sh.a3 * sh.a4; } -static __inline__ __device__ int index(DIM1 sh, DIM1 ix)+/*+ * Convert the individual dimensions of a linear array into a shape+ */+static __inline__ __device__ DIM1 shape(Ix a) {+ return a;+}++static __inline__ __device__ DIM2 shape(Ix b, Ix a)+{+ DIM2 sh = { b, a };+ return sh;+}++static __inline__ __device__ DIM3 shape(Ix c, Ix b, Ix a)+{+ DIM3 sh = { c, b, a };+ return sh;+}++static __inline__ __device__ DIM4 shape(Ix d, Ix c, Ix b, Ix a)+{+ DIM4 sh = { d, c, b, a };+ return sh;+}++static __inline__ __device__ DIM5 shape(Ix e, Ix d, Ix c, Ix b, Ix a)+{+ DIM5 sh = { e, d, c, b, a };+ return sh;+}++/*+ * Yield the index position in a linear, row-major representation of the array.+ * First argument is the shape of the array, the second the index+ */+static __inline__ __device__ Ix toIndex(DIM1 sh, DIM1 ix)+{ return ix; } -static __inline__ __device__ int index(DIM2 sh, DIM2 ix)+static __inline__ __device__ Ix toIndex(DIM2 sh, DIM2 ix) {- DIM1 sh_ = sh.a;- DIM1 ix_ = ix.a;- return index(sh_, ix_) + ix.b * size(sh_);+ DIM1 sh_ = sh.a1;+ DIM1 ix_ = ix.a1;+ return toIndex(sh_, ix_) * sh.a0 + ix.a0; } -static __inline__ __device__ int index(DIM3 sh, DIM3 ix)+static __inline__ __device__ Ix toIndex(DIM3 sh, DIM3 ix) {- DIM2 sh_ = { sh.a, sh.b };- DIM2 ix_ = { ix.a, ix.b };- return index(sh_, ix_) + ix.c * size(sh_);+ DIM2 sh_ = { sh.a2, sh.a1 };+ DIM2 ix_ = { ix.a2, ix.a1 };+ return toIndex(sh_, ix_) * sh.a0 + ix.a0; } -static __inline__ __device__ int index(DIM4 sh, DIM4 ix)+static __inline__ __device__ Ix toIndex(DIM4 sh, DIM4 ix) {- DIM3 sh_ = { sh.a, sh.b, sh.c };- DIM3 ix_ = { ix.a, ix.b, ix.c };- return index(sh_, ix_) + ix.d * size(sh_);+ DIM3 sh_ = { sh.a3, sh.a2, sh.a1 };+ DIM3 ix_ = { ix.a3, ix.a2, ix.a1 };+ return toIndex(sh_, ix_) * sh.a0 + ix.a0; } -static __inline__ __device__ int index(DIM5 sh, DIM5 ix)+static __inline__ __device__ Ix toIndex(DIM5 sh, DIM5 ix) {- DIM4 sh_ = { sh.a, sh.b, sh.c, sh.d };- DIM4 ix_ = { ix.a, ix.b, ix.c, ix.d };- return index(sh_, ix_) + ix.e * size(sh_);+ DIM4 sh_ = { sh.a4, sh.a3, sh.a2, sh.a1 };+ DIM4 ix_ = { ix.a4, ix.a3, ix.a2, ix.a1 };+ return toIndex(sh_, ix_) * sh.a0 + ix.a0; } +/*+ * Inverse of 'toIndex'+ */+static __inline__ __device__ DIM1 fromIndex(DIM1 sh, Ix ix)+{+ return ix;+}++static __inline__ __device__ DIM2 fromIndex(DIM2 sh, Ix ix)+{+ DIM1 sh_ = shape(sh.a1);+ DIM1 ix_ = fromIndex(sh_, ix / sh.a0);+ return shape(ix_, ix % sh.a0);+}++static __inline__ __device__ DIM3 fromIndex(DIM3 sh, Ix ix)+{+ DIM2 sh_ = shape(sh.a2, sh.a1);+ DIM2 ix_ = fromIndex(sh_, ix / sh.a0);+ return shape(ix_.a1, ix_.a0, ix % sh.a0);+}++static __inline__ __device__ DIM4 fromIndex(DIM4 sh, Ix ix)+{+ DIM3 sh_ = shape(sh.a3, sh.a2, sh.a1);+ DIM3 ix_ = fromIndex(sh_, ix / sh.a0);+ return shape(ix_.a2, ix_.a1, ix_.a0, ix % sh.a0);+}++static __inline__ __device__ DIM5 fromIndex(DIM5 sh, Ix ix)+{+ DIM4 sh_ = shape(sh.a4, sh.a3, sh.a2, sh.a1);+ DIM4 ix_ = fromIndex(sh_, ix / sh.a0);+ return shape(ix_.a3, ix_.a2, ix_.a1, ix_.a0, ix % sh.a0);+}+++/*+ * Test for the magic index `ignore`+ */+static __inline__ __device__ int ignore(DIM1 ix)+{+ return ix == -1;+}++static __inline__ __device__ int ignore(DIM2 ix)+{+ return ix.a0 == -1 && ix.a1 == -1;+}++static __inline__ __device__ int ignore(DIM3 ix)+{+ return ix.a0 == -1 && ix.a1 == -1 && ix.a2 == -1;+}++static __inline__ __device__ int ignore(DIM4 ix)+{+ return ix.a0 == -1 && ix.a1 == -1 && ix.a2 == -1 && ix.a3 == -1;+}++static __inline__ __device__ int ignore(DIM5 ix)+{+ return ix.a0 == -1 && ix.a1 == -1 && ix.a2 == -1 && ix.a3 == -1 && ix.a4 == -1;+}++ #else int dim(Ix sh); int size(Ix sh);-int index(Ix sh, Ix ix);+int shape(Ix sh);+int toIndex(Ix sh, Ix ix);+int fromIndex(Ix sh, Ix ix); #endif - /* ------------------------------------------------------------------------------ * Tuple Types+ * Functions * -------------------------------------------------------------------------- */++#ifdef __cplusplus+template <typename T>+static __inline__ __device__ T rotateL(const T x, int32_t i)+{+ return (i &= 8 * sizeof(x) - 1) == 0 ? x : x << i | x >> 8 * sizeof(x) - i;+}++template <typename T>+static __inline__ __device__ T rotateR(const T x, int32_t i)+{+ return (i &= 8 * sizeof(x) - 1) == 0 ? x : x >> i | x << 8 * sizeof(x) - i;+}++template <typename T>+static __inline__ __device__ T idiv(const T x, const T y)+{+ return x > 0 && y < 0 ? (x - y - 1) / y : (x < 0 && y > 0 ? (x - y + 1) / y : x / y);+}++template <typename T>+static __inline__ __device__ T mod(const T x, const T y)+{+ const T r = x % y;+ return x > 0 && y < 0 || x < 0 && y > 0 ? (r != 0 ? r + y : 0) : r;+}+#endif #endif // __ACCELERATE_CUDA_EXTRAS_H__
cubits/backpermute.inl view
@@ -18,15 +18,19 @@ ( ArrOut d_out, const ArrIn0 d_in0,- const Ix shape+ const DimOut shOut,+ const DimIn0 shIn0 ) {- Ix idx;+ Ix shapeSize = size(shOut); const Ix gridSize = __umul24(blockDim.x, gridDim.x); - for (idx = __umul24(blockDim.x, blockIdx.x) + threadIdx.x; idx < shape; idx += gridSize)+ for (Ix ix = __umul24(blockDim.x, blockIdx.x) + threadIdx.x; ix < shapeSize; ix += gridSize) {- set(d_out, idx, get0(d_in0, project(idx)));+ DimOut dst = fromIndex(shOut, ix);+ DimIn0 src = project(dst);++ set(d_out, ix, get0(d_in0, toIndex(shIn0, src))); } }
cubits/fold.inl view
@@ -8,10 +8,12 @@ * * ---------------------------------------------------------------------------*/ -#include <accelerate_cuda_utils.h>-+/*+ * We require block sizes are a power of two. This value gives maximum occupancy+ * for both 1.x and 2.x class devices.+ */ #ifndef BLOCK_SIZE-#define BLOCK_SIZE blockDim.x+#define BLOCK_SIZE 256 #endif #ifndef LENGTH_IS_POW_2@@ -33,17 +35,16 @@ const Ix shape ) {- extern __shared__ TyOut scratch[];+ extern volatile __shared__ TyOut s_data[]; /* * Calculate first level of reduction reading into shared memory */ Ix i;+ TyOut sum = identity(); const Ix tid = threadIdx.x; const Ix gridSize = BLOCK_SIZE * 2 * gridDim.x; - scratch[tid] = identity();- /* * Reduce multiple elements per thread. The number is determined by the * number of active thread blocks (via gridDim). More blocks will result in@@ -53,40 +54,45 @@ */ for (i = blockIdx.x * BLOCK_SIZE * 2 + tid; i < shape; i += gridSize) {- scratch[tid] = apply(scratch[tid], get0(d_in0, i));+ sum = apply(sum, get0(d_in0, i)); /* * Ensure we don't read out of bounds. This is optimised away if the * input length is a power of two */ if (LENGTH_IS_POW_2 || i + BLOCK_SIZE < shape)- scratch[tid] = apply(scratch[tid], get0(d_in0, i+BLOCK_SIZE));+ sum = apply(sum, get0(d_in0, i+BLOCK_SIZE)); }- __syncthreads(); /*- * Now, calculate the reduction in shared memory+ * Each thread puts its local sum into shared memory, then threads+ * cooperatively reduce the shared array to a single value. */- if (BLOCK_SIZE >= 512) { if (tid < 256) { scratch[tid] = apply(scratch[tid], scratch[tid+256]); } __syncthreads(); }- if (BLOCK_SIZE >= 256) { if (tid < 128) { scratch[tid] = apply(scratch[tid], scratch[tid+128]); } __syncthreads(); }- if (BLOCK_SIZE >= 128) { if (tid < 64) { scratch[tid] = apply(scratch[tid], scratch[tid+ 64]); } __syncthreads(); }+ s_data[tid] = sum;+ __syncthreads(); -#ifndef __DEVICE_EMULATION__+ if (BLOCK_SIZE >= 512) { if (tid < 256) { s_data[tid] = sum = apply(sum, s_data[tid+256]); } __syncthreads(); }+ if (BLOCK_SIZE >= 256) { if (tid < 128) { s_data[tid] = sum = apply(sum, s_data[tid+128]); } __syncthreads(); }+ if (BLOCK_SIZE >= 128) { if (tid < 64) { s_data[tid] = sum = apply(sum, s_data[tid+ 64]); } __syncthreads(); }+ if (tid < 32)-#endif {- if (BLOCK_SIZE >= 64) { scratch[tid] = apply(scratch[tid], scratch[tid+32]); __EMUSYNC; }- if (BLOCK_SIZE >= 32) { scratch[tid] = apply(scratch[tid], scratch[tid+16]); __EMUSYNC; }- if (BLOCK_SIZE >= 16) { scratch[tid] = apply(scratch[tid], scratch[tid+ 8]); __EMUSYNC; }- if (BLOCK_SIZE >= 8) { scratch[tid] = apply(scratch[tid], scratch[tid+ 4]); __EMUSYNC; }- if (BLOCK_SIZE >= 4) { scratch[tid] = apply(scratch[tid], scratch[tid+ 2]); __EMUSYNC; }- if (BLOCK_SIZE >= 2) { scratch[tid] = apply(scratch[tid], scratch[tid+ 1]); __EMUSYNC; }+ /*+ * Use an extra warps worth of elements of shared memory, to let threads+ * index beyond the input data without using any branch instructions.+ */+ s_data[tid] = sum = apply(sum, s_data[tid + 32]);+ s_data[tid] = sum = apply(sum, s_data[tid + 16]);+ s_data[tid] = sum = apply(sum, s_data[tid + 8]);+ s_data[tid] = sum = apply(sum, s_data[tid + 4]);+ s_data[tid] = sum = apply(sum, s_data[tid + 2]);+ s_data[tid] = sum = apply(sum, s_data[tid + 1]); } /* * Write the results of this block back to global memory */ if (tid == 0)- set(d_out, blockIdx.x, scratch[0]);+ set(d_out, blockIdx.x, sum); }
+ cubits/fold_segmented.inl view
@@ -0,0 +1,120 @@+/* -----------------------------------------------------------------------------+ *+ * Module : FoldSeg+ * Copyright : (c) [2009..2010] Trevor L. McDonell, Rami G. Mukhtar+ * License : BSD+ *+ * Reduction an array to a single value for each segment, using a binary+ * associative function+ *+ * ---------------------------------------------------------------------------*/++#define WARP_SIZE 32+++/*+ * Cooperatively reduce an array to a single value. The computation requires an+ * extra half-warp worth of elements of shared memory per block, to let threads+ * index beyond the input data without using any branch instructions.+ */+template <typename T>+static __inline__ __device__+T reduce_warp(volatile T* s_data, T sum)+{+ s_data[threadIdx.x] = sum;+ s_data[threadIdx.x] = sum = apply(sum, s_data[threadIdx.x + 16]);+ s_data[threadIdx.x] = sum = apply(sum, s_data[threadIdx.x + 8]);+ s_data[threadIdx.x] = sum = apply(sum, s_data[threadIdx.x + 4]);+ s_data[threadIdx.x] = sum = apply(sum, s_data[threadIdx.x + 2]);+ s_data[threadIdx.x] = sum = apply(sum, s_data[threadIdx.x + 1]);++ return sum;+}+++/*+ * Each segment of the vector is assigned to a warp, which computes the+ * reduction of the i-th section, in parallel.+ *+ * This division of work implies that the data arrays are accessed in a+ * contiguous manner (if not necessarily aligned). For devices of compute+ * capability 1.2 and later, these accesses will be coalesced. A single+ * transaction will be issued if all of the addresses for a half-warp happen to+ * fall within a single 128-byte boundary. Extra transactions will be made to+ * cover any spill. The same applies for 2.x devices, except that all widths are+ * doubled since transactions occur on a per-warp basis.+ *+ * Since an entire 32-thread warp is assigned for each segment, many threads+ * will remain idle when the segments are very small. This code relies on+ * implicit synchronisation among threads in a warp.+ *+ * The offset array contains the starting index for each segment in the input+ * array. The i-th warp reduces values in the input array at indices+ * [d_offset[i], d_offset[i+1]).+ */+extern "C"+__global__ void+fold_segmented+(+ ArrOut d_out,+ const ArrIn0 d_in0,+ const Int* d_offset,+ const Ix num_segments,+ const Ix length // of input array d_in0+)+{+ const Ix vectors_per_block = blockDim.x / WARP_SIZE;+ const Ix num_vectors = vectors_per_block * gridDim.x;+ const Ix thread_id = blockDim.x * blockIdx.x + threadIdx.x;+ const Ix vector_id = thread_id / WARP_SIZE;+ const Ix thread_lane = threadIdx.x & (WARP_SIZE - 1);+ const Ix vector_lane = threadIdx.x / WARP_SIZE;++ /*+ * Manually partition (dynamically-allocated) shared memory+ */+ extern __shared__ Ix s_ptrs[][2];+ TyOut* s_data = (TyOut*) &s_ptrs[vectors_per_block][2];++ for (Ix seg = vector_id; seg < num_segments; seg += num_vectors)+ {+ /*+ * Use two threads to fetch the indices of the start and end of this+ * segment. This results in single coalesced global read, instead of two+ * separate transactions. If we are the last segment, the final index is+ * taken from the overall array length.+ */+ if (seg < num_segments - 1)+ {+ if (thread_lane < 2)+ s_ptrs[vector_lane][thread_lane] = d_offset[seg + thread_lane];+ }+ else+ {+ if (thread_lane == 0)+ s_ptrs[vector_lane][0] = d_offset[seg];++ s_ptrs[vector_lane][1] = length;+ }++ const Ix start = s_ptrs[vector_lane][0];+ const Ix end = s_ptrs[vector_lane][1];++ /*+ * Have each thread read in all values for this segment, accumulating a+ * local sum. This is then reduced cooperatively in shared memory.+ */+ TyOut sum = identity();+ for (Ix i = start + thread_lane; i < end; i += WARP_SIZE)+ sum = apply(sum, get0(d_in0, i));++ sum = reduce_warp(s_data, sum);++ /*+ * Finally, the first thread writes the result for this segment+ */+ if (thread_lane == 0)+ set(d_out, seg, sum);+ }+}+
cubits/permute.inl view
@@ -20,20 +20,22 @@ ( ArrOut d_out, const ArrIn0 d_in0,- const Ix shape+ const DimOut shOut,+ const DimIn0 shIn0 ) {- Ix dst;- Ix idx;+ Ix shapeSize = size(shIn0); const Ix gridSize = __umul24(blockDim.x, gridDim.x); - for (idx = __umul24(blockDim.x, blockIdx.x) + threadIdx.x; idx < shape; idx += gridSize)+ for (Ix ix = __umul24(blockDim.x, blockIdx.x) + threadIdx.x; ix < shapeSize; ix += gridSize) {- dst = project(idx);+ DimIn0 src = fromIndex(shIn0, ix);+ DimOut dst = project(src); - if (dst != ignore)+ if (!ignore(dst)) {- set(d_out, dst, apply(get0(d_in0, idx), get0(d_out, dst)));+ Ix j = toIndex(shOut, dst);+ set(d_out, j, apply(get0(d_in0, ix), get0(d_out, j))); } } }
+ cubits/replicate.inl view
@@ -0,0 +1,30 @@+/* -----------------------------------------------------------------------------+ *+ * Module : Replicate+ * Copyright : (c) [2009..2010] Trevor L. McDonell+ * License : BSD+ *+ * ---------------------------------------------------------------------------*/++extern "C"+__global__ void+replicate+(+ ArrOut d_out,+ const ArrIn0 d_in0,+ const Slice slice,+ const SliceDim sliceDim+)+{+ Ix shapeSize = size(sliceDim);+ const Ix gridSize = __umul24(blockDim.x, gridDim.x);++ for (Ix ix = __umul24(blockDim.x, blockIdx.x) + threadIdx.x; ix < shapeSize; ix += gridSize)+ {+ SliceDim dst = fromIndex(sliceDim, ix);+ Slice src = sliceIndex(dst);++ set(d_out, ix, get0(d_in0, toIndex(slice, src)));+ }+}+
+ cubits/slice.inl view
@@ -0,0 +1,31 @@+/* -----------------------------------------------------------------------------+ *+ * Module : Slice+ * Copyright : (c) [2009..2010] Trevor L. McDonell+ * License : BSD+ *+ * ---------------------------------------------------------------------------*/++extern "C"+__global__ void+slice+(+ ArrOut d_out,+ const ArrIn0 d_in0,+ const Slice slice,+ const CoSlice slix,+ const SliceDim sliceDim+)+{+ Ix shapeSize = size(slice);+ const Ix gridSize = __umul24(blockDim.x, gridDim.x);++ for (Ix ix = __umul24(blockDim.x, blockIdx.x) + threadIdx.x; ix < shapeSize; ix += gridSize)+ {+ Slice dst = fromIndex(slice, ix);+ SliceDim src = sliceIndex(dst, slix);++ set(d_out, ix, get0(d_in0, toIndex(sliceDim, src)));+ }+}+
cubits/thrust/scan_safe.inl view
@@ -42,6 +42,7 @@ else return threadIdx.x > 0 ? array[threadIdx.x - 1] : identity(); } +#if 0 template <bool inclusive, typename T> __device__ T scan_block_n(T* array, const unsigned int n, T val)@@ -64,6 +65,7 @@ if (inclusive) return val; else return threadIdx.x > 0 ? array[threadIdx.x - 1] : identity(); }+#endif template <bool inclusive>@@ -81,12 +83,11 @@ const Ix interval_begin = interval_size * blockIdx.x; const Ix interval_end = min(interval_begin + interval_size, N); - TyOut val = identity();- Ix base = interval_begin;+ TyOut val; Ix output = reverse ? interval_end - threadIdx.x - 1 : interval_begin + threadIdx.x; - // process full blocks- for(; base + blockDim.x <= interval_end; base += blockDim.x)+ // process intervals+ for(Ix base = interval_begin + threadIdx.x; base < interval_end; base += blockDim.x) { // read data val = get0(d_in0, output);@@ -109,35 +110,9 @@ if (reverse) output -= blockDim.x; else output += blockDim.x; }-- // process partially full block at end of input (if necessary)- if (base < interval_end)- {- // read data- if (base + threadIdx.x < interval_end)- {- val = get0(d_in0, output);- }-- // carry in- if (threadIdx.x == 0 && base != interval_begin)- {- TyOut tmp = sdata[blockDim.x - 1];- val = apply(tmp, val);- }- __syncthreads();-- // scan block- val = scan_block_n<inclusive>(sdata, interval_end - base, val);-- // write data- if (base + threadIdx.x < interval_end)- {- set(d_out, output, val);- }- } __syncthreads(); + // write block sum if (threadIdx.x == 0) { if (reverse)
cubits/zipWith.inl view
@@ -16,15 +16,20 @@ ArrOut d_out, const ArrIn1 d_in1, const ArrIn0 d_in0,- const Ix shape+ const DimOut shOut,+ const DimIn1 shIn1,+ const DimIn0 shIn0 ) {- Ix idx;+ Ix shapeSize = size(shOut); const Ix gridSize = __umul24(blockDim.x, gridDim.x); - for (idx = __umul24(blockDim.x, blockIdx.x) + threadIdx.x; idx < shape; idx += gridSize)+ for (Ix ix = __umul24(blockDim.x, blockIdx.x) + threadIdx.x; ix < shapeSize; ix += gridSize) {- set(d_out, idx, apply(get1(d_in1, idx), get0(d_in0, idx)));+ Ix i1 = toIndex(shIn1, fromIndex(shOut, ix));+ Ix i0 = toIndex(shIn0, fromIndex(shOut, ix));++ set(d_out, ix, apply(get1(d_in1, i1), get0(d_in0, i0))); } }
− examples/simple/DotP.hs
@@ -1,25 +0,0 @@-{-# LANGUAGE ParallelListComp #-}--module DotP (dotp, dotp_ref) where--import Prelude hiding (replicate, zip, map, filter, max, min, not, zipWith)-import qualified Prelude--import Data.Array.Unboxed-import Data.Array.IArray--import Data.Array.Accelerate--dotp :: Vector Float -> Vector Float -> Acc (Scalar Float)-dotp xs ys - = let- xs' = use xs- ys' = use ys- in- fold (+) 0 (zipWith (*) xs' ys')--dotp_ref :: UArray Int Float - -> UArray Int Float - -> UArray () Float-dotp_ref xs ys - = listArray ((), ()) $ [sum [x * y | x <- elems xs | y <- elems ys]]
− examples/simple/Filter.hs
@@ -1,39 +0,0 @@-{-# LANGUAGE FlexibleContexts #-}--module Filter (filter, filter_ref) where--import Prelude hiding (replicate, scanl, zip, map, filter, max, min, not, zipWith)-import qualified Prelude--import Data.Array.Unboxed (UArray)-import Data.Array.IArray (IArray)-import qualified Data.Array.IArray as IArray--import Data.Array.Accelerate---filter :: Elem a - => (Exp a -> Exp Bool) - -> Acc (Vector a)- -> Acc (Vector a)-filter p arr- = let flags = map (boolToInt . p) arr- (targetIdx, length) = scanl (+) 0 flags- arr' = backpermute (length!(constant ())) id arr- in- permute const arr' (\ix -> flags!ix ==* 0 ? (ignore, targetIdx!ix)) arr- -- FIXME: This is abusing 'permute' in that the first two arguments are- -- only justified because we know the permutation function will- -- write to each location in the target exactly once.- -- Instead, we should have a primitive that directly encodes the- -- compaction pattern of the permutation function.-- -filter_ref :: IArray UArray e - => (e -> Bool)- -> UArray Int e - -> UArray Int e-filter_ref p xs- = let xs' = Prelude.filter p (IArray.elems xs)- in- IArray.listArray (0, length xs' - 1) xs'
− examples/simple/Main.hs
@@ -1,372 +0,0 @@-{-# LANGUAGE FlexibleContexts, ParallelListComp #-}--module Main where--import Prelude hiding (filter)--import Control.Exception-import Data.Array.Unboxed-import Data.Array.IArray-import System.Random--import qualified Data.Array.Accelerate as Acc-import qualified Data.Array.Accelerate.Interpreter as Interp-import qualified Data.Array.Accelerate.CUDA as CUDA--import Time-import SAXPY-import Square-import DotP-import Filter-import Sum----- Auxilliary array functions--- ------------------------------ To ensure that a singleton unboxed array is fully evaluated--- -evaluateUScalar :: (IArray UArray e) => UArray () e -> IO ()-evaluateUScalar uarr = evaluate (uarr!()) >> return ()---- To ensure that a singleton unboxed array is fully evaluated--- -evaluateScalar :: Acc.Scalar e -> IO ()-evaluateScalar arr = evaluate (arr `Acc.indexArray` ()) >> return ()---- To ensure that an unboxed array is fully evaluated, just force one element--- -evaluateUVector :: (IArray UArray e) => UArray Int e -> IO ()-evaluateUVector uarr = evaluate (uarr!0) >> return ()---- To ensure that an unboxed array is fully evaluated, just force one element--- -evaluateVector :: Acc.Vector e -> IO ()-evaluateVector arr = evaluate (arr `Acc.indexArray` 0) >> return ()--randomUVector :: (Num e, Random e, IArray UArray e) => Int -> IO (UArray Int e)-randomUVector n- = do- rg <- newStdGen- let -- The std random function is too slow to generate really big vectors- -- with. Instead, we generate a short random vector and repeat that.- randvec = take k (randomRs (-100, 100) rg)- vec = listArray (0, n - 1) - [randvec !! (i `mod` k) | i <- [0..n - 1]]- evaluateUVector vec- return vec- where- k = 1000--convertUScalar :: (IArray UArray e, Acc.Elem e) - => UArray () e -> IO (Acc.Scalar e)-convertUScalar uarr- = do- let arr = Acc.fromIArray uarr- evaluateScalar arr- return arr--convertUVector :: (IArray UArray e, Acc.Elem e) - => UArray Int e -> IO (Acc.Vector e)-convertUVector uarr- = do- let arr = Acc.fromIArray uarr- evaluateVector arr- return arr--validate :: (Eq e, IArray UArray e, Ix ix) - => UArray ix e -> UArray ix e -> IO ()-validate arr_ref arr | arr_ref == arr = putStrLn "Valid."- | otherwise = putStrLn "INVALID!"--validateFloats :: Ix ix- => UArray ix Float -> UArray ix Float -> IO ()-validateFloats arr_ref arr | arr_ref `similar` arr = putStrLn "Valid."- | otherwise = putStrLn "INVALID!" >> diff arr_ref arr- where- similar arr1 arr2 = all (< epsilon) [abs ((x - y) / (x + y + epsilon)) - | x <- elems arr1 - | y <- elems arr2]- diff arr1 arr2 = sequence_ [if abs ((x - y) / x) < epsilon- then return ()- else putStrLn $ ">>> " ++ "(" ++ show x ++ ", " ++ show y ++ ")"- | x <- elems arr1 | y <- elems arr2]- epsilon = 0.0001----- Timing--- --------timeUScalar :: IArray UArray e => (() -> UArray () e) -> IO (UArray () e)-{-# NOINLINE timeUScalar #-}-timeUScalar testee - = do- (r, time1) <- oneRun testee- (r, time2) <- oneRun testee- (r, time3) <- oneRun testee- putStrLn $ showMinAvgMax milliseconds [time1, time2, time3] ++- " (wall - cpu min/avg/max in ms)"- return r- where- oneRun testee = do- start <- getTime- let r = testee ()- evaluateUScalar r- end <- getTime- return (r, end `minus` start)--timeScalar :: (IArray UArray e, Acc.Elem e)- => (() -> Acc.Scalar e) -> IO (UArray () e)-{-# NOINLINE timeScalar #-}-timeScalar testee - = do- (r, time1) <- oneRun testee- (r, time2) <- oneRun testee- (r, time3) <- oneRun testee- putStrLn $ showMinAvgMax milliseconds [time1, time2, time3] ++- " (wall - cpu min/avg/max in ms)"- return $ Acc.toIArray r- where- oneRun testee = do- start <- getTime- let r = testee ()- evaluateScalar r- end <- getTime- return (r, end `minus` start)--timeScalar' :: (IArray UArray e, Acc.Elem e)- => (() -> IO (Acc.Scalar e)) -> IO (UArray () e)-{-# NOINLINE timeScalar' #-}-timeScalar' testee - = do- (r, time1) <- oneRun testee- (r, time2) <- oneRun testee- (r, time3) <- oneRun testee- putStrLn $ showMinAvgMax milliseconds [time1, time2, time3] ++- " (wall - cpu min/avg/max in ms)"- return $ Acc.toIArray r- where- oneRun testee = do- start <- getTime- r <- testee ()- evaluateScalar r- end <- getTime- return (r, end `minus` start)--timeUVector :: IArray UArray e => (() -> UArray Int e) -> IO (UArray Int e)-{-# NOINLINE timeUVector #-}-timeUVector testee - = do- (r, time1) <- oneRun testee- (r, time2) <- oneRun testee- (r, time3) <- oneRun testee- putStrLn $ showMinAvgMax milliseconds [time1, time2, time3] ++- " (wall - cpu min/avg/max in ms)"- return r--- where-{-# NOINLINE oneRun #-}-oneRun testee = do- start <- getTime- let r = testee ()- evaluateUVector r- end <- getTime- return (r, end `minus` start)--timeVector :: (IArray UArray e, Acc.Elem e)- => (() -> Acc.Vector e) -> IO (UArray Int e)-{-# NOINLINE timeVector #-}-timeVector testee - = do- (r, time1) <- oneRun testee- (r, time2) <- oneRun testee- (r, time3) <- oneRun testee- putStrLn $ showMinAvgMax milliseconds [time1, time2, time3] ++- " (wall - cpu min/avg/max in ms)"- return $ Acc.toIArray r- where- oneRun testee = do- start <- getTime- let r = testee ()- evaluateVector r- end <- getTime- return (r, end `minus` start)---timeVector' :: (IArray UArray e, Acc.Elem e)- => (() -> IO (Acc.Vector e)) -> IO (UArray Int e)-{-# NOINLINE timeVector' #-}-timeVector' testee - = do- (r, time1) <- oneRun testee- (r, time2) <- oneRun testee- (r, time3) <- oneRun testee- putStrLn $ showMinAvgMax milliseconds [time1, time2, time3] ++- " (wall - cpu min/avg/max in ms)"- return $ Acc.toIArray r- where- oneRun testee = do- start <- getTime- r <- testee ()- evaluateVector r- end <- getTime- return (r, end `minus` start)---- Tests--- -------test_saxpy :: Int -> IO ()-test_saxpy n- = do- putStrLn "== SAXPY"- putStrLn $ "Generating data (n = " ++ show n ++ ")..."- v1_ref <- randomUVector n- v1 <- convertUVector v1_ref- v2_ref <- randomUVector n- v2 <- convertUVector v2_ref- putStrLn "Running reference code..."- ref_result <- timeUVector $ saxpy_ref' 1.5 v1_ref v2_ref- putStrLn "[Interpreter]"- putStrLn "Running Accelerate code..."- result <- timeVector $ saxpy_interp 1.5 v1 v2- putStrLn "Validating result..."- validateFloats ref_result result- putStrLn "[CUDA]"- putStrLn "Running Accelerate code..."- result_cuda <- timeVector' $ saxpy_cuda 1.5 v1 v2- putStrLn "Validating result..."- validateFloats ref_result result_cuda- where- -- idiom with NOINLINE and extra parameter needed to prevent optimisations- -- from sharing results over multiple runs- {-# NOINLINE saxpy_ref' #-}- saxpy_ref' a arr1 arr2 () = saxpy_ref a arr1 arr2- {-# NOINLINE saxpy_interp #-}- saxpy_interp a arr1 arr2 () = Interp.run (saxpy a arr1 arr2)- {-# NOINLINE saxpy_cuda #-}- saxpy_cuda a arr1 arr2 () = CUDA.run (saxpy a arr1 arr2)--test_square :: Int -> IO ()-test_square n- = do- putStrLn "== Square"- putStrLn $ "Generating data (n = " ++ show n ++ ")..."- v1_ref <- randomUVector n- v1 <- convertUVector v1_ref- putStrLn "Running reference code..."- ref_result <- timeUVector $ square_ref' v1_ref- putStrLn "[Interpreter]"- putStrLn "Running Accelerate code..."- result <- timeVector $ square_interp v1- putStrLn "Validating result..."- validateFloats ref_result result- putStrLn "[CUDA]"- putStrLn "Running Accelerate code..."- result_cuda <- timeVector' $ square_cuda v1- putStrLn "Validating result..."- validateFloats ref_result result_cuda- where- -- idiom with NOINLINE and extra parameter needed to prevent optimisations- -- from sharing results over multiple runs- {-# NOINLINE square_ref' #-}- square_ref' arr1 () = square_ref arr1- {-# NOINLINE square_interp #-}- square_interp arr1 () = Interp.run (square arr1)- {-# NOINLINE square_cuda #-}- square_cuda arr1 () = CUDA.run (square arr1)--test_sum :: Int -> IO ()-test_sum n- = do- putStrLn "== Sum"- putStrLn $ "Generating data (n = " ++ show n ++ ")..."- v1_ref <- randomUVector n- v1 <- convertUVector v1_ref- putStrLn "Running reference code..."- ref_result <- timeUScalar $ sum_ref' v1_ref- putStrLn "[Interpreter]"- putStrLn "Running Accelerate code..."- result <- timeScalar $ sum_interp v1- putStrLn "Validating result..."- validateFloats ref_result result- putStrLn "[CUDA]"- putStrLn "Running Accelerate code..."- result_cuda <- timeScalar' $ sum_cuda v1- putStrLn "Validating result..."- validateFloats ref_result result_cuda- where- -- idiom with NOINLINE and extra parameter needed to prevent optimisations- -- from sharing results over multiple runs- {-# NOINLINE sum_ref' #-}- sum_ref' arr1 () = sum_ref arr1- {-# NOINLINE sum_interp #-}- sum_interp arr1 () = Interp.run (Sum.sum arr1)- {-# NOINLINE sum_cuda #-}- sum_cuda arr1 () = CUDA.run (Sum.sum arr1)--test_dotp :: Int -> IO ()-test_dotp n- = do- putStrLn "== Dot product"- putStrLn $ "Generating data (n = " ++ show n ++ ")..."- v1_ref <- randomUVector n- v1 <- convertUVector v1_ref- v2_ref <- randomUVector n- v2 <- convertUVector v2_ref- putStrLn "Running reference code..."- ref_result <- timeUScalar $ dotp_ref' v1_ref v2_ref- putStrLn "Running Accelerate code..."- putStrLn "[Interpreter]"- result <- timeScalar $ dotp_interp v1 v2- putStrLn "Validating result..."- validateFloats ref_result result- putStrLn "[CUDA]"- result_cuda <- timeScalar' $ dotp_cuda v1 v2- putStrLn "Validating result..."- validateFloats ref_result result_cuda- where- -- idiom with NOINLINE and extra parameter needed to prevent optimisations- -- from sharing results over multiple runs- {-# NOINLINE dotp_ref' #-}- dotp_ref' arr1 arr2 () = dotp_ref arr1 arr2- {-# NOINLINE dotp_interp #-}- dotp_interp arr1 arr2 () = Interp.run (dotp arr1 arr2)- {-# NOINLINE dotp_cuda #-}- dotp_cuda arr1 arr2 () = CUDA.run (dotp arr1 arr2)--test_filter :: Int -> IO ()-test_filter n- = do- putStrLn "== Filter"- putStrLn $ "Generating data (n = " ++ show n ++ ")..."- v_ref <- randomUVector n- v <- convertUVector v_ref- putStrLn "Running reference code..."- ref_result <- timeUVector $ filter_ref' (< 0) v_ref- putStrLn "Running Accelerate code..."- putStrLn "[Interpreter]"- result <- timeVector $ filter_interp (Acc.<* Acc.constant 0) (Acc.use v)- putStrLn "Validating result..."- validateFloats ref_result result- putStrLn "[CUDA]"- result_cuda <- timeVector' $ filter_cuda (Acc.<* Acc.constant 0) (Acc.use v)- putStrLn "Validating result..."- validateFloats ref_result result_cuda- where- -- idiom with NOINLINE and extra parameter needed to prevent optimisations- -- from sharing results over multiple runs- {-# NOINLINE filter_ref' #-}- filter_ref' p arr () = filter_ref p arr- {-# NOINLINE filter_interp #-}- filter_interp p arr () = Interp.run (filter p arr)- {-# NOINLINE filter_cuda #-}- filter_cuda p arr () = CUDA.run (filter p arr)--main :: IO ()-main- = do- putStrLn "Data.Array.Accelerate: simple examples"- putStrLn "--------------------------------------"- - test_saxpy 100000- test_dotp 100000- test_filter 1800
examples/simple/Makefile view
@@ -1,17 +1,17 @@-HCFLAGS = -O -package accelerate +GHC = ghc+HCFLAGS = -O2 -Wall -package accelerate+SRCDIR = src+BUILDDIR = dist+HSMAIN = src/Main.hs+TARGET = test++ all:- ghc $(HCFLAGS) -c Time.hs- ghc $(HCFLAGS) -c Square.hs- ghc $(HCFLAGS) -c Sum.hs- ghc $(HCFLAGS) -c SAXPY.hs- ghc $(HCFLAGS) -c SMVM.hs- ghc $(HCFLAGS) -c DotP.hs- ghc $(HCFLAGS) -c Filter.hs- ghc $(HCFLAGS) -c Main.hs- ghc $(HCFLAGS) -o test Main.o Time.o Square.o Sum.o SAXPY.o SMVM.o DotP.o \- Filter.o+ @mkdir -p $(BUILDDIR)+ $(GHC) --make $(HCFLAGS) -odir $(BUILDDIR) -hidir $(BUILDDIR) -i$(SRCDIR) $(HSMAIN) -o $(TARGET) clean:- rm -f *.o *.hi test- + $(RM) -r $(BUILDDIR)+ $(RM) $(TARGET)+
− examples/simple/SAXPY.hs
@@ -1,24 +0,0 @@-{-# LANGUAGE ParallelListComp #-}--module SAXPY (saxpy, saxpy_ref) where--import Prelude hiding (replicate, zip, map, filter, max, min, not, zipWith)-import qualified Prelude--import Data.Array.Unboxed-import Data.Array.IArray--import Data.Array.Accelerate--saxpy :: Float -> Vector Float -> Vector Float -> Acc (Vector Float)-saxpy alpha xs ys- = let- xs' = use xs- ys' = use ys- in - zipWith (\x y -> constant alpha * x + y) xs' ys'--saxpy_ref :: Float -> UArray Int Float -> UArray Int Float -> UArray Int Float-saxpy_ref alpha xs ys- = listArray (bounds xs) [alpha * x + y | x <- elems xs | y <- elems ys]-
− examples/simple/SMVM.hs
@@ -1,39 +0,0 @@-{-# LANGUAGE ParallelListComp #-}--module SMVM (smvm, smvm_ref) where--import Prelude hiding (replicate, zip, unzip, map, filter, max, min, not,- zipWith)-import qualified Prelude--import Data.Array.Unboxed hiding ((!))-import Data.Array.IArray hiding ((!))--import Data.Array.Accelerate---type SparseVector a = Vector (Int, a)-type SparseMatrix a = (Segments, SparseVector a)--smvm :: SparseMatrix Float -> Vector Float -> Acc (Vector Float)-smvm (segd', smat') vec'- = let- segd = use segd'- (inds, vals) = unzip (use smat')- vec = use vec'- ---- vecVals = backpermute (shape inds) (\i -> inds!i) vec- products = zipWith (*) vecVals vals- in- foldSeg (+) 0 products segd---type USparseMatrix a = (UArray Int Int, (UArray Int Int, UArray Int a))--smvm_ref :: USparseMatrix Float - -> UArray Int Float - -> UArray Int Float-smvm_ref (segd, (inds, values)) vec- -- FIXME: implement reference version of the algorithm- = undefined --listArray ((), ()) $ [sum [x * y | x <- elems xs | y <- elems ys]]-
− examples/simple/Square.hs
@@ -1,23 +0,0 @@-{-# LANGUAGE ParallelListComp #-}--module Square (square, square_ref) where--import Prelude hiding (replicate, zip, map, filter, max, min, not, zipWith)-import qualified Prelude--import Data.Array.Unboxed-import Data.Array.IArray--import Data.Array.Accelerate--square :: Vector Float -> Acc (Vector Float)-square xs- = let- xs' = use xs- in - map (\x -> x * x) xs'--square_ref :: UArray Int Float -> UArray Int Float-square_ref xs- = listArray (bounds xs) [x * x| x <- elems xs]-
− examples/simple/Sum.hs
@@ -1,23 +0,0 @@-{-# LANGUAGE ParallelListComp #-}--module Sum (sum, sum_ref) where--import Prelude hiding (replicate, zip, map, filter, max, min, not, zipWith, sum)-import qualified Prelude--import Data.Array.Unboxed-import Data.Array.IArray--import Data.Array.Accelerate--sum :: Vector Float -> Acc (Scalar Float)-sum xs- = let- xs' = use xs- in - fold (\x y -> x + y) (constant 0.0) xs'--sum_ref :: UArray Int Float -> UArray () Float-sum_ref xs- = listArray ((), ()) $ [Prelude.sum $ elems xs]-
− examples/simple/Time.hs
@@ -1,108 +0,0 @@--- |Auxiliary functions to time benchmarks------ Copyright (c) [2007..2009] Roman Leshchinskiy, Manuel M T Chakravarty------ License: BSD3------- Description -----------------------------------------------------------------module Time (- Time,- getTime,- wallTime, cpuTime,- picoseconds, milliseconds, seconds,-- minus, plus, div,- min, max, avg,- sum, minimum, maximum, average,- - showTime, showMinAvgMax-) where--import System.CPUTime-import System.Time--import Prelude hiding (div, min, max, sum, minimum, maximum)-import qualified Prelude as P--infixl 6 `plus`, `minus`-infixl 7 `div`--data Time = Time { cpu_time :: Integer- , wall_time :: Integer- }--type TimeUnit = Integer -> Integer--picoseconds :: TimeUnit-picoseconds = id--milliseconds :: TimeUnit-milliseconds n = n `P.div` 1000000000--seconds :: TimeUnit-seconds n = n `P.div` 1000000000000--cpuTime :: TimeUnit -> Time -> Integer-cpuTime f = f . cpu_time--wallTime :: TimeUnit -> Time -> Integer-wallTime f = f . wall_time--getTime :: IO Time-getTime =- do- cpu <- getCPUTime- TOD sec pico <- getClockTime- return $ Time cpu (pico + sec * 1000000000000)--zipT :: (Integer -> Integer -> Integer) -> Time -> Time -> Time-zipT f (Time cpu1 wall1) (Time cpu2 wall2) =- Time (f cpu1 cpu2) (f wall1 wall2)--minus :: Time -> Time -> Time-minus = zipT (-)--plus :: Time -> Time -> Time-plus = zipT (+)--div :: Time -> Int -> Time-div (Time cpu clock) n = Time (cpu `P.div` n') (clock `P.div` n')- where- n' = toInteger n--min :: Time -> Time -> Time-min = zipT P.min--max :: Time -> Time -> Time-max = zipT P.max--avg :: Time -> Time -> Time-avg t1 t2 = (t1 `plus` t2) `div` 2--sum :: [Time] -> Time-sum = foldr1 plus--minimum :: [Time] -> Time-minimum = foldr1 min--maximum :: [Time] -> Time-maximum = foldr1 max--average :: [Time] -> Time-average ts = sum ts `div` length ts--showTime :: TimeUnit -> Time -> String-showTime f t = show (wallTime f t) ++ "; " ++ show (cpuTime f t)- -showMinAvgMax :: TimeUnit -> [Time] -> String-showMinAvgMax f ts = show (wallTime f min) ++ "/" ++ - show (wallTime f avg) ++ "/" ++- show (wallTime f max) ++ " - " ++- show (cpuTime f min) ++ "/" ++ - show (cpuTime f avg) ++ "/" ++- show (cpuTime f max)- where- min = minimum ts- avg = average ts- max = maximum ts
+ examples/simple/src/DotP.hs view
@@ -0,0 +1,22 @@+{-# LANGUAGE ParallelListComp #-}++module DotP (dotp, dotp_ref) where++import Prelude hiding (zipWith)++import Data.Array.Unboxed+import Data.Array.Accelerate++dotp :: Vector Float -> Vector Float -> Acc (Scalar Float)+dotp xs ys + = let+ xs' = use xs+ ys' = use ys+ in+ fold (+) 0 (zipWith (*) xs' ys')++dotp_ref :: UArray Int Float + -> UArray Int Float + -> UArray () Float+dotp_ref xs ys + = listArray ((), ()) $ [sum [x * y | x <- elems xs | y <- elems ys]]
+ examples/simple/src/Filter.hs view
@@ -0,0 +1,40 @@+{-# LANGUAGE FlexibleContexts #-}++module Filter (filter, filter_ref) where++import Prelude hiding (length, scanl, map, filter)+import qualified Prelude++import Data.Array.Unboxed (UArray)+import Data.Array.IArray (IArray)+import qualified Data.Array.IArray as IArray++import Data.Array.Accelerate+++filter :: Elem a + => (Exp a -> Exp Bool) + -> Vector a+ -> Acc (Vector a)+filter p vec+ = let arr = use vec+ flags = map (boolToInt . p) arr+ (targetIdx, length) = scanl (+) 0 flags+ arr' = backpermute (length!(constant ())) id arr+ in+ permute const arr' (\ix -> flags!ix ==* 0 ? (ignore, targetIdx!ix)) arr+ -- FIXME: This is abusing 'permute' in that the first two arguments are+ -- only justified because we know the permutation function will+ -- write to each location in the target exactly once.+ -- Instead, we should have a primitive that directly encodes the+ -- compaction pattern of the permutation function.++ +filter_ref :: IArray UArray e + => (e -> Bool)+ -> UArray Int e + -> UArray Int e+filter_ref p xs+ = let xs' = Prelude.filter p (IArray.elems xs)+ in+ IArray.listArray (0, Prelude.length xs' - 1) xs'
+ examples/simple/src/Main.hs view
@@ -0,0 +1,123 @@+{-# LANGUAGE FlexibleContexts #-}++module Main where++import Prelude hiding (filter)++import DotP+import Filter+import Random+import SAXPY+import SMVM++import Data.Array.Accelerate (Acc)+import qualified Data.Array.Accelerate as Acc+import qualified Data.Array.Accelerate.CUDA as CUDA+import qualified Data.Array.Accelerate.Interpreter as Interp++import Data.Array.Unboxed (IArray, UArray, Ix, elems, indices, (!))+import System.Random.MWC (create, uniform, GenIO)+import Control.Exception (evaluate)+import Control.DeepSeq+import Criterion.Main+++instance (Ix dim, IArray UArray e) => NFData (UArray dim e) where+ rnf a = a ! (head (indices a)) `seq` ()+++-- Generate a benchmark test iff the reference and accelerate tests succeed.+--+benchmark+ :: (IArray UArray e, Ix dim, Acc.Ix dim, Acc.Elem e)+ => String+ -> (e -> e -> Bool)+ -> (() -> UArray dim e)+ -> (() -> Acc (Acc.Array dim e))+ -> IO Benchmark++benchmark name sim ref acc = do+ putStr "Interpreter : " ; v1 <- validate sim (ref ()) (Acc.toIArray $ Interp.run (acc ()))+ putStr "CUDA : " ; v2 <- validate sim (ref ()) . Acc.toIArray =<< (CUDA.run (acc ()))+ if not (v1 && v2)+ then return $ bgroup "" []+ else return $ bgroup name+ [ bench "ref" $ nf ref ()+ , bench "cuda" $ (CUDA.run . acc) ()+ ]+++-- Tests+--+test_dotp :: GenIO -> Int -> IO Benchmark+test_dotp gen n = do+ putStrLn $ "== Dot Product (n = " ++ shows n ") =="+ xs <- randomVector gen id n+ ys <- randomVector gen id n+ xs' <- convertVector xs+ ys' <- convertVector ys+ benchmark "dotp" similar (run_ref xs ys) (run_acc xs' ys')+ where+ run_ref x y () = dotp_ref x y+ run_acc x y () = dotp x y+++test_saxpy :: GenIO -> Int -> IO Benchmark+test_saxpy gen n = do+ putStrLn $ "== SAXPY (n = " ++ shows n ") =="+ xs <- randomVector gen id n+ ys <- randomVector gen id n+ xs' <- convertVector xs+ ys' <- convertVector ys+ alpha <- uniform gen+ benchmark "saxpy" similar (run_ref alpha xs ys) (run_acc alpha xs' ys')+ where+ run_ref alpha x y () = saxpy_ref alpha x y+ run_acc alpha x y () = saxpy alpha x y+++test_filter :: GenIO -> Int -> IO Benchmark+test_filter gen n = do+ putStrLn $ "== Filter (n = " ++ shows n ") =="+ xs <- randomVector gen id n :: IO (UArray Int Float)+ xs' <- convertVector xs+ benchmark "filter" similar (run_ref xs) (run_acc xs')+ where+ run_ref x () = filter_ref (< 0.5) x+ run_acc x () = filter (Acc.<* 0.5) x+++test_smvm :: GenIO -> (Int,Int) -> (Int,Int) -> IO Benchmark+test_smvm gen (n,m) (rows,cols) = do+ putStr $ "== SMVM (" ++ shows rows " x " ++ shows cols ", "+ vec <- randomVector gen id cols+ segd <- randomVector gen (\x -> (abs x `rem` (m-n)) + n) rows+ let nnz = sum (elems segd)+ putStrLn $ shows nnz " non-zeros) =="+ inds <- randomVector gen (\x -> abs x `rem` cols) nnz+ vals <- randomVector gen id nnz+ segd' <- convertVector segd+ vec' <- convertVector vec+ mat' <- let v = Acc.fromList nnz (zip (elems inds) (elems vals))+ in evaluate (v `Acc.indexArray` 0) >> return v+ benchmark "smvm" similar (run_ref segd inds vals vec) (run_acc segd' mat' vec')+ where+ run_ref d i x v () = smvm_ref (d, (i,x)) v+ run_acc d x v () = smvm (d,x) v+++-- Main+--+main :: IO ()+main = do+ putStrLn "Data.Array.Accelerate: simple examples"+ putStrLn "--------------------------------------"++ gen <- create+ defaultMain =<< sequence+ [ test_dotp gen 100000+ , test_saxpy gen 100000+ , test_filter gen 1800+ , test_smvm gen (0,42) (2400,400)+ ]+
+ examples/simple/src/Random.hs view
@@ -0,0 +1,60 @@+{-# LANGUAGE BangPatterns, FlexibleContexts, ScopedTypeVariables #-}+{-# LANGUAGE ParallelListComp #-}++module Random where++import System.Random.MWC+import Data.Array.IArray+import Data.Array.Unboxed (UArray)+import Data.Array.IO (MArray, IOUArray)+import Control.Exception (evaluate)+import qualified Data.Array.MArray as M+import qualified Data.Array.Accelerate as Acc+++-- Convert an unboxed array to an Accelerate array+--+convertVector :: (IArray UArray e, Acc.Elem e) => UArray Int e -> IO (Acc.Vector e)+convertVector v =+ let arr = Acc.fromIArray v+ in evaluate (arr `Acc.indexArray` 0) >> return arr+++-- Generate a random, uniformly distributed vector of specified size. The second+-- argument can be used to modify the generated value (e.g. map into a certain+-- range)+--+randomVector+ :: (Variate a, MArray IOUArray e IO, IArray UArray e)+ => GenIO -> (a -> e) -> Int -> IO (UArray Int e)++randomVector gen f n = do+ mu <- M.newArray_ (0,n-1) :: MArray IOUArray e IO => IO (IOUArray Int e)+ let go !i | i < n = uniform gen >>= (\e -> M.writeArray mu i (f e)) >> go (i+1)+ | otherwise = M.unsafeFreeze mu+ go 0+++-- Compare two vectors element-wise for equality, for a given measure of+-- similarity. The index and values are printed for pairs that fail.+--+validate+ :: (IArray UArray e, Ix ix, Show e, Show ix)+ => (e -> e -> Bool) -> UArray ix e -> UArray ix e -> IO Bool++validate f ref arr =+ let sim = filter (not . null) [ if f x y then [] else ">>> " ++ shows i ": " ++ show (x,y)+ | (i,x) <- assocs ref+ | y <- elems arr ]+ in if null sim+ then putStrLn "Valid" >> return True+ else mapM_ putStrLn ("INVALID!" : sim) >> return False+++-- Floating point equality with relative tolerance+--+similar :: (Fractional a, Ord a) => a -> a -> Bool+similar x y =+ let epsilon = 0.0001+ in abs ((x-y) / (x+y+epsilon)) < epsilon+
+ examples/simple/src/SAXPY.hs view
@@ -0,0 +1,21 @@+{-# LANGUAGE ParallelListComp #-}++module SAXPY (saxpy, saxpy_ref) where++import Prelude hiding (replicate, zip, map, filter, max, min, not, zipWith)++import Data.Array.Unboxed+import Data.Array.Accelerate++saxpy :: Float -> Vector Float -> Vector Float -> Acc (Vector Float)+saxpy alpha xs ys+ = let+ xs' = use xs+ ys' = use ys+ in + zipWith (\x y -> constant alpha * x + y) xs' ys'++saxpy_ref :: Float -> UArray Int Float -> UArray Int Float -> UArray Int Float+saxpy_ref alpha xs ys+ = listArray (bounds xs) [alpha * x + y | x <- elems xs | y <- elems ys]+
+ examples/simple/src/SMVM.hs view
@@ -0,0 +1,39 @@+{-# LANGUAGE ParallelListComp #-}++module SMVM (smvm, smvm_ref) where++import Data.Array.Unboxed++import Data.Array.Accelerate (Vector, Segments, Acc)+import qualified Data.Array.Accelerate as Acc+++type SparseVector a = Vector (Int, a)+type SparseMatrix a = (Segments, SparseVector a)++smvm :: SparseMatrix Float -> Vector Float -> Acc (Vector Float)+smvm (segd', smat') vec'+ = let+ segd = Acc.use segd'+ (inds, vals) = Acc.unzip (Acc.use smat')+ vec = Acc.use vec'+ ---+ vecVals = Acc.backpermute (Acc.shape inds) (\i -> inds Acc.! i) vec+ products = Acc.zipWith (*) vecVals vals+ in+ Acc.foldSeg (+) 0 products segd+++type USparseMatrix a = (UArray Int Int, (UArray Int Int, UArray Int a))++smvm_ref :: USparseMatrix Float + -> UArray Int Float + -> UArray Int Float+smvm_ref (segd, (inds, values)) vec+ = listArray (0, rangeSize (bounds segd) - 1)+ $ [sum [ values!i * vec!(inds!i) | i <- range seg] | seg <- segd' ]+ where+ segbegin = scanl (+) 0 $ elems segd+ segend = scanl1 (+) $ elems segd+ segd' = zipWith (\x y -> (x,y-1)) segbegin segend+
+ examples/simple/src/Square.hs view
@@ -0,0 +1,23 @@+{-# LANGUAGE ParallelListComp #-}++module Square (square, square_ref) where++import Prelude hiding (replicate, zip, map, filter, max, min, not, zipWith)+import qualified Prelude++import Data.Array.Unboxed+import Data.Array.IArray++import Data.Array.Accelerate++square :: Vector Float -> Acc (Vector Float)+square xs+ = let+ xs' = use xs+ in + map (\x -> x * x) xs'++square_ref :: UArray Int Float -> UArray Int Float+square_ref xs+ = listArray (bounds xs) [x * x| x <- elems xs]+
+ examples/simple/src/Stencil.hs view
@@ -0,0 +1,77 @@+{-# LANGUAGE ParallelListComp #-}++module Stencil (stencil_test, stencil_test_ref) where++import Prelude hiding (replicate, zip, map, filter, max, min, not, zipWith)+import qualified Prelude++import Data.Array.Unboxed hiding (Array)+import Data.Array.IArray as IArray hiding (Array)++import Data.Array.Accelerate++stencil_test :: Array DIM2 Float -> Acc (Array DIM2 Float)+stencil_test = stencil stencil2D5 Clamp . use++stencil_test_ref :: UArray (Int, Int) Float + -> UArray (Int, Int) Float+stencil_test_ref arr+ = array (bounds arr) [(xy, stencilFun xy) | xy <- indices arr]+ where+ stencilFun (x, y) = (arr IArray.! (clamp (x-1,y)) + + arr IArray.! (clamp (x+1,y)) + + arr IArray.! (clamp (x,y-1)) + + arr IArray.! (clamp (x,y+1)) - + 4*arr IArray.! (clamp (x,y))) / 4+ clamp (x, y) = (minx `Prelude.max` x `Prelude.min` maxx, + miny `Prelude.max` y `Prelude.min` maxy) + where+ ((minx, miny), (maxx, maxy)) = bounds arr++-- some example stencils++stencil1D :: (Elem a, IsFloating a) + => Stencil3 a -> Exp a+stencil1D (x, y, z) = (x + z - 2 * y) / 2++stencil2D5 :: (Elem a, IsFloating a) + => Stencil3x3 a -> Exp a+stencil2D5 ( (_, t, _)+ , (l, m, r)+ , (_, b, _)+ ) + = (t + l + r + b - 4 * m) / 4++stencil2D :: (Elem a, IsFloating a) + => Stencil3x3 a -> Exp a+stencil2D ( (t1, t2, t3)+ , (l , m, r )+ , (b1, b2, b3)+ ) + = (t1/2 + t2 + t3/2 + l + r + b1/2 + b2 + b3/2 - 4 * m) / 4++stencil3D :: (Elem a, IsNum a)+ => Stencil3x3x3 a -> Exp a+stencil3D (front, back, _) = -- 'b4' is the focal point+ let ((f1, f2, _),+ (f3, f4, _),+ _ ) = front+ ((b1, b2, _),+ (b3, b4, _),+ _ ) = back+ in+ f1 + f2 + f3 + f4 + b1 + b2 + b3 + b4++-- usaging them to ensure the types of the example fit the 'stencil' function++use1D :: Acc (Array DIM1 Float) -> Acc (Array DIM1 Float)+use1D arr = stencil stencil1D Clamp arr++use2D5 :: Acc (Array DIM2 Float) -> Acc (Array DIM2 Float)+use2D5 arr = stencil stencil2D5 Clamp arr++use2D :: Acc (Array DIM2 Float) -> Acc (Array DIM2 Float)+use2D arr = stencil stencil2D Clamp arr++use3D :: Acc (Array DIM3 Float) -> Acc (Array DIM3 Float)+use3D arr = stencil stencil3D Clamp arr
+ examples/simple/src/Sum.hs view
@@ -0,0 +1,23 @@+{-# LANGUAGE ParallelListComp #-}++module Sum (sum, sum_ref) where++import Prelude hiding (replicate, zip, map, filter, max, min, not, zipWith, sum)+import qualified Prelude++import Data.Array.Unboxed+import Data.Array.IArray++import Data.Array.Accelerate++sum :: Vector Float -> Acc (Scalar Float)+sum xs+ = let+ xs' = use xs+ in + fold (\x y -> x + y) (constant 0.0) xs'++sum_ref :: UArray Int Float -> UArray () Float+sum_ref xs+ = listArray ((), ()) $ [Prelude.sum $ elems xs]+
+ include/accelerate.h view
@@ -0,0 +1,25 @@++#ifndef NOT_ACCELERATE_MODULE+import qualified Data.Array.Accelerate.Internal.Check as Ck+#endif++#define ERROR(f) (Ck.f __FILE__ __LINE__)+#define ASSERT (Ck.assert __FILE__ __LINE__)+#define ENSURE (Ck.f __FILE__ __LINE__)+#define CHECK(f) (Ck.f __FILE__ __LINE__)++#define BOUNDS_ERROR(f) (ERROR(f) Ck.Bounds)+#define BOUNDS_ASSERT (ASSERT Ck.Bounds)+#define BOUNDS_ENSURE (ENSURE Ck.Bounds)+#define BOUNDS_CHECK(f) (CHECK(f) Ck.Bounds)++#define UNSAFE_ERROR(f) (ERROR(f) Ck.Unsafe)+#define UNSAFE_ASSERT (ASSERT Ck.Unsafe)+#define UNSAFE_ENSURE (ENSURE Ck.Unsafe)+#define UNSAFE_CHECK(f) (CHECK(f) Ck.Unsafe)++#define INTERNAL_ERROR(f) (ERROR(f) Ck.Internal)+#define INTERNAL_ASSERT (ASSERT Ck.Internal)+#define INTERNAL_ENSURE (ENSURE Ck.Internal)+#define INTERNAL_CHECK(f) (CHECK(f) Ck.Internal)+