Skip to content

HTTPS clone URL

Subversion checkout URL

You can clone with
or
.
Download ZIP
Browse files

Quite a few bug fixes. replicate and array slicing now seems to work

  • Loading branch information...
commit 1435f8408e23e7e57cf696736d734c4a057a5b2c 1 parent ec3cb5c
@dybber dybber authored
View
2  Data/Array/Accelerate/OpenCL/CodeGen.hs
@@ -280,7 +280,7 @@ codeGenExp (Cond p t e) =
codeGenExp (Size a) = return $ ccall "size" (codeGenExp (Shape a))
codeGenExp (Shape a)
- | OpenAcc (Avar var) <- a = return $ cvar ("sh" ++ show (idxToInt var))
+ | OpenAcc (Avar var) <- a = return $ cvar ("sh" ++ idxToString var)
| otherwise = INTERNAL_ERROR(error) "codeGenExp" "expected array variable"
idxToString :: forall env t. Idx env t -> String
View
74 Data/Array/Accelerate/OpenCL/CodeGen/Skeleton.hs
@@ -18,8 +18,7 @@ module Data.Array.Accelerate.OpenCL.CodeGen.Skeleton
mkMap, mkZipWith,
-- mkStencil, mkStencil2,
-- mkScanl, mkScanr, mkScanl', mkScanr', mkScanl1, mkScanr1,
- mkPermute, mkBackpermute, mkReplicate
---, mkIndex
+ mkPermute, mkBackpermute, mkReplicate, mkIndex
)
where
@@ -30,8 +29,6 @@ import Language.C.Quote.OpenCL
import Data.Loc
import Data.Symbol
-
---import System.FilePath
--import Data.Array.Accelerate.Type
import Data.Array.Accelerate.OpenCL.CodeGen.Data
import Data.Array.Accelerate.OpenCL.CodeGen.Util
@@ -46,8 +43,8 @@ import Data.Array.Accelerate.OpenCL.CodeGen.Monad
mkGenerate :: ([C.Type],Int) -> C.Exp -> CUTranslSkel
mkGenerate (tyOut, dimOut) apply = runCGM $ do
d_out <- mkOutputTuple tyOut
- mkShape "DimOut" dimOut
- mkShape "TyInA" dimOut
+ mkDim "DimOut" dimOut
+ mkDim "TyInA" dimOut
mkApply 1 apply
@@ -158,9 +155,9 @@ mkZipWith (tyOut,dimOut) (tyInB, dimInB) (tyInA, dimInA) apply =
d_inB <- mkInputTuple "B" tyInB
mkApply 2 apply
- mkShape "DimOut" dimOut
- mkShape "DimInB" dimInB
- mkShape "DimInA" dimInA
+ mkDim "DimOut" dimOut
+ mkDim "DimInB" dimInB
+ mkDim "DimInA" dimInA
ps <- getParams
addDefinitions
@@ -281,8 +278,8 @@ mkZipWith (tyOut,dimOut) (tyInB, dimInB) (tyInA, dimInA) apply =
mkPermute :: [C.Type] -> Int -> Int -> C.Exp -> C.Exp -> CUTranslSkel
mkPermute ty dimOut dimInA combinefn indexfn = runCGM $ do
(d_out, d_inA : _) <- mkTupleTypeAsc 2 ty
- mkShape "DimOut" dimOut
- mkShape "DimInA" dimInA
+ mkDim "DimOut" dimOut
+ mkDim "DimInA" dimInA
mkApply 2 combinefn
mkProject Forward indexfn
@@ -315,8 +312,8 @@ mkPermute ty dimOut dimInA combinefn indexfn = runCGM $ do
mkBackpermute :: [C.Type] -> Int -> Int -> C.Exp -> CUTranslSkel
mkBackpermute ty dimOut dimInA indexFn = runCGM $ do
(d_out, d_inA : _) <- mkTupleTypeAsc 1 ty
- mkShape "DimOut" dimOut
- mkShape "DimInA" dimInA
+ mkDim "DimOut" dimOut
+ mkDim "DimInA" dimInA
mkProject Backward indexFn
@@ -340,34 +337,51 @@ mkBackpermute ty dimOut dimInA indexFn = runCGM $ do
|]
--- -- Multidimensional Index and Replicate
--- -- ------------------------------------
+-- Multidimensional Index and Replicate
+-- ------------------------------------
--- mkIndex :: [CType] -> Int -> Int -> Int -> [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))
+mkIndex :: [C.Type] -> Int -> Int -> Int -> C.Exp -> CUTranslSkel
+mkIndex ty dimSl dimCo dimInA slix = runCGM $ do
+ (d_out, d_inA : _) <- mkTupleTypeAsc 1 ty
+ mkDim "Slice" dimSl
+ mkDim "SliceDim" dimCo
+ mkDim "SliceDim" dimInA
+
+ mkSliceIndex slix
+
+ ps <- getParams
+ addDefinitions
+ [cunit|
+ __kernel void slice (const typename Slice slice,
+ const typename CoSlice slix,
+ const typename SliceDim sliceDim,
+ $params:ps) {
+ const typename Ix shapeSize = $id:(size dimSl)(slice);
+ const typename Ix gridSize = get_global_size(0);
+
+ for (typename Ix ix = get_global_id(0); ix < shapeSize; ix += gridSize) {
+ typename Slice dst = $id:(fromIndex dimSl)(slice, ix);
+ typename SliceDim src = sliceIndex(dst);
+
+ typename Ix j = $id:(toIndex dimInA)(sliceDim, src);
+ set(ix, getA(j, $args:d_inA), $args:d_out) ;
+ }
+ }
+ |]
mkReplicate :: [C.Type] -> Int -> Int -> C.Exp -> CUTranslSkel
mkReplicate ty dimSl dimOut slix = runCGM $ do
(d_out, d_inA : _) <- mkTupleTypeAsc 1 ty
- mkShape "Slice" dimSl
- mkShape "SliceDim" dimOut
+ mkDim "Slice" dimSl
+ mkDim "SliceDim" dimOut
mkSliceReplicate slix
ps <- getParams
addDefinitions
[cunit|
- __kernel void replicate (const typename Slice shOut,
- const typename SliceDim shInA,
+ __kernel void replicate (const typename Slice slice,
+ const typename SliceDim sliceDim,
$params:ps) {
const typename Ix shapeSize = $id:(size dimOut)(sliceDim);
const typename Ix gridSize = get_global_size(0);
View
18 Data/Array/Accelerate/OpenCL/CodeGen/Tuple.hs
@@ -18,6 +18,7 @@ module Data.Array.Accelerate.OpenCL.CodeGen.Tuple
where
import Data.Maybe
+import Data.Char
-- Quasiquotation library
import Language.C.Quote.OpenCL
@@ -58,16 +59,11 @@ mkTupleType subscript types = do
return args
mkTupleTypeAsc :: Int -> [Type] -> CGM (Arguments, [Arguments])
-mkTupleTypeAsc n types = do
- argsOut <- mkOutputTuple types
- argsIn <- mkInputTuples (n-1)
+mkTupleTypeAsc n typ = do
+ argsOut <- mkOutputTuple typ
+ let names = [ [chr $ ord 'A' + i] | i <- [0..n-1]]
+ argsIn <- mapM (flip mkInputTuple typ) names
return $ (argsOut, argsIn)
- where
- mkInputTuples 0 = return []
- mkInputTuples n = do
- as <- mkInputTuples (n-1)
- a <- mkInputTuple (show $ n-1) types
- return $ a : as
mkParameterList :: Maybe String -> Int -> [String] -> CGM (Arguments, [Param])
mkParameterList subscript n tynames = do
@@ -82,10 +78,6 @@ mkParameterList subscript n tynames = do
types' = map (mkPtr . mkGlobal . typename) tynames
args = map (\p -> [cexp|$id:p|]) param_names
- -- accessorCall =
- -- case subscript of
- -- Nothing -> Set $ \idx val -> [cexp|set($id:idx, $id:val, $args:args)|]
- -- Just x -> Get $ \idx -> [cexp|$id:("get" ++ x)($id:idx, $args:args)|]
mkGet :: String -> Int -> [Param] -> CGM ()
mkGet prj n params = do
View
20 Data/Array/Accelerate/OpenCL/CodeGen/Util.hs
@@ -37,14 +37,15 @@ mkApply argc exp
mkProject :: Direction -> Exp -> CGM ()
mkProject Forward exp =
addDefinition $
- (mkDeviceFun "project" (typename "DimOut") $ params [(typename "DimIn0","x0")]) exp
+ (mkDeviceFun "project" (typename "DimOut") $ params [(typename "DimInA","xA")]) exp
mkProject Backward exp =
addDefinition $
- (mkDeviceFun "project" (typename "DimIn0") $ params [(typename "DimOut","x0")]) exp
+ (mkDeviceFun "project" (typename "DimInA") $ params [(typename "DimOut","xA")]) exp
-mkSliceIndex :: Exp -> Definition
-mkSliceIndex =
- mkDeviceFun "sliceIndex" (typename "SliceDim") $ params [(typename "Slice","sl"), (typename "CoSlice","co")]
+mkSliceIndex :: Exp -> CGM ()
+mkSliceIndex exp =
+ addDefinition $
+ (mkDeviceFun "sliceIndex" (typename "SliceDim") $ params [(typename "Slice","sl"), (typename "CoSlice","co")]) exp
mkSliceReplicate :: Exp -> CGM ()
mkSliceReplicate exp =
@@ -83,8 +84,8 @@ fromBool :: Bool -> Exp
fromBool True = [cexp|1|]
fromBool False = [cexp|0|]
-mkDim :: String -> Int -> Definition
-mkDim name n = [cedecl|typedef $ty:dim $id:name;|]
+mkDim :: String -> Int -> CGM ()
+mkDim name n = addDefinition [cedecl|typedef $ty:dim $id:name;|]
where dim = typename $ "DIM" ++ show n
mkVolatile :: Type -> Type
@@ -108,11 +109,6 @@ mkTypedef volatile tyname typ | volatile = let typ' = mkVolatile typ
in [cedecl|typedef $ty:typ' $id:tyname;|]
| otherwise = [cedecl|typedef $ty:typ $id:tyname;|]
-mkShape :: String -> Int -> CGM ()
-mkShape name 0 = addDefinition [cedecl| typedef void* $id:name; |]
-mkShape name 1 = addDefinition [cedecl| typedef $ty:ix $id:name; |]
-mkShape name dim = addDefinition $ mkStruct name False (replicate dim ix)
-
toIndex :: Int -> String
toIndex dim = "toIndexDIM" ++ show dim
View
2  Data/Array/Accelerate/OpenCL/Compile.hs
@@ -595,7 +595,7 @@ compile table key acc fvar = do
-- Compile in another thread
_ <- liftIO . forkIO $ do
let p = (show $ codeGenAcc acc fvar)
--- putStrLn p
+ putStrLn p
prog <- OpenCL.createProgram ctx p
OpenCL.buildProgram prog (map fst devices) =<< compileFlags
putMVar mvar prog
View
9 Data/Array/Accelerate/OpenCL/Execute.hs
@@ -44,6 +44,7 @@ import Control.Monad
import Control.Monad.Trans
import System.IO.Unsafe
+import Foreign.Storable
import Foreign.Ptr (Ptr)
import qualified Foreign.OpenCL.Bindings as OpenCL
@@ -286,7 +287,7 @@ replicateOp :: (Shape dim, Elt slix)
-> CIO (Array dim e)
replicateOp c kernel bindings acc aenv sliceIndex slix (Array sh0 in0) = do
res@(Array sh out) <- newArray c (toElt $ extend sliceIndex (fromElt slix) sh0)
- execute kernel bindings acc aenv (size sh) (((((),out),in0),convertIx sh0),convertIx sh)
+ execute kernel bindings acc aenv (size sh) (((((), convertIx sh0), convertIx sh), out), in0)
freeArray in0
return res
where
@@ -309,7 +310,7 @@ indexOp :: (Shape sl, Elt slix)
indexOp c kernel bindings acc aenv sliceIndex (Array sh0 in0) slix = do
res@(Array sh out) <- newArray c (toElt $ restrict sliceIndex (fromElt slix) sh0)
execute kernel bindings acc aenv (size sh)
- ((((((),out),in0),convertIx sh),convertSlix sliceIndex (fromElt slix)),convertIx sh0)
+ ((((((),convertIx sh),convertSlix sliceIndex (fromElt slix)),convertIx sh0),in0),out)
freeArray in0
return res
where
@@ -700,8 +701,8 @@ instance Marshalable OpenCL.KernelArg where
instance AD.ArrayElt e => Marshalable (AD.ArrayData e) where
marshal = marshalArrayData -- Marshalable (DevicePtrs a) does not type )=
-instance Marshalable a => Marshalable [a] where
- marshal = concatMapM marshal
+instance Storable a => Marshalable [a] where
+ marshal args = return $ [OpenCL.StructArg args]
instance (Marshalable a, Marshalable b) => Marshalable (a,b) where
marshal (a,b) = (++) <$> marshal a <*> marshal b
Please sign in to comment.
Something went wrong with that request. Please try again.