{-# LANGUAGE FlexibleContexts #-}
module Futhark.CodeGen.Backends.CSOpenCL
  ( compileProg
  ) where

import Control.Monad
import Data.List (intersperse)

import Futhark.Representation.KernelsMem (Prog, KernelsMem, int32)
import Futhark.CodeGen.Backends.CSOpenCL.Boilerplate
import qualified Futhark.CodeGen.Backends.GenericCSharp as CS
import qualified Futhark.CodeGen.ImpCode.OpenCL as Imp
import qualified Futhark.CodeGen.ImpGen.OpenCL as ImpGen
import Futhark.CodeGen.Backends.GenericCSharp.AST
import Futhark.CodeGen.Backends.GenericCSharp.Options
import Futhark.CodeGen.Backends.GenericCSharp.Definitions
import Futhark.Util (zEncodeString)
import Futhark.MonadFreshNames


compileProg :: MonadFreshNames m =>
               Maybe String -> Prog KernelsMem -> m String
compileProg :: Maybe String -> Prog KernelsMem -> m String
compileProg Maybe String
module_name Prog KernelsMem
prog = do
  Imp.Program String
opencl_code String
opencl_prelude Map String Safety
kernel_names [PrimType]
types Map Name SizeClass
sizes [FailureMsg]
failures Definitions OpenCL
prog' <-
    Prog KernelsMem -> m Program
forall (m :: * -> *).
MonadFreshNames m =>
Prog KernelsMem -> m Program
ImpGen.compileProg Prog KernelsMem
prog
  Maybe String
-> Constructor
-> [CSStmt]
-> [CSStmt]
-> Operations OpenCL ()
-> ()
-> CompilerM OpenCL () ()
-> [CSStmt]
-> [Space]
-> [Option]
-> Definitions OpenCL
-> m String
forall (m :: * -> *) op s.
MonadFreshNames m =>
Maybe String
-> Constructor
-> [CSStmt]
-> [CSStmt]
-> Operations op s
-> s
-> CompilerM op s ()
-> [CSStmt]
-> [Space]
-> [Option]
-> Definitions op
-> m String
CS.compileProg
    Maybe String
module_name
    Constructor
CS.emptyConstructor
    [CSStmt]
imports
    [CSStmt]
defines
    Operations OpenCL ()
operations
    ()
    (String
-> String
-> Map String Safety
-> [PrimType]
-> Map Name SizeClass
-> [FailureMsg]
-> CompilerM OpenCL () ()
generateBoilerplate String
opencl_code String
opencl_prelude Map String Safety
kernel_names [PrimType]
types Map Name SizeClass
sizes [FailureMsg]
failures)
    []
    [String -> Space
Imp.Space String
"device", String -> Space
Imp.Space String
"local", Space
Imp.DefaultSpace]
    [Option]
cliOptions
    Definitions OpenCL
prog'

  where operations :: CS.Operations Imp.OpenCL ()
        operations :: Operations OpenCL ()
operations = Operations Any Any
forall op s. Operations op s
CS.defaultOperations
                     { opsCompiler :: OpCompiler OpenCL ()
CS.opsCompiler = OpCompiler OpenCL ()
callKernel
                     , opsWriteScalar :: WriteScalar OpenCL ()
CS.opsWriteScalar = WriteScalar OpenCL ()
writeOpenCLScalar
                     , opsReadScalar :: ReadScalar OpenCL ()
CS.opsReadScalar = ReadScalar OpenCL ()
readOpenCLScalar
                     , opsAllocate :: Allocate OpenCL ()
CS.opsAllocate = Allocate OpenCL ()
allocateOpenCLBuffer
                     , opsCopy :: Copy OpenCL ()
CS.opsCopy = Copy OpenCL ()
copyOpenCLMemory
                     , opsStaticArray :: StaticArray OpenCL ()
CS.opsStaticArray = StaticArray OpenCL ()
staticOpenCLArray
                     , opsEntryInput :: EntryInput OpenCL ()
CS.opsEntryInput = EntryInput OpenCL ()
unpackArrayInput
                     , opsEntryOutput :: EntryOutput OpenCL ()
CS.opsEntryOutput = EntryOutput OpenCL ()
packArrayOutput
                     , opsSyncRun :: CSStmt
CS.opsSyncRun = CSStmt
futharkSyncContext
                     }
        imports :: [CSStmt]
imports = [ Maybe String -> String -> CSStmt
Using Maybe String
forall a. Maybe a
Nothing String
"System.Runtime.CompilerServices"
                  , Maybe String -> String -> CSStmt
Using Maybe String
forall a. Maybe a
Nothing String
"Cloo"
                  , Maybe String -> String -> CSStmt
Using Maybe String
forall a. Maybe a
Nothing String
"Cloo.Bindings" ]
        defines :: [CSStmt]
defines = [ String -> CSStmt
Escape String
csOpenCL
                  , String -> CSStmt
Escape String
csMemoryOpenCL ]
cliOptions :: [Option]
cliOptions :: [Option]
cliOptions = [ Option :: String -> Maybe Char -> OptionArgument -> [CSStmt] -> Option
Option { optionLongName :: String
optionLongName = String
"platform"
                      , optionShortName :: Maybe Char
optionShortName = Char -> Maybe Char
forall a. a -> Maybe a
Just Char
'p'
                      , optionArgument :: OptionArgument
optionArgument = OptionArgument
RequiredArgument
                      , optionAction :: [CSStmt]
optionAction = [String -> CSStmt
Escape String
"FutharkContextConfigSetPlatform(ref Cfg, optarg);"]
                      }
             , Option :: String -> Maybe Char -> OptionArgument -> [CSStmt] -> Option
Option { optionLongName :: String
optionLongName = String
"device"
                      , optionShortName :: Maybe Char
optionShortName = Char -> Maybe Char
forall a. a -> Maybe a
Just Char
'd'
                      , optionArgument :: OptionArgument
optionArgument = OptionArgument
RequiredArgument
                      , optionAction :: [CSStmt]
optionAction = [String -> CSStmt
Escape String
"FutharkContextConfigSetDevice(ref Cfg, optarg);"]
                      }
             , Option :: String -> Maybe Char -> OptionArgument -> [CSStmt] -> Option
Option { optionLongName :: String
optionLongName = String
"dump-opencl"
                      , optionShortName :: Maybe Char
optionShortName = Maybe Char
forall a. Maybe a
Nothing
                      , optionArgument :: OptionArgument
optionArgument = OptionArgument
RequiredArgument
                      , optionAction :: [CSStmt]
optionAction = [String -> CSStmt
Escape String
"FutharkContextConfigDumpProgramTo(ref Cfg, optarg);"]
                      }
             , Option :: String -> Maybe Char -> OptionArgument -> [CSStmt] -> Option
Option { optionLongName :: String
optionLongName = String
"load-opencl"
                      , optionShortName :: Maybe Char
optionShortName = Maybe Char
forall a. Maybe a
Nothing
                      , optionArgument :: OptionArgument
optionArgument = OptionArgument
RequiredArgument
                      , optionAction :: [CSStmt]
optionAction = [String -> CSStmt
Escape String
"FutharkContextConfigLoadProgramFrom(ref Cfg, optarg);"]
                      }
             , Option :: String -> Maybe Char -> OptionArgument -> [CSStmt] -> Option
Option { optionLongName :: String
optionLongName = String
"debugging"
                      , optionShortName :: Maybe Char
optionShortName = Char -> Maybe Char
forall a. a -> Maybe a
Just Char
'D'
                      , optionArgument :: OptionArgument
optionArgument = OptionArgument
NoArgument
                      , optionAction :: [CSStmt]
optionAction = [String -> CSStmt
Escape String
"FutharkContextConfigSetDebugging(ref Cfg, true);"]
                      }
             , Option :: String -> Maybe Char -> OptionArgument -> [CSStmt] -> Option
Option { optionLongName :: String
optionLongName = String
"default-group-size"
                      , optionShortName :: Maybe Char
optionShortName = Maybe Char
forall a. Maybe a
Nothing
                      , optionArgument :: OptionArgument
optionArgument = OptionArgument
RequiredArgument
                      , optionAction :: [CSStmt]
optionAction = [String -> CSStmt
Escape String
"FutharkContextConfigSetDefaultGroupSize(ref Cfg, Convert.ToInt32(optarg));"]
                      }
             , Option :: String -> Maybe Char -> OptionArgument -> [CSStmt] -> Option
Option { optionLongName :: String
optionLongName = String
"default-num-groups"
                      , optionShortName :: Maybe Char
optionShortName = Maybe Char
forall a. Maybe a
Nothing
                      , optionArgument :: OptionArgument
optionArgument = OptionArgument
RequiredArgument
                      , optionAction :: [CSStmt]
optionAction = [String -> CSStmt
Escape String
"FutharkContextConfigSetDefaultNumGroups(ref Cfg, Convert.ToInt32(optarg));"]
                      }
             , Option :: String -> Maybe Char -> OptionArgument -> [CSStmt] -> Option
Option { optionLongName :: String
optionLongName = String
"default-tile-size"
                      , optionShortName :: Maybe Char
optionShortName = Maybe Char
forall a. Maybe a
Nothing
                      , optionArgument :: OptionArgument
optionArgument = OptionArgument
RequiredArgument
                      , optionAction :: [CSStmt]
optionAction = [String -> CSStmt
Escape String
"FutharkContextConfigSetDefaultTileSize(ref Cfg, Convert.ToInt32(optarg));"]
                      }
             , Option :: String -> Maybe Char -> OptionArgument -> [CSStmt] -> Option
Option { optionLongName :: String
optionLongName = String
"default-threshold"
                      , optionShortName :: Maybe Char
optionShortName = Maybe Char
forall a. Maybe a
Nothing
                      , optionArgument :: OptionArgument
optionArgument = OptionArgument
RequiredArgument
                      , optionAction :: [CSStmt]
optionAction = [String -> CSStmt
Escape String
"FutharkContextConfigSetDefaultThreshold(ref Cfg, Convert.ToInt32(optarg));"]
                      }
             , Option :: String -> Maybe Char -> OptionArgument -> [CSStmt] -> Option
Option { optionLongName :: String
optionLongName = String
"print-sizes"
                      , optionShortName :: Maybe Char
optionShortName = Maybe Char
forall a. Maybe a
Nothing
                      , optionArgument :: OptionArgument
optionArgument = OptionArgument
NoArgument
                      , optionAction :: [CSStmt]
optionAction = [String -> CSStmt
Escape String
"FutharkConfigPrintSizes();"]
                      }
             , Option :: String -> Maybe Char -> OptionArgument -> [CSStmt] -> Option
Option { optionLongName :: String
optionLongName = String
"size"
                      , optionShortName :: Maybe Char
optionShortName = Maybe Char
forall a. Maybe a
Nothing
                      , optionArgument :: OptionArgument
optionArgument = OptionArgument
RequiredArgument
                      , optionAction :: [CSStmt]
optionAction = [String -> CSStmt
Escape String
"FutharkConfigSetSize(ref Cfg, optarg);"]
                      }
             , Option :: String -> Maybe Char -> OptionArgument -> [CSStmt] -> Option
Option { optionLongName :: String
optionLongName = String
"tuning"
                      , optionShortName :: Maybe Char
optionShortName = Maybe Char
forall a. Maybe a
Nothing
                      , optionArgument :: OptionArgument
optionArgument = OptionArgument
RequiredArgument
                      , optionAction :: [CSStmt]
optionAction = [String -> CSStmt
Escape String
"FutharkConfigLoadTuning(ref Cfg, optarg);"]
                      }
             ]

callKernel :: CS.OpCompiler Imp.OpenCL ()
callKernel :: OpCompiler OpenCL ()
callKernel (Imp.GetSize VName
v Name
key) =
  CSStmt -> CompilerM OpenCL () ()
forall op s. CSStmt -> CompilerM op s ()
CS.stm (CSStmt -> CompilerM OpenCL () ())
-> CSStmt -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ CSExp -> CSExp -> CSStmt
Reassign (String -> CSExp
Var (VName -> String
CS.compileName VName
v)) (CSExp -> CSStmt) -> CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$
    CSExp -> String -> CSExp
Field (String -> CSExp
Var String
"Ctx.Sizes") (String -> CSExp) -> String -> CSExp
forall a b. (a -> b) -> a -> b
$ String -> String
zEncodeString (String -> String) -> String -> String
forall a b. (a -> b) -> a -> b
$ Name -> String
forall a. Pretty a => a -> String
pretty Name
key

callKernel (Imp.GetSizeMax VName
v SizeClass
size_class) = do
  CSExp
v' <- VName -> CompilerM OpenCL () CSExp
forall op s. VName -> CompilerM op s CSExp
CS.compileVar VName
v
  CSStmt -> CompilerM OpenCL () ()
forall op s. CSStmt -> CompilerM op s ()
CS.stm (CSStmt -> CompilerM OpenCL () ())
-> CSStmt -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ CSExp -> CSExp -> CSStmt
Reassign CSExp
v' (CSExp -> CSStmt) -> CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$
    CSExp -> String -> CSExp
Field (String -> CSExp
Var String
"Ctx.OpenCL") (String -> CSExp) -> String -> CSExp
forall a b. (a -> b) -> a -> b
$
    case SizeClass
size_class of SizeClass
Imp.SizeGroup -> String
"MaxGroupSize"
                       SizeClass
Imp.SizeNumGroups -> String
"MaxNumGroups"
                       SizeClass
Imp.SizeTile -> String
"MaxTileSize"
                       Imp.SizeThreshold{} -> String
"MaxThreshold"
                       SizeClass
Imp.SizeLocalMemory -> String
"MaxLocalMemory"
                       Imp.SizeBespoke{} -> String
"MaxBespoke"

callKernel (Imp.LaunchKernel Safety
safety String
name [KernelArg]
args [Exp]
num_workgroups [Exp]
workgroup_size) = do
  [CSExp]
num_workgroups' <- (Exp -> CompilerM OpenCL () CSExp)
-> [Exp] -> CompilerM OpenCL () [CSExp]
forall (t :: * -> *) (m :: * -> *) a b.
(Traversable t, Monad m) =>
(a -> m b) -> t a -> m (t b)
mapM Exp -> CompilerM OpenCL () CSExp
forall op s. Exp -> CompilerM op s CSExp
CS.compileExp [Exp]
num_workgroups
  [CSExp]
workgroup_size' <- (Exp -> CompilerM OpenCL () CSExp)
-> [Exp] -> CompilerM OpenCL () [CSExp]
forall (t :: * -> *) (m :: * -> *) a b.
(Traversable t, Monad m) =>
(a -> m b) -> t a -> m (t b)
mapM Exp -> CompilerM OpenCL () CSExp
forall op s. Exp -> CompilerM op s CSExp
CS.compileExp [Exp]
workgroup_size
  let kernel_size :: [CSExp]
kernel_size = (CSExp -> CSExp -> CSExp) -> [CSExp] -> [CSExp] -> [CSExp]
forall a b c. (a -> b -> c) -> [a] -> [b] -> [c]
zipWith CSExp -> CSExp -> CSExp
mult_exp [CSExp]
num_workgroups' [CSExp]
workgroup_size'
      total_elements :: CSExp
total_elements = (CSExp -> CSExp -> CSExp) -> CSExp -> [CSExp] -> CSExp
forall (t :: * -> *) b a.
Foldable t =>
(b -> a -> b) -> b -> t a -> b
foldl CSExp -> CSExp -> CSExp
mult_exp (Integer -> CSExp
Integer Integer
1) [CSExp]
kernel_size
      cond :: CSExp
cond = String -> CSExp -> CSExp -> CSExp
BinOp String
"!=" CSExp
total_elements (Integer -> CSExp
Integer Integer
0)
  [CSStmt]
body <- CompilerM OpenCL () () -> CompilerM OpenCL () [CSStmt]
forall op s. CompilerM op s () -> CompilerM op s [CSStmt]
CS.collect (CompilerM OpenCL () () -> CompilerM OpenCL () [CSStmt])
-> CompilerM OpenCL () () -> CompilerM OpenCL () [CSStmt]
forall a b. (a -> b) -> a -> b
$ Safety
-> String
-> [CSExp]
-> [CSExp]
-> [KernelArg]
-> CompilerM OpenCL () ()
forall op s.
Safety
-> String -> [CSExp] -> [CSExp] -> [KernelArg] -> CompilerM op s ()
launchKernel Safety
safety String
name [CSExp]
kernel_size [CSExp]
workgroup_size' [KernelArg]
args
  CSStmt -> CompilerM OpenCL () ()
forall op s. CSStmt -> CompilerM op s ()
CS.stm (CSStmt -> CompilerM OpenCL () ())
-> CSStmt -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ CSExp -> [CSStmt] -> [CSStmt] -> CSStmt
If CSExp
cond [CSStmt]
body []
  where mult_exp :: CSExp -> CSExp -> CSExp
mult_exp = String -> CSExp -> CSExp -> CSExp
BinOp String
"*"

callKernel (Imp.CmpSizeLe VName
v Name
key Exp
x) = do
  CSExp
v' <- VName -> CompilerM OpenCL () CSExp
forall op s. VName -> CompilerM op s CSExp
CS.compileVar VName
v
  CSExp
x' <- Exp -> CompilerM OpenCL () CSExp
forall op s. Exp -> CompilerM op s CSExp
CS.compileExp Exp
x
  CSStmt -> CompilerM OpenCL () ()
forall op s. CSStmt -> CompilerM op s ()
CS.stm (CSStmt -> CompilerM OpenCL () ())
-> CSStmt -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ CSExp -> CSExp -> CSStmt
Reassign CSExp
v' (CSExp -> CSStmt) -> CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$
    String -> CSExp -> CSExp -> CSExp
BinOp String
"<=" (CSExp -> String -> CSExp
Field (String -> CSExp
Var String
"Ctx.Sizes") (String -> String
zEncodeString (String -> String) -> String -> String
forall a b. (a -> b) -> a -> b
$ Name -> String
forall a. Pretty a => a -> String
pretty Name
key)) CSExp
x'

launchKernel :: Imp.Safety -> String -> [CSExp] -> [CSExp] -> [Imp.KernelArg] -> CS.CompilerM op s ()
launchKernel :: Safety
-> String -> [CSExp] -> [CSExp] -> [KernelArg] -> CompilerM op s ()
launchKernel Safety
safety String
kernel_name [CSExp]
kernel_dims [CSExp]
workgroup_dims [KernelArg]
args = do
  let kernel_name' :: String
kernel_name' = String
"Ctx."String -> String -> String
forall a. [a] -> [a] -> [a]
++String
kernel_name

  let failure_args :: [CompilerM op s [CSStmt]]
failure_args =
        [ String -> Integer -> CSExp -> CompilerM op s [CSStmt]
forall op s. String -> Integer -> CSExp -> CompilerM op s [CSStmt]
processMemArg String
kernel_name' Integer
0 (CSExp -> CompilerM op s [CSStmt])
-> CSExp -> CompilerM op s [CSStmt]
forall a b. (a -> b) -> a -> b
$ String -> CSExp
Var String
"Ctx.GlobalFailure"
        , String -> Integer -> PrimType -> CSExp -> CompilerM op s [CSStmt]
forall (m :: * -> *).
MonadFreshNames m =>
String -> Integer -> PrimType -> CSExp -> m [CSStmt]
processValueArg String
kernel_name' Integer
1 PrimType
int32 (CSExp -> CompilerM op s [CSStmt])
-> CSExp -> CompilerM op s [CSStmt]
forall a b. (a -> b) -> a -> b
$ String -> CSExp
Var String
"Ctx.GlobalFailureIsAnOption"
        , String -> Integer -> CSExp -> CompilerM op s [CSStmt]
forall op s. String -> Integer -> CSExp -> CompilerM op s [CSStmt]
processMemArg String
kernel_name' Integer
2 (CSExp -> CompilerM op s [CSStmt])
-> CSExp -> CompilerM op s [CSStmt]
forall a b. (a -> b) -> a -> b
$ String -> CSExp
Var String
"Ctx.GlobalFailureArgs"]

  [CSStmt]
failure_args' <- [[CSStmt]] -> [CSStmt]
forall (t :: * -> *) a. Foldable t => t [a] -> [a]
concat ([[CSStmt]] -> [CSStmt])
-> CompilerM op s [[CSStmt]] -> CompilerM op s [CSStmt]
forall (f :: * -> *) a b. Functor f => (a -> b) -> f a -> f b
<$> [CompilerM op s [CSStmt]] -> CompilerM op s [[CSStmt]]
forall (t :: * -> *) (m :: * -> *) a.
(Traversable t, Monad m) =>
t (m a) -> m (t a)
sequence (Int -> [CompilerM op s [CSStmt]] -> [CompilerM op s [CSStmt]]
forall a. Int -> [a] -> [a]
take (Safety -> Int
Imp.numFailureParams Safety
safety) [CompilerM op s [CSStmt]]
forall op s. [CompilerM op s [CSStmt]]
failure_args)

  [[CSStmt]]
args_stms <- (Integer -> KernelArg -> CompilerM op s [CSStmt])
-> [Integer] -> [KernelArg] -> CompilerM op s [[CSStmt]]
forall (m :: * -> *) a b c.
Applicative m =>
(a -> b -> m c) -> [a] -> [b] -> m [c]
zipWithM (String -> Integer -> KernelArg -> CompilerM op s [CSStmt]
forall op s.
String -> Integer -> KernelArg -> CompilerM op s [CSStmt]
processKernelArg String
kernel_name')
               [Int -> Integer
forall a. Integral a => a -> Integer
toInteger (Safety -> Int
Imp.numFailureParams Safety
safety)..] [KernelArg]
args
  CSStmt -> CompilerM op s ()
forall op s. CSStmt -> CompilerM op s ()
CS.stm (CSStmt -> CompilerM op s ()) -> CSStmt -> CompilerM op s ()
forall a b. (a -> b) -> a -> b
$ [CSStmt] -> CSStmt
Unsafe ([CSStmt] -> CSStmt) -> [CSStmt] -> CSStmt
forall a b. (a -> b) -> a -> b
$ [CSStmt]
failure_args' [CSStmt] -> [CSStmt] -> [CSStmt]
forall a. [a] -> [a] -> [a]
++ [[CSStmt]] -> [CSStmt]
forall (t :: * -> *) a. Foldable t => t [a] -> [a]
concat [[CSStmt]]
args_stms

  String
global_work_size <- String -> CompilerM op s String
forall (f :: * -> *). MonadFreshNames f => String -> f String
newVName' String
"GlobalWorkSize"
  String
local_work_size <- String -> CompilerM op s String
forall (f :: * -> *). MonadFreshNames f => String -> f String
newVName' String
"LocalWorkSize"
  String
stop_watch <- String -> CompilerM op s String
forall (f :: * -> *). MonadFreshNames f => String -> f String
newVName' String
"StopWatch"
  String
time_diff <- String -> CompilerM op s String
forall (f :: * -> *). MonadFreshNames f => String -> f String
newVName' String
"TimeDiff"

  let debugStartStmts :: [CSStmt]
debugStartStmts =
        (CSExp -> CSStmt) -> [CSExp] -> [CSStmt]
forall a b. (a -> b) -> [a] -> [b]
map CSExp -> CSStmt
Exp ([CSExp] -> [CSStmt]) -> [CSExp] -> [CSStmt]
forall a b. (a -> b) -> a -> b
$ [String -> [CSExp] -> CSExp
CS.consoleErrorWrite String
"Launching {0} with global work size [" [String -> CSExp
String String
kernel_name]] [CSExp] -> [CSExp] -> [CSExp]
forall a. [a] -> [a] -> [a]
++
                  String -> [CSExp]
printKernelSize String
global_work_size [CSExp] -> [CSExp] -> [CSExp]
forall a. [a] -> [a] -> [a]
++
                  [ String -> [CSExp] -> CSExp
CS.consoleErrorWrite String
"] and local work size [" []] [CSExp] -> [CSExp] -> [CSExp]
forall a. [a] -> [a] -> [a]
++
                  String -> [CSExp]
printKernelSize String
local_work_size [CSExp] -> [CSExp] -> [CSExp]
forall a. [a] -> [a] -> [a]
++
                  [ String -> [CSExp] -> CSExp
CS.consoleErrorWrite String
"].\n" []
                  , CSExp -> CSExp -> [CSArg] -> CSExp
CallMethod (String -> CSExp
Var String
stop_watch) (String -> CSExp
Var String
"Start") []]

  let ctx :: String -> String
ctx = String -> String -> String
forall a. [a] -> [a] -> [a]
(++) String
"Ctx."
  let debugEndStmts :: [CSStmt]
debugEndStmts =
          [ CSExp -> CSStmt
Exp (CSExp -> CSStmt) -> CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$ String -> [CSExp] -> CSExp
CS.simpleCall String
"FutharkContextSync" []
          , CSExp -> CSStmt
Exp (CSExp -> CSStmt) -> CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$ CSExp -> CSExp -> [CSArg] -> CSExp
CallMethod (String -> CSExp
Var String
stop_watch) (String -> CSExp
Var String
"Stop") []
          , CSExp -> CSExp -> CSStmt
Assign (String -> CSExp
Var String
time_diff) (CSExp -> CSStmt) -> CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$ CSExp -> CSExp
asMicroseconds (String -> CSExp
Var String
stop_watch)
          , String -> CSExp -> CSExp -> CSStmt
AssignOp String
"+" (String -> CSExp
Var (String -> CSExp) -> String -> CSExp
forall a b. (a -> b) -> a -> b
$ String -> String
ctx (String -> String) -> String -> String
forall a b. (a -> b) -> a -> b
$ String -> String
kernelRuntime String
kernel_name) (String -> CSExp
Var String
time_diff)
          , String -> CSExp -> CSExp -> CSStmt
AssignOp String
"+" (String -> CSExp
Var (String -> CSExp) -> String -> CSExp
forall a b. (a -> b) -> a -> b
$ String -> String
ctx (String -> String) -> String -> String
forall a b. (a -> b) -> a -> b
$ String -> String
kernelRuns String
kernel_name) (Integer -> CSExp
Integer Integer
1)
          , CSExp -> CSStmt
Exp (CSExp -> CSStmt) -> CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$ String -> [CSExp] -> CSExp
CS.consoleErrorWriteLine String
"kernel {0} runtime: {1}" [String -> CSExp
String String
kernel_name, String -> CSExp
Var String
time_diff]
          ]


  CSStmt -> CompilerM op s ()
forall op s. CSStmt -> CompilerM op s ()
CS.stm (CSStmt -> CompilerM op s ()) -> CSStmt -> CompilerM op s ()
forall a b. (a -> b) -> a -> b
$ CSExp -> [CSStmt] -> [CSStmt] -> CSStmt
If (String -> CSExp -> CSExp -> CSExp
BinOp String
"!=" CSExp
total_elements (Integer -> CSExp
Integer Integer
0))
    ([ CSExp -> CSExp -> CSStmt
Assign (String -> CSExp
Var String
global_work_size) (String -> [CSExp] -> CSExp
Collection String
"IntPtr[]" ([CSExp] -> CSExp) -> [CSExp] -> CSExp
forall a b. (a -> b) -> a -> b
$ (CSExp -> CSExp) -> [CSExp] -> [CSExp]
forall a b. (a -> b) -> [a] -> [b]
map CSExp -> CSExp
CS.toIntPtr [CSExp]
kernel_dims)
     , CSExp -> CSExp -> CSStmt
Assign (String -> CSExp
Var String
local_work_size) (String -> [CSExp] -> CSExp
Collection String
"IntPtr[]" ([CSExp] -> CSExp) -> [CSExp] -> CSExp
forall a b. (a -> b) -> a -> b
$ (CSExp -> CSExp) -> [CSExp] -> [CSExp]
forall a b. (a -> b) -> [a] -> [b]
map CSExp -> CSExp
CS.toIntPtr [CSExp]
workgroup_dims)
     , CSExp -> CSExp -> CSStmt
Assign (String -> CSExp
Var String
stop_watch) (CSExp -> CSStmt) -> CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$ String -> [CSExp] -> CSExp
CS.simpleInitClass String
"Stopwatch" []
     , CSExp -> [CSStmt] -> [CSStmt] -> CSStmt
If (String -> CSExp
Var String
"Ctx.Debugging") [CSStmt]
debugStartStmts []
     ]
     [CSStmt] -> [CSStmt] -> [CSStmt]
forall a. [a] -> [a] -> [a]
++
     [ CSExp -> CSStmt
Exp (CSExp -> CSStmt) -> CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$ String -> [CSExp] -> CSExp
CS.simpleCall String
"OPENCL_SUCCEED" [
         String -> [CSExp] -> CSExp
CS.simpleCall String
"CL10.EnqueueNDRangeKernel"
           [ String -> CSExp
Var String
"Ctx.OpenCL.Queue", String -> CSExp
Var String
kernel_name', Integer -> CSExp
Integer Integer
kernel_rank, CSExp
Null
           , String -> CSExp
Var String
global_work_size, String -> CSExp
Var String
local_work_size, Integer -> CSExp
Integer Integer
0, CSExp
Null, CSExp
Null]]]
     [CSStmt] -> [CSStmt] -> [CSStmt]
forall a. [a] -> [a] -> [a]
++
     [ CSExp -> [CSStmt] -> [CSStmt] -> CSStmt
If (String -> CSExp
Var String
"Ctx.Debugging") [CSStmt]
debugEndStmts [] ]) []

  Bool -> CompilerM op s () -> CompilerM op s ()
forall (f :: * -> *). Applicative f => Bool -> f () -> f ()
when (Safety
safety Safety -> Safety -> Bool
forall a. Ord a => a -> a -> Bool
>= Safety
Imp.SafetyFull) (CompilerM op s () -> CompilerM op s ())
-> CompilerM op s () -> CompilerM op s ()
forall a b. (a -> b) -> a -> b
$
    CSStmt -> CompilerM op s ()
forall op s. CSStmt -> CompilerM op s ()
CS.stm (CSStmt -> CompilerM op s ()) -> CSStmt -> CompilerM op s ()
forall a b. (a -> b) -> a -> b
$ CSExp -> CSExp -> CSStmt
Reassign (String -> CSExp
Var String
"Ctx.GlobalFailureIsAnOption") (Integer -> CSExp
Integer Integer
1)

  CompilerM op s ()
forall op s. CompilerM op s ()
finishIfSynchronous

  where processMemArg :: String -> Integer -> CSExp -> CompilerM op s [CSStmt]
processMemArg String
kernel Integer
argnum CSExp
mem = do
          String
err <- String -> CompilerM op s String
forall (f :: * -> *). MonadFreshNames f => String -> f String
newVName' String
"setargErr"
          VName
dest <- String -> CompilerM op s VName
forall (m :: * -> *). MonadFreshNames m => String -> m VName
newVName String
"kArgDest"
          let err_var :: CSExp
err_var = String -> CSExp
Var String
err
          CSExp
dest' <- VName -> CompilerM op s CSExp
forall op s. VName -> CompilerM op s CSExp
CS.compileVar VName
dest
          [CSStmt] -> CompilerM op s [CSStmt]
forall (m :: * -> *) a. Monad m => a -> m a
return [ CSExp -> CSExp -> [CSStmt] -> CSStmt
Fixed CSExp
dest' (CSExp -> CSExp
Addr CSExp
mem)
                   [ CSExp -> CSExp -> CSStmt
Assign CSExp
err_var (CSExp -> CSStmt) -> CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$ String -> Integer -> CSExp -> CSExp -> CSExp
getKernelCall String
kernel Integer
argnum
                     (CSType -> CSExp
CS.sizeOf (CSType -> CSExp) -> CSType -> CSExp
forall a b. (a -> b) -> a -> b
$ CSPrim -> CSType
Primitive CSPrim
IntPtrT) CSExp
dest']
                 ]

        processValueArg :: String -> Integer -> PrimType -> CSExp -> m [CSStmt]
processValueArg String
kernel Integer
argnum PrimType
et CSExp
e = do
          let t :: CSType
t = PrimType -> CSType
CS.compilePrimTypeToAST PrimType
et
          String
tmp <- String -> m String
forall (f :: * -> *). MonadFreshNames f => String -> f String
newVName' String
"kernelArg"
          String
err <- String -> m String
forall (f :: * -> *). MonadFreshNames f => String -> f String
newVName' String
"setargErr"
          let err_var :: CSExp
err_var = String -> CSExp
Var String
err
          [CSStmt] -> m [CSStmt]
forall (m :: * -> *) a. Monad m => a -> m a
return [ CSType -> CSExp -> Maybe CSExp -> CSStmt
AssignTyped CSType
t (String -> CSExp
Var String
tmp) (CSExp -> Maybe CSExp
forall a. a -> Maybe a
Just CSExp
e)
                 , CSExp -> CSExp -> CSStmt
Assign CSExp
err_var (CSExp -> CSStmt) -> CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$ String -> Integer -> CSExp -> CSExp -> CSExp
getKernelCall String
kernel Integer
argnum (CSType -> CSExp
CS.sizeOf CSType
t) (CSExp -> CSExp
Addr (CSExp -> CSExp) -> CSExp -> CSExp
forall a b. (a -> b) -> a -> b
$ String -> CSExp
Var String
tmp)]

        processKernelArg :: String
                         -> Integer
                         -> Imp.KernelArg
                         -> CS.CompilerM op s [CSStmt]
        processKernelArg :: String -> Integer -> KernelArg -> CompilerM op s [CSStmt]
processKernelArg String
kernel Integer
argnum (Imp.ValueKArg Exp
e PrimType
et) =
          String -> Integer -> PrimType -> CSExp -> CompilerM op s [CSStmt]
forall (m :: * -> *).
MonadFreshNames m =>
String -> Integer -> PrimType -> CSExp -> m [CSStmt]
processValueArg String
kernel Integer
argnum PrimType
et (CSExp -> CompilerM op s [CSStmt])
-> CompilerM op s CSExp -> CompilerM op s [CSStmt]
forall (m :: * -> *) a b. Monad m => (a -> m b) -> m a -> m b
=<< Exp -> CompilerM op s CSExp
forall op s. Exp -> CompilerM op s CSExp
CS.compileExp Exp
e
        processKernelArg String
kernel Integer
argnum (Imp.MemKArg VName
v) =
          String -> Integer -> CSExp -> CompilerM op s [CSStmt]
forall op s. String -> Integer -> CSExp -> CompilerM op s [CSStmt]
processMemArg String
kernel Integer
argnum (CSExp -> CompilerM op s [CSStmt])
-> (CSExp -> CSExp) -> CSExp -> CompilerM op s [CSStmt]
forall b c a. (b -> c) -> (a -> b) -> a -> c
. CSExp -> CSExp
memblockFromMem (CSExp -> CompilerM op s [CSStmt])
-> CompilerM op s CSExp -> CompilerM op s [CSStmt]
forall (m :: * -> *) a b. Monad m => (a -> m b) -> m a -> m b
=<< VName -> CompilerM op s CSExp
forall op s. VName -> CompilerM op s CSExp
CS.compileVar VName
v

        processKernelArg String
kernel Integer
argnum (Imp.SharedMemoryKArg (Imp.Count Exp
num_bytes)) = do
          String
err <- String -> CompilerM op s String
forall (f :: * -> *). MonadFreshNames f => String -> f String
newVName' String
"setargErr"
          let err_var :: CSExp
err_var = String -> CSExp
Var String
err
          CSExp
num_bytes' <- Exp -> CompilerM op s CSExp
forall op s. Exp -> CompilerM op s CSExp
CS.compileExp Exp
num_bytes
          [CSStmt] -> CompilerM op s [CSStmt]
forall (m :: * -> *) a. Monad m => a -> m a
return [ CSExp -> CSExp -> CSStmt
Assign CSExp
err_var (CSExp -> CSStmt) -> CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$ String -> Integer -> CSExp -> CSExp -> CSExp
getKernelCall String
kernel Integer
argnum CSExp
num_bytes' CSExp
Null ]

        kernel_rank :: Integer
kernel_rank = Int -> Integer
forall a. Integral a => a -> Integer
toInteger (Int -> Integer) -> Int -> Integer
forall a b. (a -> b) -> a -> b
$ [CSExp] -> Int
forall (t :: * -> *) a. Foldable t => t a -> Int
length [CSExp]
kernel_dims
        total_elements :: CSExp
total_elements = (CSExp -> CSExp -> CSExp) -> CSExp -> [CSExp] -> CSExp
forall (t :: * -> *) b a.
Foldable t =>
(b -> a -> b) -> b -> t a -> b
foldl (String -> CSExp -> CSExp -> CSExp
BinOp String
"*") (Integer -> CSExp
Integer Integer
1) [CSExp]
kernel_dims

        printKernelSize :: String -> [CSExp]
        printKernelSize :: String -> [CSExp]
printKernelSize String
work_size =
          CSExp -> [CSExp] -> [CSExp]
forall a. a -> [a] -> [a]
intersperse (String -> [CSExp] -> CSExp
CS.consoleErrorWrite String
", " []) ([CSExp] -> [CSExp]) -> [CSExp] -> [CSExp]
forall a b. (a -> b) -> a -> b
$ (Integer -> CSExp) -> [Integer] -> [CSExp]
forall a b. (a -> b) -> [a] -> [b]
map (String -> Integer -> CSExp
forall a. Integral a => String -> a -> CSExp
printKernelDim String
work_size) [Integer
0..Integer
kernel_rankInteger -> Integer -> Integer
forall a. Num a => a -> a -> a
-Integer
1]

        printKernelDim :: String -> a -> CSExp
printKernelDim String
global_work_size a
i =
          String -> [CSExp] -> CSExp
CS.consoleErrorWrite String
"{0}" [CSExp -> CSIdx -> CSExp
Index (String -> CSExp
Var String
global_work_size) (CSExp -> CSIdx
IdxExp (Integer -> CSExp
Integer (Integer -> CSExp) -> Integer -> CSExp
forall a b. (a -> b) -> a -> b
$ a -> Integer
forall a. Integral a => a -> Integer
toInteger a
i))]

        asMicroseconds :: CSExp -> CSExp
asMicroseconds CSExp
watch =
          String -> CSExp -> CSExp -> CSExp
BinOp String
"/" (CSExp -> String -> CSExp
Field CSExp
watch String
"ElapsedTicks")
          (String -> CSExp -> CSExp -> CSExp
BinOp String
"/" (CSExp -> String -> CSExp
Field (String -> CSExp
Var String
"TimeSpan") String
"TicksPerMillisecond") (Integer -> CSExp
Integer Integer
1000))



getKernelCall :: String -> Integer -> CSExp -> CSExp -> CSExp
getKernelCall :: String -> Integer -> CSExp -> CSExp -> CSExp
getKernelCall String
kernel Integer
arg_num CSExp
size CSExp
Null =
  String -> [CSExp] -> CSExp
CS.simpleCall String
"CL10.SetKernelArg" [ String -> CSExp
Var String
kernel, Integer -> CSExp
Integer Integer
arg_num, CSExp -> CSExp
CS.toIntPtr CSExp
size, String -> CSExp
Var String
"Ctx.NULL"]
getKernelCall String
kernel Integer
arg_num CSExp
size CSExp
e =
  String -> [CSExp] -> CSExp
CS.simpleCall String
"CL10.SetKernelArg" [ String -> CSExp
Var String
kernel, Integer -> CSExp
Integer Integer
arg_num, CSExp -> CSExp
CS.toIntPtr CSExp
size, CSExp -> CSExp
CS.toIntPtr CSExp
e]

writeOpenCLScalar :: CS.WriteScalar Imp.OpenCL ()
writeOpenCLScalar :: WriteScalar OpenCL ()
writeOpenCLScalar CSExp
mem CSExp
i PrimType
bt String
"device" CSExp
val = do
  let bt' :: CSType
bt' = PrimType -> CSType
CS.compilePrimTypeToAST PrimType
bt
  String
scalar <- String -> CompilerM OpenCL () String
forall (f :: * -> *). MonadFreshNames f => String -> f String
newVName' String
"scalar"
  String
ptr <- String -> CompilerM OpenCL () String
forall (f :: * -> *). MonadFreshNames f => String -> f String
newVName' String
"ptr"
  CSStmt -> CompilerM OpenCL () ()
forall op s. CSStmt -> CompilerM op s ()
CS.stm (CSStmt -> CompilerM OpenCL () ())
-> CSStmt -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ [CSStmt] -> CSStmt
Unsafe
    [ CSType -> CSExp -> Maybe CSExp -> CSStmt
AssignTyped CSType
bt' (String -> CSExp
Var String
scalar) (CSExp -> Maybe CSExp
forall a. a -> Maybe a
Just CSExp
val)
    , CSType -> CSExp -> Maybe CSExp -> CSStmt
AssignTyped (CSType -> CSType
PointerT CSType
VoidT) (String -> CSExp
Var String
ptr) (CSExp -> Maybe CSExp
forall a. a -> Maybe a
Just (CSExp -> Maybe CSExp) -> CSExp -> Maybe CSExp
forall a b. (a -> b) -> a -> b
$ CSExp -> CSExp
Addr (CSExp -> CSExp) -> CSExp -> CSExp
forall a b. (a -> b) -> a -> b
$ String -> CSExp
Var String
scalar)
    , CSExp -> CSStmt
Exp (CSExp -> CSStmt) -> CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$ String -> [CSExp] -> CSExp
CS.simpleCall String
"CL10.EnqueueWriteBuffer"
        [ String -> CSExp
Var String
"Ctx.OpenCL.Queue", CSExp -> CSExp
memblockFromMem CSExp
mem, Bool -> CSExp
Bool Bool
True
        , CSExp -> CSExp
CS.toIntPtr (CSExp -> CSExp) -> CSExp -> CSExp
forall a b. (a -> b) -> a -> b
$ String -> CSExp -> CSExp -> CSExp
BinOp String
"*" CSExp
i (CSType -> CSExp
CS.sizeOf CSType
bt')
        , CSExp -> CSExp
CS.toIntPtr (CSExp -> CSExp) -> CSExp -> CSExp
forall a b. (a -> b) -> a -> b
$ CSType -> CSExp
CS.sizeOf CSType
bt',CSExp -> CSExp
CS.toIntPtr (CSExp -> CSExp) -> CSExp -> CSExp
forall a b. (a -> b) -> a -> b
$ String -> CSExp
Var String
ptr
    , Integer -> CSExp
Integer Integer
0, CSExp
Null, CSExp
Null]
    ]

writeOpenCLScalar CSExp
_ CSExp
_ PrimType
_ String
space CSExp
_ =
  String -> CompilerM OpenCL () ()
forall a. HasCallStack => String -> a
error (String -> CompilerM OpenCL () ())
-> String -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ String
"Cannot write to '" String -> String -> String
forall a. [a] -> [a] -> [a]
++ String
space String -> String -> String
forall a. [a] -> [a] -> [a]
++ String
"' memory space."

readOpenCLScalar :: CS.ReadScalar Imp.OpenCL ()
readOpenCLScalar :: ReadScalar OpenCL ()
readOpenCLScalar CSExp
mem CSExp
i PrimType
bt String
"device" = do
  String
val <- String -> CompilerM OpenCL () String
forall (f :: * -> *). MonadFreshNames f => String -> f String
newVName' String
"read_res"
  String
ptr <- String -> CompilerM OpenCL () String
forall (f :: * -> *). MonadFreshNames f => String -> f String
newVName' String
"ptr"
  let bt' :: CSType
bt' = PrimType -> CSType
CS.compilePrimTypeToAST PrimType
bt
  CSStmt -> CompilerM OpenCL () ()
forall op s. CSStmt -> CompilerM op s ()
CS.stm (CSStmt -> CompilerM OpenCL () ())
-> CSStmt -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ CSType -> CSExp -> Maybe CSExp -> CSStmt
AssignTyped CSType
bt' (String -> CSExp
Var String
val) (CSExp -> Maybe CSExp
forall a. a -> Maybe a
Just (CSExp -> Maybe CSExp) -> CSExp -> Maybe CSExp
forall a b. (a -> b) -> a -> b
$ String -> [CSExp] -> CSExp
CS.simpleInitClass (CSType -> String
forall a. Pretty a => a -> String
pretty CSType
bt') [])
  CSStmt -> CompilerM OpenCL () ()
forall op s. CSStmt -> CompilerM op s ()
CS.stm (CSStmt -> CompilerM OpenCL () ())
-> CSStmt -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ [CSStmt] -> CSStmt
Unsafe
    [ CSExp -> CSExp -> CSStmt
CS.assignScalarPointer (String -> CSExp
Var String
val) (String -> CSExp
Var String
ptr)
    , CSExp -> CSStmt
Exp (CSExp -> CSStmt) -> CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$ String -> [CSExp] -> CSExp
CS.simpleCall String
"CL10.EnqueueReadBuffer"
      [ String -> CSExp
Var String
"Ctx.OpenCL.Queue", CSExp -> CSExp
memblockFromMem CSExp
mem , Bool -> CSExp
Bool Bool
True
      , CSExp -> CSExp
CS.toIntPtr (CSExp -> CSExp) -> CSExp -> CSExp
forall a b. (a -> b) -> a -> b
$ String -> CSExp -> CSExp -> CSExp
BinOp String
"*" CSExp
i (CSType -> CSExp
CS.sizeOf CSType
bt')
      , CSExp -> CSExp
CS.toIntPtr (CSExp -> CSExp) -> CSExp -> CSExp
forall a b. (a -> b) -> a -> b
$ CSType -> CSExp
CS.sizeOf CSType
bt', CSExp -> CSExp
CS.toIntPtr (CSExp -> CSExp) -> CSExp -> CSExp
forall a b. (a -> b) -> a -> b
$ String -> CSExp
Var String
ptr
      , Integer -> CSExp
Integer Integer
0, CSExp
Null, CSExp
Null]
    ]
  CSExp -> CompilerM OpenCL () CSExp
forall (m :: * -> *) a. Monad m => a -> m a
return (CSExp -> CompilerM OpenCL () CSExp)
-> CSExp -> CompilerM OpenCL () CSExp
forall a b. (a -> b) -> a -> b
$ String -> CSExp
Var String
val

readOpenCLScalar CSExp
_ CSExp
_ PrimType
_ String
space =
  String -> CompilerM OpenCL () CSExp
forall a. HasCallStack => String -> a
error (String -> CompilerM OpenCL () CSExp)
-> String -> CompilerM OpenCL () CSExp
forall a b. (a -> b) -> a -> b
$ String
"Cannot read from '" String -> String -> String
forall a. [a] -> [a] -> [a]
++ String
space String -> String -> String
forall a. [a] -> [a] -> [a]
++ String
"' memory space."

computeErrCodeT :: CSType
computeErrCodeT :: CSType
computeErrCodeT = String -> CSType
CustomT String
"ComputeErrorCode"

allocateOpenCLBuffer :: CS.Allocate Imp.OpenCL ()
allocateOpenCLBuffer :: Allocate OpenCL ()
allocateOpenCLBuffer CSExp
mem CSExp
size String
"device" = do
  String
errcode <- VName -> String
CS.compileName (VName -> String)
-> CompilerM OpenCL () VName -> CompilerM OpenCL () String
forall (f :: * -> *) a b. Functor f => (a -> b) -> f a -> f b
<$> String -> CompilerM OpenCL () VName
forall (m :: * -> *). MonadFreshNames m => String -> m VName
newVName String
"errCode"
  CSStmt -> CompilerM OpenCL () ()
forall op s. CSStmt -> CompilerM op s ()
CS.stm (CSStmt -> CompilerM OpenCL () ())
-> CSStmt -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ CSType -> CSExp -> Maybe CSExp -> CSStmt
AssignTyped CSType
computeErrCodeT (String -> CSExp
Var String
errcode) Maybe CSExp
forall a. Maybe a
Nothing
  CSStmt -> CompilerM OpenCL () ()
forall op s. CSStmt -> CompilerM op s ()
CS.stm (CSStmt -> CompilerM OpenCL () ())
-> CSStmt -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ CSExp -> CSExp -> CSStmt
Reassign CSExp
mem (String -> [CSExp] -> CSExp
CS.simpleCall String
"MemblockAllocDevice" [CSExp -> CSExp
Ref (CSExp -> CSExp) -> CSExp -> CSExp
forall a b. (a -> b) -> a -> b
$ String -> CSExp
Var String
"Ctx", CSExp
mem, CSExp
size, String -> CSExp
String (String -> CSExp) -> String -> CSExp
forall a b. (a -> b) -> a -> b
$ CSExp -> String
forall a. Pretty a => a -> String
pretty CSExp
mem])

allocateOpenCLBuffer CSExp
_ CSExp
_ String
space =
  String -> CompilerM OpenCL () ()
forall a. HasCallStack => String -> a
error (String -> CompilerM OpenCL () ())
-> String -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ String
"Cannot allocate in '" String -> String -> String
forall a. [a] -> [a] -> [a]
++ String
space String -> String -> String
forall a. [a] -> [a] -> [a]
++ String
"' space"

copyOpenCLMemory :: CS.Copy Imp.OpenCL ()
copyOpenCLMemory :: Copy OpenCL ()
copyOpenCLMemory CSExp
destmem CSExp
destidx Space
Imp.DefaultSpace CSExp
srcmem CSExp
srcidx (Imp.Space String
"device") CSExp
nbytes PrimType
_ = do
  String
ptr <- String -> CompilerM OpenCL () String
forall (f :: * -> *). MonadFreshNames f => String -> f String
newVName' String
"ptr"
  CSStmt -> CompilerM OpenCL () ()
forall op s. CSStmt -> CompilerM op s ()
CS.stm (CSStmt -> CompilerM OpenCL () ())
-> CSStmt -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ CSExp -> CSExp -> [CSStmt] -> CSStmt
Fixed (String -> CSExp
Var String
ptr) (CSExp -> CSExp
Addr (CSExp -> CSExp) -> CSExp -> CSExp
forall a b. (a -> b) -> a -> b
$ CSExp -> CSIdx -> CSExp
Index CSExp
destmem (CSIdx -> CSExp) -> CSIdx -> CSExp
forall a b. (a -> b) -> a -> b
$ CSExp -> CSIdx
IdxExp (CSExp -> CSIdx) -> CSExp -> CSIdx
forall a b. (a -> b) -> a -> b
$ Integer -> CSExp
Integer Integer
0)
    [ CSExp -> CSStmt -> CSStmt
ifNotZeroSize CSExp
nbytes (CSStmt -> CSStmt) -> CSStmt -> CSStmt
forall a b. (a -> b) -> a -> b
$
      CSExp -> CSStmt
Exp (CSExp -> CSStmt) -> CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$ String -> [CSExp] -> CSExp
CS.simpleCall String
"CL10.EnqueueReadBuffer"
      [ String -> CSExp
Var String
"Ctx.Opencl.Queue", CSExp -> CSExp
memblockFromMem CSExp
srcmem, Bool -> CSExp
Bool Bool
True
      , CSExp -> CSExp
CS.toIntPtr CSExp
srcidx, CSExp
nbytes,CSExp -> CSExp
CS.toIntPtr (CSExp -> CSExp) -> CSExp -> CSExp
forall a b. (a -> b) -> a -> b
$ String -> CSExp
Var String
ptr
      , CSExp -> CSExp
CS.toIntPtr CSExp
destidx, CSExp
Null, CSExp
Null]
    ]

copyOpenCLMemory CSExp
destmem CSExp
destidx (Imp.Space String
"device") CSExp
srcmem CSExp
srcidx Space
Imp.DefaultSpace CSExp
nbytes PrimType
_ = do
  String
ptr <- String -> CompilerM OpenCL () String
forall (f :: * -> *). MonadFreshNames f => String -> f String
newVName' String
"ptr"
  CSStmt -> CompilerM OpenCL () ()
forall op s. CSStmt -> CompilerM op s ()
CS.stm (CSStmt -> CompilerM OpenCL () ())
-> CSStmt -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ CSExp -> CSExp -> [CSStmt] -> CSStmt
Fixed (String -> CSExp
Var String
ptr) (CSExp -> CSExp
Addr (CSExp -> CSExp) -> CSExp -> CSExp
forall a b. (a -> b) -> a -> b
$ CSExp -> CSIdx -> CSExp
Index CSExp
srcmem (CSIdx -> CSExp) -> CSIdx -> CSExp
forall a b. (a -> b) -> a -> b
$ CSExp -> CSIdx
IdxExp (CSExp -> CSIdx) -> CSExp -> CSIdx
forall a b. (a -> b) -> a -> b
$ Integer -> CSExp
Integer Integer
0)
    [ CSExp -> CSStmt -> CSStmt
ifNotZeroSize CSExp
nbytes (CSStmt -> CSStmt) -> CSStmt -> CSStmt
forall a b. (a -> b) -> a -> b
$
      CSExp -> CSStmt
Exp (CSExp -> CSStmt) -> CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$ String -> [CSExp] -> CSExp
CS.simpleCall String
"CL10.EnqueueWriteBuffer"
        [ String -> CSExp
Var String
"Ctx.OpenCL.Queue", CSExp -> CSExp
memblockFromMem CSExp
destmem, Bool -> CSExp
Bool Bool
True
        , CSExp -> CSExp
CS.toIntPtr CSExp
destidx, CSExp -> CSExp
CS.toIntPtr CSExp
nbytes, CSExp -> CSExp
CS.toIntPtr (CSExp -> CSExp) -> CSExp -> CSExp
forall a b. (a -> b) -> a -> b
$ String -> CSExp
Var String
ptr
        , CSExp
srcidx, CSExp
Null, CSExp
Null]
    ]

copyOpenCLMemory CSExp
destmem CSExp
destidx (Imp.Space String
"device") CSExp
srcmem CSExp
srcidx (Imp.Space String
"device") CSExp
nbytes PrimType
_ = do
  CSStmt -> CompilerM OpenCL () ()
forall op s. CSStmt -> CompilerM op s ()
CS.stm (CSStmt -> CompilerM OpenCL () ())
-> CSStmt -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ CSExp -> CSStmt -> CSStmt
ifNotZeroSize CSExp
nbytes (CSStmt -> CSStmt) -> CSStmt -> CSStmt
forall a b. (a -> b) -> a -> b
$
    CSExp -> CSStmt
Exp (CSExp -> CSStmt) -> CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$ String -> [CSExp] -> CSExp
CS.simpleCall String
"CL10.EnqueueCopyBuffer"
      [ String -> CSExp
Var String
"Ctx.OpenCL.Queue", CSExp -> CSExp
memblockFromMem CSExp
srcmem, CSExp -> CSExp
memblockFromMem CSExp
destmem
      , CSExp -> CSExp
CS.toIntPtr CSExp
srcidx, CSExp -> CSExp
CS.toIntPtr CSExp
destidx, CSExp -> CSExp
CS.toIntPtr CSExp
nbytes
      , Integer -> CSExp
Integer Integer
0, CSExp
Null, CSExp
Null]
  CompilerM OpenCL () ()
forall op s. CompilerM op s ()
finishIfSynchronous

copyOpenCLMemory CSExp
destmem CSExp
destidx Space
Imp.DefaultSpace CSExp
srcmem CSExp
srcidx Space
Imp.DefaultSpace CSExp
nbytes PrimType
_ =
  CSExp -> CSExp -> CSExp -> CSExp -> CSExp -> CompilerM OpenCL () ()
forall op s.
CSExp -> CSExp -> CSExp -> CSExp -> CSExp -> CompilerM op s ()
CS.copyMemoryDefaultSpace CSExp
destmem CSExp
destidx CSExp
srcmem CSExp
srcidx CSExp
nbytes

copyOpenCLMemory CSExp
_ CSExp
_ Space
destspace CSExp
_ CSExp
_ Space
srcspace CSExp
_ PrimType
_=
  String -> CompilerM OpenCL () ()
forall a. HasCallStack => String -> a
error (String -> CompilerM OpenCL () ())
-> String -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ String
"Cannot copy to " String -> String -> String
forall a. [a] -> [a] -> [a]
++ Space -> String
forall a. Show a => a -> String
show Space
destspace String -> String -> String
forall a. [a] -> [a] -> [a]
++ String
" from " String -> String -> String
forall a. [a] -> [a] -> [a]
++ Space -> String
forall a. Show a => a -> String
show Space
srcspace

staticOpenCLArray :: CS.StaticArray Imp.OpenCL ()
staticOpenCLArray :: StaticArray OpenCL ()
staticOpenCLArray VName
name String
"device" PrimType
t ArrayContents
vs = do
  CSExp
name' <- VName -> CompilerM OpenCL () CSExp
forall op s. VName -> CompilerM op s CSExp
CS.compileVar VName
name
  CSStmt -> CompilerM OpenCL () ()
forall op s. CSStmt -> CompilerM op s ()
CS.staticMemDecl (CSStmt -> CompilerM OpenCL () ())
-> CSStmt -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ CSType -> CSExp -> Maybe CSExp -> CSStmt
AssignTyped (String -> CSType
CustomT String
"OpenCLMemblock") CSExp
name' Maybe CSExp
forall a. Maybe a
Nothing

  -- Create host-side C# array with intended values.
  String
tmp_arr <- String -> CompilerM OpenCL () String
forall (f :: * -> *). MonadFreshNames f => String -> f String
newVName' String
"tmpArr"
  let t' :: CSType
t' = PrimType -> CSType
CS.compilePrimTypeToAST PrimType
t
  CSStmt -> CompilerM OpenCL () ()
forall op s. CSStmt -> CompilerM op s ()
CS.staticMemDecl (CSStmt -> CompilerM OpenCL () ())
-> CSStmt -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ CSType -> CSExp -> Maybe CSExp -> CSStmt
AssignTyped (CSComp -> CSType
Composite (CSComp -> CSType) -> CSComp -> CSType
forall a b. (a -> b) -> a -> b
$ CSType -> CSComp
ArrayT CSType
t') (String -> CSExp
Var String
tmp_arr) (Maybe CSExp -> CSStmt) -> Maybe CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$ CSExp -> Maybe CSExp
forall a. a -> Maybe a
Just (CSExp -> Maybe CSExp) -> CSExp -> Maybe CSExp
forall a b. (a -> b) -> a -> b
$
    case ArrayContents
vs of Imp.ArrayValues [PrimValue]
vs' ->
                 CSType -> Either Int [CSExp] -> CSExp
CreateArray (PrimType -> CSType
CS.compilePrimTypeToAST PrimType
t) (Either Int [CSExp] -> CSExp) -> Either Int [CSExp] -> CSExp
forall a b. (a -> b) -> a -> b
$ [CSExp] -> Either Int [CSExp]
forall a b. b -> Either a b
Right ([CSExp] -> Either Int [CSExp]) -> [CSExp] -> Either Int [CSExp]
forall a b. (a -> b) -> a -> b
$ (PrimValue -> CSExp) -> [PrimValue] -> [CSExp]
forall a b. (a -> b) -> [a] -> [b]
map PrimValue -> CSExp
CS.compilePrimValue [PrimValue]
vs'
               Imp.ArrayZeros Int
n ->
                 CSType -> Either Int [CSExp] -> CSExp
CreateArray (PrimType -> CSType
CS.compilePrimTypeToAST PrimType
t) (Either Int [CSExp] -> CSExp) -> Either Int [CSExp] -> CSExp
forall a b. (a -> b) -> a -> b
$ Int -> Either Int [CSExp]
forall a b. a -> Either a b
Left Int
n

  -- Create memory block on the device.
  String
ptr <- String -> CompilerM OpenCL () String
forall (f :: * -> *). MonadFreshNames f => String -> f String
newVName' String
"ptr"
  let num_elems :: Int
num_elems = case ArrayContents
vs of Imp.ArrayValues [PrimValue]
vs' -> [PrimValue] -> Int
forall (t :: * -> *) a. Foldable t => t a -> Int
length [PrimValue]
vs'
                             Imp.ArrayZeros Int
n -> Int
n
      size :: CSExp
size = Integer -> CSExp
Integer (Integer -> CSExp) -> Integer -> CSExp
forall a b. (a -> b) -> a -> b
$ Int -> Integer
forall a. Integral a => a -> Integer
toInteger Int
num_elems Integer -> Integer -> Integer
forall a. Num a => a -> a -> a
* PrimType -> Integer
forall a. Num a => PrimType -> a
Imp.primByteSize PrimType
t

  CSStmt -> CompilerM OpenCL () ()
forall op s. CSStmt -> CompilerM op s ()
CS.staticMemAlloc (CSStmt -> CompilerM OpenCL () ())
-> CSStmt -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ CSExp -> CSExp -> CSStmt
Reassign CSExp
name' (CSExp -> CSStmt) -> CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$
    String -> [CSExp] -> CSExp
CS.simpleCall String
"EmptyMemblock" [String -> CSExp
Var String
"Ctx.EMPTY_MEM_HANDLE"]
  String
errcode <- VName -> String
CS.compileName (VName -> String)
-> CompilerM OpenCL () VName -> CompilerM OpenCL () String
forall (f :: * -> *) a b. Functor f => (a -> b) -> f a -> f b
<$> String -> CompilerM OpenCL () VName
forall (m :: * -> *). MonadFreshNames m => String -> m VName
newVName String
"errCode"
  CSStmt -> CompilerM OpenCL () ()
forall op s. CSStmt -> CompilerM op s ()
CS.staticMemAlloc (CSStmt -> CompilerM OpenCL () ())
-> CSStmt -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ CSType -> CSExp -> Maybe CSExp -> CSStmt
AssignTyped CSType
computeErrCodeT (String -> CSExp
Var String
errcode) Maybe CSExp
forall a. Maybe a
Nothing
  CSStmt -> CompilerM OpenCL () ()
forall op s. CSStmt -> CompilerM op s ()
CS.staticMemAlloc (CSStmt -> CompilerM OpenCL () ())
-> CSStmt -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ CSExp -> CSExp -> CSStmt
Reassign CSExp
name' (CSExp -> CSStmt) -> CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$
    String -> [CSExp] -> CSExp
CS.simpleCall String
"MemblockAllocDevice"
    [CSExp -> CSExp
Ref (CSExp -> CSExp) -> CSExp -> CSExp
forall a b. (a -> b) -> a -> b
$ String -> CSExp
Var String
"Ctx", CSExp
name', CSExp
size, String -> CSExp
String (String -> CSExp) -> String -> CSExp
forall a b. (a -> b) -> a -> b
$ CSExp -> String
forall a. Pretty a => a -> String
pretty CSExp
name']

  -- Copy Numpy array to the device memory block.
  CSStmt -> CompilerM OpenCL () ()
forall op s. CSStmt -> CompilerM op s ()
CS.staticMemAlloc (CSStmt -> CompilerM OpenCL () ())
-> CSStmt -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ [CSStmt] -> CSStmt
Unsafe [
    CSExp -> CSExp -> [CSStmt] -> CSStmt
Fixed (String -> CSExp
Var String
ptr) (CSExp -> CSExp
Addr (CSExp -> CSExp) -> CSExp -> CSExp
forall a b. (a -> b) -> a -> b
$ CSExp -> CSIdx -> CSExp
Index (String -> CSExp
Var String
tmp_arr) (CSIdx -> CSExp) -> CSIdx -> CSExp
forall a b. (a -> b) -> a -> b
$ CSExp -> CSIdx
IdxExp (CSExp -> CSIdx) -> CSExp -> CSIdx
forall a b. (a -> b) -> a -> b
$ Integer -> CSExp
Integer Integer
0)
      [ CSExp -> CSStmt -> CSStmt
ifNotZeroSize CSExp
size (CSStmt -> CSStmt) -> CSStmt -> CSStmt
forall a b. (a -> b) -> a -> b
$
        CSExp -> CSStmt
Exp (CSExp -> CSStmt) -> CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$ String -> [CSExp] -> CSExp
CS.simpleCall String
"CL10.EnqueueWriteBuffer"
          [ String -> CSExp
Var String
"Ctx.OpenCL.Queue", CSExp -> CSExp
memblockFromMem CSExp
name', Bool -> CSExp
Bool Bool
True
          , CSExp -> CSExp
CS.toIntPtr (Integer -> CSExp
Integer Integer
0),CSExp -> CSExp
CS.toIntPtr CSExp
size
          , CSExp -> CSExp
CS.toIntPtr (CSExp -> CSExp) -> CSExp -> CSExp
forall a b. (a -> b) -> a -> b
$ String -> CSExp
Var String
ptr, Integer -> CSExp
Integer Integer
0, CSExp
Null, CSExp
Null ]
      ]
    ]

staticOpenCLArray VName
_ String
space PrimType
_ ArrayContents
_ =
  String -> CompilerM OpenCL () ()
forall a. HasCallStack => String -> a
error (String -> CompilerM OpenCL () ())
-> String -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ String
"CSOpenCL backend cannot create static array in memory space '" String -> String -> String
forall a. [a] -> [a] -> [a]
++ String
space String -> String -> String
forall a. [a] -> [a] -> [a]
++ String
"'"

memblockFromMem :: CSExp -> CSExp
memblockFromMem :: CSExp -> CSExp
memblockFromMem CSExp
mem = CSExp -> String -> CSExp
Field CSExp
mem String
"Mem"

packArrayOutput :: CS.EntryOutput Imp.OpenCL ()
packArrayOutput :: EntryOutput OpenCL ()
packArrayOutput CSExp
mem String
"device" PrimType
bt Signedness
ept [DimSize]
dims = do
  let size :: CSExp
size = (CSExp -> CSExp -> CSExp) -> CSExp -> [CSExp] -> CSExp
forall (t :: * -> *) a b.
Foldable t =>
(a -> b -> b) -> b -> t a -> b
foldr (String -> CSExp -> CSExp -> CSExp
BinOp String
"*") (Integer -> CSExp
Integer Integer
1) [CSExp]
dims'
  let bt' :: CSType
bt' = PrimType -> Signedness -> CSType
CS.compilePrimTypeToASText PrimType
bt Signedness
ept
  let nbytes :: CSExp
nbytes = String -> CSExp -> CSExp -> CSExp
BinOp String
"*" (CSType -> CSExp
CS.sizeOf CSType
bt') CSExp
size
  let createTuple :: String
createTuple = String
"createTuple_"String -> String -> String
forall a. [a] -> [a] -> [a]
++ CSType -> String
forall a. Pretty a => a -> String
pretty CSType
bt'

  CSExp -> CompilerM OpenCL () CSExp
forall (m :: * -> *) a. Monad m => a -> m a
return (CSExp -> CompilerM OpenCL () CSExp)
-> CSExp -> CompilerM OpenCL () CSExp
forall a b. (a -> b) -> a -> b
$ String -> [CSExp] -> CSExp
CS.simpleCall String
createTuple [ CSExp -> CSExp
memblockFromMem CSExp
mem, String -> CSExp
Var String
"Ctx.OpenCL.Queue", CSExp
nbytes
                                     , CSType -> Either Int [CSExp] -> CSExp
CreateArray (CSPrim -> CSType
Primitive (CSPrim -> CSType) -> CSPrim -> CSType
forall a b. (a -> b) -> a -> b
$ CSInt -> CSPrim
CSInt CSInt
Int64T) (Either Int [CSExp] -> CSExp) -> Either Int [CSExp] -> CSExp
forall a b. (a -> b) -> a -> b
$ [CSExp] -> Either Int [CSExp]
forall a b. b -> Either a b
Right [CSExp]
dims']
  where dims' :: [CSExp]
dims' = (DimSize -> CSExp) -> [DimSize] -> [CSExp]
forall a b. (a -> b) -> [a] -> [b]
map DimSize -> CSExp
CS.compileDim [DimSize]
dims

packArrayOutput CSExp
_ String
sid PrimType
_ Signedness
_ [DimSize]
_ =
  String -> CompilerM OpenCL () CSExp
forall a. HasCallStack => String -> a
error (String -> CompilerM OpenCL () CSExp)
-> String -> CompilerM OpenCL () CSExp
forall a b. (a -> b) -> a -> b
$ String
"Cannot return array from " String -> String -> String
forall a. [a] -> [a] -> [a]
++ String
sid String -> String -> String
forall a. [a] -> [a] -> [a]
++ String
" space."

unpackArrayInput :: CS.EntryInput Imp.OpenCL ()
unpackArrayInput :: EntryInput OpenCL ()
unpackArrayInput VName
mem String
"device" PrimType
t Signedness
_ [DimSize]
dims CSExp
e = do
  let size :: CSExp
size = (CSExp -> CSExp -> CSExp) -> CSExp -> [CSExp] -> CSExp
forall (t :: * -> *) a b.
Foldable t =>
(a -> b -> b) -> b -> t a -> b
foldr (String -> CSExp -> CSExp -> CSExp
BinOp String
"*") (Integer -> CSExp
Integer Integer
1) [CSExp]
dims'
  let t' :: CSType
t' = PrimType -> CSType
CS.compilePrimTypeToAST PrimType
t
  let nbytes :: CSExp
nbytes = String -> CSExp -> CSExp -> CSExp
BinOp String
"*" (CSType -> CSExp
CS.sizeOf CSType
t') CSExp
size
  (DimSize -> Int32 -> CompilerM OpenCL () ())
-> [DimSize] -> [Int32] -> CompilerM OpenCL () ()
forall (m :: * -> *) a b c.
Applicative m =>
(a -> b -> m c) -> [a] -> [b] -> m ()
zipWithM_ (CSExp -> DimSize -> Int32 -> CompilerM OpenCL () ()
forall op s. CSExp -> DimSize -> Int32 -> CompilerM op s ()
CS.unpackDim CSExp
e) [DimSize]
dims [Int32
0..]
  String
ptr <- VName -> String
forall a. Pretty a => a -> String
pretty (VName -> String)
-> CompilerM OpenCL () VName -> CompilerM OpenCL () String
forall (f :: * -> *) a b. Functor f => (a -> b) -> f a -> f b
<$> String -> CompilerM OpenCL () VName
forall (m :: * -> *). MonadFreshNames m => String -> m VName
newVName String
"ptr"

  CSExp
mem' <- VName -> CompilerM OpenCL () CSExp
forall op s. VName -> CompilerM op s CSExp
CS.compileVar VName
mem
  CSStmt -> CompilerM OpenCL () ()
forall op s. CSStmt -> CompilerM op s ()
CS.stm (CSStmt -> CompilerM OpenCL () ())
-> CSStmt -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ Param -> CSStmt
CS.getDefaultDecl (VName -> Space -> Param
Imp.MemParam VName
mem (String -> Space
Imp.Space String
"device"))
  Allocate OpenCL ()
allocateOpenCLBuffer CSExp
mem' CSExp
nbytes String
"device"
  CSStmt -> CompilerM OpenCL () ()
forall op s. CSStmt -> CompilerM op s ()
CS.stm (CSStmt -> CompilerM OpenCL () ())
-> CSStmt -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ [CSStmt] -> CSStmt
Unsafe [CSExp -> CSExp -> [CSStmt] -> CSStmt
Fixed (String -> CSExp
Var String
ptr) (CSExp -> CSExp
Addr (CSExp -> CSExp) -> CSExp -> CSExp
forall a b. (a -> b) -> a -> b
$ CSExp -> CSIdx -> CSExp
Index (CSExp -> String -> CSExp
Field CSExp
e String
"Item1") (CSIdx -> CSExp) -> CSIdx -> CSExp
forall a b. (a -> b) -> a -> b
$ CSExp -> CSIdx
IdxExp (CSExp -> CSIdx) -> CSExp -> CSIdx
forall a b. (a -> b) -> a -> b
$ Integer -> CSExp
Integer Integer
0)
      [ CSExp -> CSStmt -> CSStmt
ifNotZeroSize CSExp
nbytes (CSStmt -> CSStmt) -> CSStmt -> CSStmt
forall a b. (a -> b) -> a -> b
$
        CSExp -> CSStmt
Exp (CSExp -> CSStmt) -> CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$ String -> [CSExp] -> CSExp
CS.simpleCall String
"CL10.EnqueueWriteBuffer"
        [ String -> CSExp
Var String
"Ctx.OpenCL.Queue", CSExp -> CSExp
memblockFromMem CSExp
mem', Bool -> CSExp
Bool Bool
True
        , CSExp -> CSExp
CS.toIntPtr (Integer -> CSExp
Integer Integer
0), CSExp -> CSExp
CS.toIntPtr CSExp
nbytes, CSExp -> CSExp
CS.toIntPtr (String -> CSExp
Var String
ptr)
        , Integer -> CSExp
Integer Integer
0, CSExp
Null, CSExp
Null]
      ]]

  where dims' :: [CSExp]
dims' = (DimSize -> CSExp) -> [DimSize] -> [CSExp]
forall a b. (a -> b) -> [a] -> [b]
map DimSize -> CSExp
CS.compileDim [DimSize]
dims

unpackArrayInput VName
_ String
sid PrimType
_ Signedness
_ [DimSize]
_ CSExp
_ =
  String -> CompilerM OpenCL () ()
forall a. HasCallStack => String -> a
error (String -> CompilerM OpenCL () ())
-> String -> CompilerM OpenCL () ()
forall a b. (a -> b) -> a -> b
$ String
"Cannot accept array from " String -> String -> String
forall a. [a] -> [a] -> [a]
++ String
sid String -> String -> String
forall a. [a] -> [a] -> [a]
++ String
" space."

futharkSyncContext :: CSStmt
futharkSyncContext :: CSStmt
futharkSyncContext = CSExp -> CSStmt
Exp (CSExp -> CSStmt) -> CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$ String -> [CSExp] -> CSExp
CS.simpleCall String
"FutharkContextSync" []

ifNotZeroSize :: CSExp -> CSStmt -> CSStmt
ifNotZeroSize :: CSExp -> CSStmt -> CSStmt
ifNotZeroSize CSExp
e CSStmt
s =
  CSExp -> [CSStmt] -> [CSStmt] -> CSStmt
If (String -> CSExp -> CSExp -> CSExp
BinOp String
"!=" CSExp
e (Integer -> CSExp
Integer Integer
0)) [CSStmt
s] []

finishIfSynchronous :: CS.CompilerM op s ()
finishIfSynchronous :: CompilerM op s ()
finishIfSynchronous =
  CSStmt -> CompilerM op s ()
forall op s. CSStmt -> CompilerM op s ()
CS.stm (CSStmt -> CompilerM op s ()) -> CSStmt -> CompilerM op s ()
forall a b. (a -> b) -> a -> b
$ CSExp -> [CSStmt] -> [CSStmt] -> CSStmt
If (String -> CSExp
Var String
"Synchronous") [CSExp -> CSStmt
Exp (CSExp -> CSStmt) -> CSExp -> CSStmt
forall a b. (a -> b) -> a -> b
$ String -> [CSExp] -> CSExp
CS.simpleCall String
"CL10.Finish" [String -> CSExp
Var String
"Ctx.OpenCL.Queue"]] []

newVName' :: MonadFreshNames f => String -> f String
newVName' :: String -> f String
newVName' String
s = VName -> String
CS.compileName (VName -> String) -> f VName -> f String
forall (f :: * -> *) a b. Functor f => (a -> b) -> f a -> f b
<$> String -> f VName
forall (m :: * -> *). MonadFreshNames m => String -> m VName
newVName String
s