Home My Page Projects Code Snippets Project Openings diderot
Summary Activity Tracker Tasks SCM

SCM Repository

[diderot] Diff of /branches/pure-cfg/src/compiler/cl-target/cl-target.sml
ViewVC logotype

Diff of /branches/pure-cfg/src/compiler/cl-target/cl-target.sml

Parent Directory Parent Directory | Revision Log Revision Log | View Patch Patch

revision 1273, Mon Jun 6 10:46:20 2011 UTC revision 1322, Sun Jun 12 17:13:33 2011 UTC
# Line 1  Line 1 
1  (* c-target.sml  (* cl-target.sml
2   *   *
3   * COPYRIGHT (c) 2011 The Diderot Project (http://diderot-language.cs.uchicago.edu)   * COPYRIGHT (c) 2011 The Diderot Project (http://diderot-language.cs.uchicago.edu)
4   * All rights reserved.   * All rights reserved.
# Line 13  Line 13 
13      structure CL = CLang      structure CL = CLang
14      structure RN = RuntimeNames      structure RN = RuntimeNames
15      structure ToCL = TreeToCL      structure ToCL = TreeToCL
16        structure N = CNames
17    
18      type var = ToCL.var    (* C variable translation *)
19        structure TrCVar =
20          struct
21            type env = CL.typed_var TreeIL.Var.Map.map
22            fun lookup (env, x) = (case V.Map.find (env, x)
23                   of SOME(CL.V(_, x')) => x'
24                    | NONE => raise Fail(concat["TrCVar.lookup(_, ", V.name x, ")"])
25                  (* end case *))
26          (* translate a variable that occurs in an l-value context (i.e., as the target of an assignment) *)
27            fun lvalueVar (env, x) = (case V.kind x
28                   of IL.VK_Global => CL.mkIndirect(CL.mkVar RN.globalsVarName, lookup(env, x))
29                    | IL.VK_State strand => CL.mkIndirect(CL.mkVar "selfOut", lookup(env, x))
30                    | IL.VK_Local => CL.mkVar(lookup(env, x))
31                  (* end case *))
32          (* translate a variable that occurs in an r-value context *)
33            fun rvalueVar (env, x) = (case V.kind x
34                   of IL.VK_Global => CL.mkIndirect(CL.mkVar RN.globalsVarName, lookup(env, x))
35                    | IL.VK_State strand => CL.mkIndirect(CL.mkVar "selfIn", lookup(env, x))
36                    | IL.VK_Local => CL.mkVar(lookup(env, x))
37                  (* end case *))
38          end
39    
40        structure ToC = TreeToCFn (TrCVar)
41    
42        type var = CL.typed_var
43      type exp = CL.exp      type exp = CL.exp
44      type stm = CL.stm      type stm = CL.stm
45    
46      (* OpenCL specific types *)
47        val clIntTy = CL.T_Named "cl_int"
48        val clProgramTy = CL.T_Named "cl_program"
49        val clKernelTy  = CL.T_Named "cl_kernel"
50        val clCmdQueueTy = CL.T_Named "cl_command_queue"
51        val clContextTy = CL.T_Named "cl_context"
52        val clDeviceIdTy = CL.T_Named "cl_device_id"
53        val clPlatformIdTy = CL.T_Named "cl_platform_id"
54        val clMemoryTy = CL.T_Named "cl_mem"
55        val globPtrTy = CL.T_Ptr(CL.T_Named RN.globalsTy)
56    
57      (* variable or field that is mirrored between host and GPU *)
58        type mirror_var = {
59                hostTy : CL.ty,             (* variable type on Host (i.e., C type) *)
60                gpuTy : CL.ty,              (* variable's type on GPU (i.e., OpenCL type) *)
61                var : CL.var                (* variable name *)
62              }
63    
64      datatype strand = Strand of {      datatype strand = Strand of {
65          name : string,          name : string,
66          tyName : string,          tyName : string,
67          state : var list ref,          state : mirror_var list ref,
68          output : (Ty.ty * CL.var) option ref,   (* the strand's output variable (only one for now) *)          output : (Ty.ty * CL.var) option ref,   (* the strand's output variable (only one for now) *)
69          code : CL.decl list ref,          code : CL.decl list ref,
70          init_code: CL.decl ref          init_code: CL.decl ref
71        }        }
72    
73      datatype program = Prog of {      datatype program = Prog of {
74            name : string,                  (* stem of source file *)
75          double : bool,                  (* true for double-precision support *)          double : bool,                  (* true for double-precision support *)
76          parallel : bool,                (* true for multithreaded (or multi-GPU) target *)          parallel : bool,                (* true for multithreaded (or multi-GPU) target *)
77          debug : bool,                   (* true for debug support in executable *)          debug : bool,                   (* true for debug support in executable *)
78          globals : CL.decl list ref,          globals : mirror_var list ref,
79          topDecls : CL.decl list ref,          topDecls : CL.decl list ref,
80          strands : strand AtomTable.hash_table,          strands : strand AtomTable.hash_table,
81          initially : CL.stm list ref,          initially :  CL.decl ref,
82          numDims: int ref,          numDims: int ref,
83          imgGlobals: (string * int) list ref,          imgGlobals: (string * int) list ref,
84          prFn: CL.decl ref          prFn: CL.decl ref
# Line 69  Line 113 
113      structure Tr =      structure Tr =
114        struct        struct
115          fun fragment (ENV{info, vMap, scope}, blk) = let          fun fragment (ENV{info, vMap, scope}, blk) = let
116                val (vMap, stms) = ToCL.trFragment (vMap, blk)                val (vMap, stms) = (case scope
117                         of GlobalScope => ToC.trFragment (vMap, blk)
118    (* NOTE: if we move strand initialization to the GPU, then we'll have to change the following code! *)
119                          | InitiallyScope => ToC.trFragment (vMap, blk)
120                          | _ => ToCL.trFragment (vMap, blk)
121                        (* end case *))
122                in                in
123                  (ENV{info=info, vMap=vMap, scope=scope}, stms)                  (ENV{info=info, vMap=vMap, scope=scope}, stms)
124                end                end
125          fun saveState cxt stateVars (env, args, stm) = (          fun block (ENV{vMap, scope, ...}, blk) = let
126                  fun saveState cxt stateVars trAssign (env, args, stm) = (
127                ListPair.foldrEq                ListPair.foldrEq
128                  (fn (x, e, stms) => ToCL.trAssign(env, x, e)@stms)                        (fn (x, e, stms) => trAssign(env, x, e)@stms)
129                    [stm]                    [stm]
130                      (stateVars, args)                      (stateVars, args)
131                ) handle ListPair.UnequalLengths => (                ) handle ListPair.UnequalLengths => (
132                  print(concat["saveState ", cxt, ": length mismatch; ", Int.toString(List.length args), " args\n"]);                  print(concat["saveState ", cxt, ": length mismatch; ", Int.toString(List.length args), " args\n"]);
133                  raise Fail(concat["saveState ", cxt, ": length mismatch"]))                  raise Fail(concat["saveState ", cxt, ": length mismatch"]))
134          fun block (ENV{vMap, scope, ...}, blk) = (case scope                in
135                 of StrandScope stateVars => ToCL.trBlock (vMap, saveState "StrandScope" stateVars, blk)                  case scope
136                  | MethodScope stateVars => ToCL.trBlock (vMap, saveState "MethodScope" stateVars, blk)  (* NOTE: if we move strand initialization to the GPU, then we'll have to change the following code! *)
137                  | _ => ToCL.trBlock (vMap, fn (_, _, stm) => [stm], blk)                   of StrandScope stateVars =>
138                (* end case *))                        ToCL.trBlock (vMap, saveState "StrandScope" stateVars ToC.trAssign, blk)
139                      | MethodScope stateVars =>
140                          ToCL.trBlock (vMap, saveState "MethodScope" stateVars ToCL.trAssign, blk)
141                      | InitiallyScope => ToCL.trBlock (vMap, fn (_, _, stm) => [stm], blk)
142                      | _ => ToC.trBlock (vMap, fn (_, _, stm) => [stm], blk)
143                    (* end case *)
144                  end
145          fun exp (ENV{vMap, ...}, e) = ToCL.trExp(vMap, e)          fun exp (ENV{vMap, ...}, e) = ToCL.trExp(vMap, e)
146        end        end
147    
# Line 94  Line 150 
150        struct        struct
151          fun name (ToCL.V(_, name)) = name          fun name (ToCL.V(_, name)) = name
152          fun global (Prog{globals,imgGlobals, ...}, name, ty) = let          fun global (Prog{globals,imgGlobals, ...}, name, ty) = let
153                val ty' = ToCL.trType ty                val x = {hostTy = ToC.trType ty, gpuTy = ToCL.trType ty, var = name}
154                fun isImgGlobal (imgGlobals, Ty.ImageTy(ImageInfo.ImgInfo{dim, ...}), name) =  imgGlobals  := (name,dim):: !imgGlobals                fun isImgGlobal (Ty.ImageTy(ImageInfo.ImgInfo{dim, ...}), name) =
155                  | isImgGlobal (imgGlobals, _, _) =  ()                      imgGlobals  := (name,dim) :: !imgGlobals
156                in                  | isImgGlobal _ =  ()
157                  globals := CL.D_Var([], ty', name, NONE) :: !globals;                in
158                  isImgGlobal(imgGlobals,ty,name);                  globals := x :: !globals;
159               ToCL.V(ty', name)                  isImgGlobal (ty, name);
160                    ToCL.V(#gpuTy x, name)
161                end                end
162          fun param x = ToCL.V(ToCL.trType(V.ty x), V.name x)          fun param x = ToCL.V(ToCL.trType(V.ty x), V.name x)
163          fun state (Strand{state, ...}, x) = let          fun state (Strand{state, ...}, x) = let
164                val ty' = ToCL.trType(V.ty x)                val ty = V.ty x
165                val x' = ToCL.V(ty', V.name x)                val x' = {hostTy = ToC.trType ty, gpuTy = ToCL.trType ty, var = V.name x}
166                in                in
167                  state := x' :: !state;                  state := x' :: !state;
168                  x'                  ToCL.V(#gpuTy x', #var x')
169                end                end
170        end        end
171    
# Line 138  Line 195 
195    (* programs *)    (* programs *)
196      structure Program =      structure Program =
197        struct        struct
198          fun new {double, parallel, debug} = (          fun new {name, double, parallel, debug} = (
199                RN.initTargetSpec double;                RN.initTargetSpec double;
200                  CNames.initTargetSpec double;
201                Prog{                Prog{
202                      name = name,
203                    double = double, parallel = parallel, debug = debug,                    double = double, parallel = parallel, debug = debug,
204                    globals = ref [],                    globals = ref [],
205                    topDecls = ref [],                    topDecls = ref [],
206                    strands = AtomTable.mkTable (16, Fail "strand table"),                    strands = AtomTable.mkTable (16, Fail "strand table"),
207                    initially = ref([CL.S_Comment["missing initially"]]),                    initially = ref(CL.D_Comment["missing initially"]),
208                                    numDims = ref(0),                                    numDims = ref(0),
209                                    imgGlobals = ref[],                                    imgGlobals = ref[],
210                                    prFn = ref(CL.D_Comment(["No Print Function"]))                                    prFn = ref(CL.D_Comment(["No Print Function"]))
211                  })                  })
212        (* register the global initialization part of a program *)        (* register the global initialization part of a program *)
213    (* FIXME: unused code; can this be removed??
214            fun globalIndirects (globals,stms) = let            fun globalIndirects (globals,stms) = let
215                  fun getGlobals(CL.D_Var(_,_,globalVar,_)::rest) = CL.mkAssign(CL.mkIndirect(CL.E_Var RN.globalsVarName,globalVar),CL.E_Var globalVar)::getGlobals(rest)                  fun getGlobals ({name,target as TargetUtil.TARGET_CL}::rest) =
216                    | getGlobals([]) = []                        CL.mkAssign(CL.mkIndirect(CL.mkVar RN.globalsVarName,name),CL.mkVar name)
217                    | getGlobals(_::rest) = getGlobals(rest)                          ::getGlobals rest
218                      | getGlobals [] = []
219                      | getGlobals (_::rest) = getGlobals rest
220                  in                  in
221                    stms @ getGlobals(globals)                    stms @ getGlobals globals
222                  end                  end
223    *)
224        (* register the code that is used to register command-line options for input variables *)        (* register the code that is used to register command-line options for input variables *)
225          fun inputs (Prog{topDecls, ...}, stm) = let          fun inputs (Prog{topDecls, ...}, stm) = let
226                val inputsFn = CL.D_Func(                val inputsFn = CL.D_Func(
# Line 169  Line 231 
231                  topDecls := inputsFn :: !topDecls                  topDecls := inputsFn :: !topDecls
232                end                end
233    
234          fun init (Prog{globals,topDecls,...}, CL.S_Block(init)) = let        (* register the global initialization part of a program *)
235                val params = [          fun init (Prog{topDecls, ...}, init) = let
236                        CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName)                                    val globalsDecl = CL.mkAssign(CL.E_Var RN.globalsVarName,
237                      ]                      CL.mkApply("malloc", [CL.mkApply("sizeof",[CL.mkVar RN.globalsTy])]))
238                val body = CL.S_Block(globalIndirects(!globals,init))  
239                val initFn = CL.D_Func([], CL.voidTy, RN.initGlobals, params, body)                val initGlobalsCall = CL.mkCall(RN.initGlobalsHelper,[])
240                in  
241                  topDecls := initFn :: !topDecls                                          val initFn = CL.D_Func(
242                end                      [], CL.voidTy, RN.initGlobals, [],
243            | init (Prog{globals,topDecls,...}, init) = let                      CL.mkBlock([globalsDecl,initGlobalsCall]))
244                val params = [                val initFn_helper = CL.D_Func(
245                        CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName)                      [], CL.voidTy, RN.initGlobalsHelper, [],
246                      ]                      init)
247                val initFn = CL.D_Func([], CL.voidTy, RN.initGlobals, params, init)                val shutdownFn = CL.D_Func(
248                        [], CL.voidTy, RN.shutdown,
249                        [CL.PARAM([], CL.T_Ptr(CL.T_Named RN.worldTy), "wrld")],
250                        CL.S_Block[])
251                in                in
252                  topDecls := initFn :: !topDecls                  topDecls := shutdownFn :: initFn :: initFn_helper :: !topDecls
253                end                end
   
254        (* create and register the initially function for a program *)        (* create and register the initially function for a program *)
255          fun initially {          fun initially {
256                prog = Prog{strands, initially,numDims,...},                prog = Prog{name=progName, strands, initially, ...},
257                isArray : bool,                isArray : bool,
258                iterPrefix : stm list,                iterPrefix : stm list,
259                iters : (var * exp * exp) list,                iters : (var * exp * exp) list,
# Line 199  Line 263 
263              } = let              } = let
264                val name = Atom.toString strand                val name = Atom.toString strand
265                val nDims = List.length iters                val nDims = List.length iters
266                  val worldTy = CL.T_Ptr(CL.T_Named N.worldTy)
267                fun mapi f xs = let                fun mapi f xs = let
268                      fun mapf (_, []) = []                      fun mapf (_, []) = []
269                        | mapf (i, x::xs) = f(i, x) :: mapf(i+1, xs)                        | mapf (i, x::xs) = f(i, x) :: mapf(i+1, xs)
# Line 207  Line 272 
272                      end                      end
273                val baseInit = mapi (fn (i, (_, e, _)) => (i, CL.I_Exp e)) iters                val baseInit = mapi (fn (i, (_, e, _)) => (i, CL.I_Exp e)) iters
274                val sizeInit = mapi                val sizeInit = mapi
275                      (fn (i, (ToCL.V(ty, _), lo, hi)) =>                      (fn (i, (CL.V(ty, _), lo, hi)) =>
276                          (i, CL.I_Exp(CL.mkBinOp(CL.mkBinOp(hi, CL.#-, lo), CL.#+, CL.E_Int(1, ty))))                          (i, CL.I_Exp(CL.mkBinOp(CL.mkBinOp(hi, CL.#-, lo), CL.#+, CL.E_Int(1, ty))))
277                      ) iters                      ) iters
278                    val numStrandsVar = "numStrandsVar"              (* code to allocate the world and initial strands *)
279                val allocCode = iterPrefix @ [                val wrld = "wrld"
280                  val allocCode = [
281                        CL.mkComment["allocate initial block of strands"],                        CL.mkComment["allocate initial block of strands"],
282                        CL.mkDecl(CL.T_Array(CL.int32, SOME nDims), "base", SOME(CL.I_Array baseInit)),                        CL.mkDecl(CL.T_Array(CL.int32, SOME nDims), "base", SOME(CL.I_Array baseInit)),
283                        CL.mkDecl(CL.T_Array(CL.uint32, SOME nDims), "size", SOME(CL.I_Array sizeInit)),                        CL.mkDecl(CL.T_Array(CL.uint32, SOME nDims), "size", SOME(CL.I_Array sizeInit)),
284                        CL.mkDecl(CL.int32,"numDims",SOME(CL.I_Exp(CL.E_Int(IntInf.fromInt nDims, CL.int32))))                        CL.mkDecl(worldTy, wrld,
285                      ]                          SOME(CL.I_Exp(CL.E_Apply(N.allocInitially, [
286                val numStrandsLoopBody = CL.mkExpStm(CL.mkAssignOp(CL.E_Var numStrandsVar, CL.*=,CL.mkSubscript(CL.E_Var "size",CL.E_Var "i")))                              CL.mkVar "ProgramName",
287                val numStrandsLoop =  CL.mkFor([(CL.intTy, "i", CL.E_Int(0,CL.intTy))],                              CL.mkUnOp(CL.%&, CL.E_Var(N.strandDesc name)),
288                      CL.mkBinOp(CL.E_Var "i", CL.#<, CL.E_Var "numDims"),                              CL.E_Bool isArray,
289                      [CL.mkPostOp(CL.E_Var "i", CL.^++)], numStrandsLoopBody)                              CL.E_Int(IntInf.fromInt nDims, CL.int32),
290                in                              CL.E_Var "base",
291                  numDims := nDims;                              CL.E_Var "size"
292                  initially := allocCode @ [numStrandsLoop]                            ]))))
               end  
   
   
       (***** OUTPUT *****)  
         fun genStrandInit(Strand{name,tyName,state,output,code,...},nDims) = let  
                 val params = [  
                           CL.PARAM([], CL.T_Ptr(CL.uint32), "sizes" ),  
                           CL.PARAM([], CL.intTy, "width"),  
                           CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "strands")  
293                          ]                          ]
294                (* create the loop nest for the initially iterations
295                  val body = let                val indexVar = "ix"
296                              fun loopParams (3) =                val strandTy = CL.T_Ptr(CL.T_Named(N.strandTy name))
297                                   "x"::"y"::"k"::[]                fun mkLoopNest [] = CL.mkBlock(createPrefix @ [
298                                | loopParams (2) =                        CL.mkDecl(strandTy, "sp",
299                                   "x"::"y"::[]                          SOME(CL.I_Exp(
300                                | loopParams (1) =                            CL.E_Cast(strandTy,
301                                   "x"::[]                            CL.E_Apply(N.inState, [CL.E_Var "wrld", CL.E_Var indexVar]))))),
302                                | loopParams (_) =                        CL.mkCall(N.strandInit name,
303                                  raise Fail("genStrandInit: missing size dim")                          CL.E_Var RN.globalsVarName :: CL.E_Var "sp" :: args),
304                             fun mkLoopNest ([],_,nDims) =  if nDims = 1 then                        CL.mkAssign(CL.E_Var indexVar, CL.mkBinOp(CL.E_Var indexVar, CL.#+, CL.E_Int(1, CL.uint32)))
305                        ])
306                                          CL.mkBlock ([CL.mkCall(RN.strandInit name, [CL.E_UnOp(CL.%&,CL.mkSubscript(CL.E_Var "strands",CL.E_Str "x")),                  | mkLoopNest ((CL.V(ty, param), lo, hi)::iters) = let
307                                                  CL.E_Var "x"])])                      val body = mkLoopNest iters
                                         else let  
                                                 val index = CL.mkBinOp(CL.mkBinOp(CL.E_Var "x",CL.#*,CL.E_Var "width"),CL.#+,CL.E_Var "y")  
                                         in  
                                                 CL.mkBlock([CL.mkCall(RN.strandInit name, [CL.E_UnOp(CL.%&,CL.mkSubscript(CL.E_Var "strands",index)),  
                                                 CL.E_Var "x", CL.E_Var"y"])])  
                                         end  
   
                                 | mkLoopNest (param::rest,count,nDims) = let  
                                         val body = mkLoopNest (rest, count + 1,nDims)  
308                                     in                                     in
309                                          CL.mkFor(                                          CL.mkFor(
310                                                          [(CL.intTy, param, CL.E_Int(0,CL.intTy))],                          [(ty, param, lo)],
311                                                  CL.mkBinOp(CL.E_Var param, CL.#<, CL.mkSubscript(CL.E_Var "sizes",CL.E_Int(count,CL.intTy))),                          CL.mkBinOp(CL.E_Var param, CL.#<=, hi),
312                                                  [CL.mkPostOp(CL.E_Var param, CL.^++)],                                                  [CL.mkPostOp(CL.E_Var param, CL.^++)],
313                                                  body)                                                  body)
314                                     end                                     end
315                  val iterCode = [
316                          CL.mkComment["initially"],
317                          CL.mkDecl(CL.uint32, indexVar, SOME(CL.I_Exp(CL.E_Int(0, CL.uint32)))),
318                          mkLoopNest iters
319                        ] *)
320                  val body = CL.mkBlock(
321                        iterPrefix @
322                        allocCode @
323                        [CL.mkReturn(SOME(CL.E_Var "wrld"))])
324                  val initFn = CL.D_Func([], worldTy, N.initially, [], body)
325                          in                          in
326                                  [mkLoopNest ((loopParams nDims),0,nDims)]                  initially := initFn
                         end  
                 in  
                         CL.D_Func(["static"], CL.voidTy, RN.strandInitSetup, params,CL.mkBlock(body))  
327                  end                  end
328          fun genStrandPrint (Strand{name, tyName, state, output, code,...},nDims) = let  
329          (***** OUTPUT *****)
330            fun genStrandPrint (Strand{name, tyName, state, output, code,...}) = let
331              (* the print function *)              (* the print function *)
332                val prFnName = concat[name, "_print"]                val prFnName = concat[name, "_print"]
   
333                val prFn = let                val prFn = let
334                      val params = [                      val params = [
335                            CL.PARAM([], CL.T_Ptr(CL.T_Named "FILE"), "outS"),                            CL.PARAM([], CL.T_Ptr(CL.T_Named "FILE"), "outS"),
                           CL.PARAM([], CL.T_Ptr(CL.uint32), "sizes" ),  
                           CL.PARAM([], CL.intTy, "width"),  
336                            CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "self")                            CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "self")
337                          ]                          ]
   
338                     val SOME(ty, x) = !output                     val SOME(ty, x) = !output
339                     val outState = if nDims = 1 then                      val outState = CL.mkIndirect(CL.mkVar "self", x)
                           CL.mkSelect(CL.mkSubscript(CL.mkVar "self",CL.E_Var "x"), x)  
                         else if nDims = 2 then  
                                 CL.mkSelect(CL.mkSubscript(CL.mkVar "self",  
                                    CL.mkBinOp(CL.mkBinOp(CL.E_Var "x",CL.#*,CL.E_Var "width"),CL.#+,CL.E_Var "y")), x)  
   
                         else CL.mkSelect(CL.mkVar "self",x)  
   
340                      val prArgs = (case ty                      val prArgs = (case ty
341                             of Ty.IVecTy 1 => [CL.E_Str(!RN.gIntFormat ^ "\n"), outState]                             of Ty.IVecTy 1 => [CL.E_Str(!N.gIntFormat ^ "\n"), outState]
342                              | Ty.IVecTy d => let                              | Ty.IVecTy d => let
343                                  val fmt = CL.E_Str(                                  val fmt = CL.mkStr(
344                                        String.concatWith " " (List.tabulate(d, fn _ => !RN.gIntFormat))                                        String.concatWith " " (List.tabulate(d, fn _ => !N.gIntFormat))
345                                        ^ "\n")                                        ^ "\n")
346                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))                                  val args = List.tabulate (d, fn i => ToC.ivecIndex(outState, d, i))
347                                  in                                  in
348                                    fmt :: args                                    fmt :: args
349                                  end                                  end
350                              | Ty.TensorTy[] => [CL.E_Str "%f\n", outState]                              | Ty.TensorTy[] => [CL.mkStr "%f\n", outState]
351                              | Ty.TensorTy[d] => let                              | Ty.TensorTy[d] => let
352                                  val fmt = CL.E_Str(                                  val fmt = CL.mkStr(
353                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))
354                                        ^ "\n")                                        ^ "\n")
355                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))                                  val args = List.tabulate (d, fn i => ToC.vecIndex(outState, d, i))
356                                  in                                  in
357                                    fmt :: args                                    fmt :: args
358                                  end                                  end
359                              | _ => raise Fail("genStrand: unsupported output type " ^ Ty.toString ty)                              | _ => raise Fail("genStrand: unsupported output type " ^ Ty.toString ty)
360                            (* end case *))                            (* end case *))
   
                           val body = let  
   
                             fun loopParams (3) =  
                                  "x"::"y"::"k"::[]  
                               | loopParams (2) =  
                                  "x"::"y"::[]  
                               | loopParams (1) =  
                                  "x"::[]  
                               | loopParams (_) =  
                                 raise Fail("genStrandPrint: unsupported output type " ^ Ty.toString ty)  
   
                            fun mkLoopNest ([],_) =  
                                                 CL.mkCall("fprintf", CL.mkVar "outS" :: prArgs)  
                                 | mkLoopNest (param::rest,count) = let  
                                         val body = mkLoopNest (rest, count + 1)  
                                    in  
                                                 CL.mkFor(  
                                                         [(CL.intTy, param, CL.E_Int(0,CL.intTy))],  
                                                 CL.mkBinOp(CL.E_Var param, CL.#<, CL.mkSubscript(CL.E_Var "sizes",CL.E_Int(count,CL.intTy))),  
                                                 [CL.mkPostOp(CL.E_Var param, CL.^++)],  
                                                 body)  
                                    end  
361                          in                          in
362                                  [mkLoopNest ((loopParams nDims),0)]                        CL.D_Func(["static"], CL.voidTy, prFnName, params,
363                          end                          CL.mkCall("fprintf", CL.mkVar "outS" :: prArgs))
   
                     in  
                       CL.D_Func(["static"], CL.voidTy, prFnName, params,CL.mkBlock(body))  
364                      end                      end
365                in                in
366                                   prFn                                   prFn
367                end                end
368          fun genStrandTyDef (Strand{tyName, state,...}) =  
369            fun genStrandTyDef (targetTy, Strand{tyName, state,...}) =
370              (* the type declaration for the strand's state struct *)              (* the type declaration for the strand's state struct *)
371                CL.D_StructDef(                CL.D_StructDef(
372                        List.rev (List.map (fn ToCL.V(ty, x) => (ty, x)) (!state)),                  List.rev (List.map (fn x => (targetTy x, #var x)) (!state)),
373                        tyName)                        tyName)
374    
   
375          (* generates the load kernel function *)          (* generates the load kernel function *)
 (* FIXME: this code might be part of the runtime system *)  
         fun genKernelLoader() =  
                 CL.D_Verbatim ( ["/* Loads the Kernel from a file */",  
                                                 "char * loadKernel (const char * filename) {",  
                                                 "struct stat statbuf;",  
                                                 "FILE *fh;",  
                                                 "char *source;",  
                                                 "fh = fopen(filename, \"r\");",  
                                                 "if (fh == 0)",  
                                                 "   return 0;",  
                                                 "stat(filename, &statbuf);",  
                                                 "source = (char *) malloc(statbuf.st_size + 1);",  
                                                 "fread(source, statbuf.st_size, 1, fh);",  
                                                 "fread(source, statbuf.st_size, 1, fh);",  
                                                 "return source;",  
                                                 "}"])  
 (* generates the opencl buffers for the image data *)  
         fun getGlobalDataBuffers(globals,count,contextVar,errVar) = let  
                 val globalBufferDecl =  CL.mkDecl(CL.clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)  
                 val globalBuffer = CL.mkAssign(CL.E_Var(concat[RN.globalsVarName,"_cl"]), CL.mkApply("clCreateBuffer",  
                                                                 [CL.E_Var contextVar,  
                                                                 CL.E_Var "CL_MEM_COPY_HOST_PTR",  
                                                                 CL.mkApply("sizeof",[CL.E_Var RN.globalsTy]),  
                                                                 CL.E_Var RN.globalsVarName,  
                                                                 CL.E_UnOp(CL.%&,CL.E_Var errVar)]))  
376    
377          fun genDataBuffers([],_,_,_) = []        (* generates the opencl buffers for the image data *)
378            | genDataBuffers((var,nDims)::globals,count,contextVar,errVar) = let          fun getGlobalDataBuffers (globals,contextVar,errVar) = let
379                  val globalBufferDecl =  CL.mkDecl(clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)
380                  val globalBuffer = CL.mkAssign(CL.mkVar(concat[RN.globalsVarName,"_cl"]),
381                        CL.mkApply("clCreateBuffer", [
382                            CL.mkVar contextVar,
383                            CL.mkVar "CL_MEM_COPY_HOST_PTR",
384                            CL.mkApply("sizeof",[CL.mkVar RN.globalsTy]),
385                            CL.mkVar RN.globalsVarName,
386                            CL.mkUnOp(CL.%&,CL.mkVar errVar)
387                          ]))
388                  fun genDataBuffers ([],_,_) = []
389                    | genDataBuffers ((var,nDims)::globals, contextVar, errVar) = let
390                        val hostVar = CL.mkIndirect(CL.mkVar RN.globalsVarName, var)
391          (* FIXME: use CL constructors to  build expressions (not strings) *)          (* FIXME: use CL constructors to  build expressions (not strings) *)
392                    val size = if nDims = 1 then                      fun sizeExp i = CL.mkSubscript(CL.mkIndirect(hostVar, "size"), CL.mkInt i)
393                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.E_Var "float"]), CL.#*,                      val size = CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*, sizeExp 0)
394                                           CL.mkIndirect(CL.E_Var var, "size[0]"))                      val size = if (nDims > 1)
395                                          else if nDims = 2 then                            then CL.mkBinOp(size, CL.#*, sizeExp 1)
396                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.E_Var "float"]), CL.#*,                            else size
397                                            CL.mkIndirect(CL.E_Var var, concat["size[0]", " * ", var, "->size[1]"]))                      val size = if (nDims > 2)
398                                          else                            then CL.mkBinOp(size, CL.#*, sizeExp 2)
399                                           CL.mkBinOp(CL.mkApply("sizeof",[CL.E_Var "float"]), CL.#*,                            else size
400                                            CL.mkIndirect(CL.E_Var var,concat["size[0]", " * ", var, "->size[1] * ", var, "->size[2]"]))                      in
401                          CL.mkDecl(clMemoryTy, RN.addBufferSuffix var ,NONE)::
402                   in                        CL.mkDecl(clMemoryTy, RN.addBufferSuffixData var ,NONE)::
403                     CL.mkDecl(CL.clMemoryTy,RN.addBufferSuffix var ,NONE)::                        CL.mkAssign(CL.mkVar(RN.addBufferSuffix var),
404                     CL.mkDecl(CL.clMemoryTy,RN.addBufferSuffixData var ,NONE)::                          CL.mkApply("clCreateBuffer", [
405                     CL.mkAssign(CL.E_Var(RN.addBufferSuffix var), CL.mkApply("clCreateBuffer",                              CL.mkVar contextVar,
406                                                                  [CL.E_Var contextVar,                              CL.mkVar "CL_MEM_COPY_HOST_PTR",
407                                                                  CL.E_Var "CL_MEM_COPY_HOST_PTR",                              CL.mkApply("sizeof",[CL.mkVar (RN.imageTy nDims)]),
408                                                                  CL.mkApply("sizeof",[CL.E_Var (RN.imageTy nDims)]),                              hostVar,
409                                                                  CL.E_Var var,                              CL.mkUnOp(CL.%&,CL.mkVar errVar)
410                                                                  CL.E_UnOp(CL.%&,CL.E_Var errVar)])) ::                            ])) ::
411                          CL.mkAssign(CL.E_Var(RN.addBufferSuffixData var), CL.mkApply("clCreateBuffer",                        CL.mkAssign(CL.mkVar(RN.addBufferSuffixData var),
412                                                                  [CL.E_Var contextVar,                          CL.mkApply("clCreateBuffer", [
413                                                                   CL.E_Var "CL_MEM_COPY_HOST_PTR",                              CL.mkVar contextVar,
414                                CL.mkVar "CL_MEM_COPY_HOST_PTR",
415                                                                  size,                                                                  size,
416                                                                  CL.mkIndirect(CL.E_Var var,"data"),                              CL.mkIndirect(hostVar, "data"),
417                                                                  CL.E_UnOp(CL.%&,CL.E_Var errVar)])):: genDataBuffers(globals,count + 2,contextVar,errVar)                              CL.mkUnOp(CL.%&,CL.mkVar errVar)
418                              ])) :: genDataBuffers(globals,contextVar,errVar)
419                  end                  end
420          in          in
421                  [globalBufferDecl] @ [globalBuffer] @ genDataBuffers(globals,count + 2,contextVar,errVar)                  globalBufferDecl :: globalBuffer :: genDataBuffers(globals,contextVar,errVar)
422          end          end
423    
   
424  (* generates the kernel arguments for the image data *)  (* generates the kernel arguments for the image data *)
425          fun genGlobalArguments(globals,count,kernelVar,errVar) = let          fun genGlobalArguments(globals,count,kernelVar,errVar) = let
426          val globalArgument = CL.mkExpStm(CL.mkAssignOp(CL.E_Var errVar,CL.|=,CL.mkApply("clSetKernelArg",                val globalArgument = CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,
427                                                                  [CL.E_Var kernelVar,                      CL.mkApply("clSetKernelArg",
428                                                                   CL.E_Int(count,CL.intTy),                        [CL.mkVar kernelVar,
429                                                                   CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),                         CL.mkPostOp(CL.E_Var count, CL.^++),
430                                                                   CL.E_UnOp(CL.%&,CL.E_Var(concat[RN.globalsVarName,"_cl"]))])))                         CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
431                           CL.mkUnOp(CL.%&,CL.mkVar(concat[RN.globalsVarName,"_cl"]))])))
432          fun genDataArguments([],_,_,_) = []          fun genDataArguments([],_,_,_) = []
433            | genDataArguments((var,nDims)::globals,count,kernelVar,errVar) =            | genDataArguments((var,nDims)::globals,count,kernelVar,errVar) =
434                        CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,
435                  CL.mkExpStm(CL.mkAssignOp(CL.E_Var errVar,CL.|=, CL.mkApply("clSetKernelArg",                        CL.mkApply("clSetKernelArg",
436                                                                  [CL.E_Var kernelVar,                          [CL.mkVar kernelVar,
437                                                                   CL.E_Int(count,CL.intTy),                           CL.mkPostOp(CL.E_Var count, CL.^++),
438                                                                   CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),                           CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
439                                                                   CL.E_UnOp(CL.%&,CL.E_Var(RN.addBufferSuffix var))])))::                           CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffix var))]))) ::
440                        CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,
441                          CL.mkExpStm(CL.mkAssignOp(CL.E_Var errVar,CL.|=,CL.mkApply("clSetKernelArg",                        CL.mkApply("clSetKernelArg",
442                                                                  [CL.E_Var kernelVar,                          [CL.mkVar kernelVar,
443                                                                   CL.E_Int((count + 1),CL.intTy),                           CL.mkPostOp(CL.E_Var count, CL.^++),
444                                                                   CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),                           CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
445                                                                   CL.E_UnOp(CL.%&,CL.E_Var(RN.addBufferSuffixData var))]))):: genDataArguments (globals, count + 2,kernelVar,errVar)                           CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffixData var))]))) ::
446                        genDataArguments (globals,count,kernelVar,errVar)
         in  
   
                 [globalArgument] @ genDataArguments(globals,count + 1,kernelVar,errVar)  
   
         end  
   
         (* generates the main function of host code *)  
         fun genHostMain() = let  
               val setupCall = [CL.mkCall(RN.setupFName,[CL.E_Var RN.globalsVarName])]  
               val globalsDecl = CL.mkDecl(  
                     CL.T_Ptr(CL.T_Named RN.globalsTy),  
                     RN.globalsVarName,  
                     SOME(CL.I_Exp(CL.mkApply("malloc", [CL.mkApply("sizeof",[CL.E_Var RN.globalsTy])]))))  
               val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.E_Var RN.globalsVarName])  
               val returnStm = [CL.mkReturn(SOME(CL.E_Int(0,CL.intTy)))]  
               val params = [  
                      CL.PARAM([],CL.intTy, "argc"),  
                      CL.PARAM([],CL.charArrayPtr,"argv")  
                    ]  
               val body = CL.mkBlock([globalsDecl] @ [initGlobalsCall]  @ setupCall @ returnStm)  
447                in                in
448                  CL.D_Func([],CL.intTy,"main",params,body)                  globalArgument :: genDataArguments(globals, count, kernelVar, errVar)
449                end                end
450    
451        (* generates the host-side setup function *)        (* generates the globals buffers and arguments function *)
452          fun genHostSetupFunc (strand as Strand{name,tyName,...}, filename, nDims, initially, imgGlobals) = let          fun genGlobalBuffersArgs (imgGlobals) = let
453              (* Delcare opencl setup objects *)              (* Delcare opencl setup objects *)
               val programVar= "program"  
               val kernelVar = "kernel"  
               val cmdVar = "queue"  
               val inStateVar = "selfin"  
               val outStateVar = "selfout"  
               val stateSizeVar= "state_mem_size"  
               val clInstateVar = "clSelfIn"  
               val clOutStateVar = "clSelfOut"  
               val clGlobals = "clGlobals"  
               val sourcesVar = "sources"  
               val contextVar = "context"  
454                val errVar = "err"                val errVar = "err"
455                val imgDataSizeVar = "image_dataSize"                val imgDataSizeVar = "image_dataSize"
               val globalVar = "global_work_size"  
               val localVar = "local_work_size"  
               val clFNVar = "filename"  
               val numStrandsVar = "numStrandsVar"  
               val headerFNVar = "header"  
               val deviceVar = "device"  
               val platformsVar = "platforms"  
               val numPlatformsVar = "num_platforms"  
               val numDevicesVar = "num_devices"  
               val assertStm = CL.mkCall("assert",[CL.mkBinOp(CL.E_Var errVar, CL.#==, CL.E_Var "CL_SUCCESS")])  
456                val params = [                val params = [
457                        CL.PARAM([],CL.T_Named("cl_device_id"), deviceVar)                        CL.PARAM([],CL.T_Named("cl_context"), "context"),
458                      ]                        CL.PARAM([],CL.T_Named("cl_kernel"), "kernel"),
459                val declarations = [                        CL.PARAM([],CL.T_Named("int"), "argStart")
                     CL.mkDecl(CL.clProgramTy, programVar, NONE),  
                     CL.mkDecl(CL.clKernelTy, kernelVar, NONE),  
                     CL.mkDecl(CL.clCmdQueueTy, cmdVar, NONE),  
                     CL.mkDecl(CL.clContextTy, contextVar, NONE),  
                     CL.mkDecl(CL.intTy, errVar, NONE),  
                     CL.mkDecl(CL.intTy, numStrandsVar, SOME(CL.I_Exp(CL.E_Int(1,CL.intTy)))),  
                     CL.mkDecl(CL.intTy, stateSizeVar, NONE),  
                     CL.mkDecl(CL.intTy, "width", NONE),  
                     CL.mkDecl(CL.intTy, imgDataSizeVar, NONE),  
                     (*CL.mkDecl(CL.clDeviceIdTy, deviceVar, NONE), *)  
                     CL.mkDecl(CL.T_Ptr(CL.T_Named tyName), inStateVar,NONE),  
                     CL.mkDecl(CL.clMemoryTy,clInstateVar,NONE),  
                     CL.mkDecl(CL.clMemoryTy,clOutStateVar,NONE),  
                     CL.mkDecl(CL.T_Ptr(CL.T_Named tyName), outStateVar,NONE),  
                     CL.mkDecl(CL.charPtr, clFNVar,SOME(CL.I_Exp(CL.E_Str filename))),  
 (* FIXME:  use Paths.diderotInclude *)  
                     CL.mkDecl(CL.charPtr, headerFNVar,SOME(CL.I_Exp(CL.E_Str "../src/include/Diderot/cl-types.h"))),  
                     CL.mkDecl(CL.T_Array(CL.charPtr,SOME(2)),sourcesVar,NONE),  
                     CL.mkDecl(CL.T_Array(CL.T_Named "size_t",SOME(nDims)),globalVar,NONE),  
                     CL.mkDecl(CL.T_Array(CL.T_Named "size_t",SOME(nDims)),localVar,NONE),  
                     CL.mkDecl(CL.intTy,numDevicesVar,SOME(CL.I_Exp(CL.E_Int(~1,CL.intTy)))),  
                     CL.mkDecl(CL.T_Array(CL.T_Named "cl_platform_id", SOME(1)), platformsVar, NONE),  
                     CL.mkDecl(CL.intTy,"num_platforms",SOME(CL.I_Exp(CL.E_Int(~1,CL.intTy))))  
460                  ]                  ]
461              (* Setup Global Variables *)                val clGlobalBuffers = getGlobalDataBuffers(!imgGlobals, "context", errVar)
462                val globalsDecl = CL.mkDecl(                val clGlobalArguments = genGlobalArguments(!imgGlobals, "argStart", "kernel", errVar)
                     CL.T_Ptr(CL.T_Named RN.globalsTy),  
                     RN.globalsVarName,  
                     SOME(CL.I_Exp(CL.mkApply("malloc", [CL.mkApply("sizeof",[CL.E_Var RN.globalsTy])]))))  
               val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.E_Var RN.globalsVarName])  
   
                 (* Retrieve the platforms  
                 val platformStm = [CL.mkAssign(CL.E_Var errVar, CL.mkApply("clGetPlatformIDs",  
                                                   [CL.E_Int(10,CL.intTy),  
                                                    CL.E_Var platformsVar,  
                                                    CL.E_UnOp(CL.%&,CL.E_Var numPlatformsVar)])),  
                                                    assertStm]  
   
                 val devicesStm = [CL.mkAssign(CL.E_Var errVar, CL.mkApply("clGetDeviceIDs",  
                                                   [CL.mkSubscript(CL.E_Var platformsVar,CL.E_Int(0,CL.intTy)),  
                                                    CL.E_Var "CL_DEVICE_TYPE_GPU",  
                                                    CL.E_Int(1,CL.intTy),  
                                                    CL.E_UnOp(CL.%&,CL.E_Var deviceVar),  
                                                    CL.E_UnOp(CL.%&,CL.E_Var numDevicesVar)])),  
                                                    assertStm] *)  
   
                 (* Create Context *)  
                 val contextStm = [CL.mkAssign(CL.E_Var contextVar, CL.mkApply("clCreateContext",  
                                                   [CL.E_Int(0,CL.intTy),  
                                                   CL.E_Int(1,CL.intTy),  
                                                   CL.E_UnOp(CL.%&,CL.E_Var deviceVar),  
                                                   CL.E_Var "NULL",  
                                                   CL.E_Var "NULL",  
                                                   CL.E_UnOp(CL.%&,CL.E_Var errVar)])),  
                                                   assertStm]  
   
                 (* Create Command Queue *)  
                 val commandStm = [CL.mkAssign(CL.E_Var cmdVar, CL.mkApply("clCreateCommandQueue",  
                                                   [CL.E_Var contextVar,  
                                                   CL.E_Var deviceVar,  
                                                   CL.E_Int(0,CL.intTy),  
                                                   CL.E_UnOp(CL.%&,CL.E_Var errVar)])),  
                                                   assertStm]  
   
   
                 (*Create Program/Build/Kernel with Source statement *)  
                 val createProgStm = CL.mkAssign(CL.E_Var programVar, CL.mkApply("clCreateProgramWithSource",  
                                                                                                                 [CL.E_Var contextVar,  
                                                                                                                  CL.E_Int(2,CL.intTy),  
                                                                                                                  CL.E_Cast(CL.T_Ptr(CL.T_Named("const char *")),CL.E_UnOp(CL.%&,CL.E_Var sourcesVar)),  
                                                                                                                  CL.E_Var "NULL",  
                                                                                                                  CL.E_UnOp(CL.%&,CL.E_Var errVar)]))  
   
                 (* FIXME: Remove after testing purposes, Build Log for OpenCL*)  
                 val buildLog = [CL.mkAssign(CL.E_Var errVar, CL.mkApply("clBuildProgram",  
                                                                                                                 [CL.E_Var programVar,  
                                                                                                                  CL.E_Int(0,CL.intTy),  
                                                                                                                  CL.E_Var "NULL",  
                                                                                                                  CL.E_Var "NULL",  
                                                                                                                  CL.E_Var "NULL",  
                                                                                                                  CL.E_Var "NULL"])),  
                                           CL.mkDecl(CL.charPtr, "build", NONE),  
                                           CL.mkDecl(CL.T_Named("size_t"),"ret_val_size",NONE),  
                                            CL.mkAssign(CL.E_Var errVar, CL.mkApply("clGetProgramBuildInfo",  
                                                                                                                 [CL.E_Var programVar,  
                                                                                                                 CL.E_Var deviceVar,  
                                                                                                                  CL.E_Var "CL_PROGRAM_BUILD_LOG",  
                                                                                                                  CL.E_Int(0,CL.intTy),  
                                                                                                                  CL.E_Var "NULL",  
                                                                                                                  CL.E_UnOp(CL.%&,CL.E_Var "ret_val_size")])),  
                                           CL.mkAssign(CL.E_Var "build", CL.mkApply("malloc", [CL.E_Var "ret_val_size"])),  
                                                 CL.mkAssign(CL.E_Var errVar, CL.mkApply("clGetProgramBuildInfo",  
                                                                                                                 [CL.E_Var programVar,  
                                                                                                                 CL.E_Var deviceVar,  
                                                                                                                  CL.E_Var "CL_PROGRAM_BUILD_LOG",  
                                                                                                                  CL.E_Var "ret_val_size",  
                                                                                                                  CL.E_Var "build",  
                                                                                                                  CL.E_Var "NULL"])),  
                                                 CL.mkAssign(CL.mkSubscript(CL.E_Var "build",CL.E_Var "ret_val_size"),CL.E_Var ("'\\" ^ "0'")),  
                                                 CL.mkCall("printf",[CL.E_Str ( "Build Log:" ^ "\n" ^ "%s" ^ "\n"), CL.E_Var "build"])]  
   
   
   
   
                 val createKernel = CL.mkAssign(CL.E_Var kernelVar, CL.mkApply("clCreateKernel",  
                                                                                                                 [CL.E_Var programVar,  
                                                                                                                  CL.E_Str RN.kernelFuncName,  
                                                                                                                  CL.E_UnOp(CL.%&,CL.E_Var errVar)]))  
   
   
                 val create_build_stms = [createProgStm,assertStm] @ buildLog @ [assertStm,createKernel,assertStm]  
   
   
   
                 (* Create Memory Buffers for Strand States and Globals *)  
                 val strandSize = CL.mkAssign(CL.E_Var stateSizeVar,CL.mkBinOp(CL.mkApply("sizeof",  
                                                                         [CL.E_Var tyName]), CL.#*,CL.E_Var numStrandsVar))  
   
                 val clStrandObjects = [CL.mkAssign(CL.E_Var clInstateVar, CL.mkApply("clCreateBuffer",  
                                                                 [CL.E_Var contextVar,  
                                                                 CL.E_Var "CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR",  
                                                                 CL.E_Var stateSizeVar,  
                                                                 CL.E_Var "NULL",  
                                                                 CL.E_UnOp(CL.%&,CL.E_Var errVar)])),  
                                                          CL.mkAssign(CL.E_Var clOutStateVar, CL.mkApply("clCreateBuffer",  
                                                                 [CL.E_Var contextVar,  
                                                                 CL.E_Var "CL_MEM_READ_WRITE",  
                                                                 CL.E_Var stateSizeVar,  
                                                                 CL.E_Var "NULL",  
                                                                 CL.E_UnOp(CL.%&,CL.E_Var errVar)]))]  
   
   
                 (* Setup up selfOut variable *)  
                 val strandsArrays = [CL.mkAssign(CL.E_Var outStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.E_Var numStrandsVar,  
                                                                         CL.#*, CL.mkApply("sizeof",[CL.E_Var tyName]))])),  
                                                                 CL.mkAssign(CL.E_Var inStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.E_Var numStrandsVar,  
                                                                         CL.#*, CL.mkApply("sizeof",[CL.E_Var tyName]))]))]  
   
   
                 (* Initialize Width Parameter *)  
                 val widthDel = if nDims = 2 then  
                           CL.mkAssign(CL.E_Var "width",CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(1, CL.intTy)))  
                    else  
                           CL.mkAssign(CL.E_Var "width",CL.E_Int(0,CL.intTy))  
   
   
                 val strands_init =      CL.mkCall(RN.strandInitSetup,  
                                                                          [CL.E_Var "size",  
                                                                          CL.E_Var "width",  
                                                                          CL.E_Var inStateVar])  
   
             val clGlobalBuffers = getGlobalDataBuffers(!imgGlobals,3,contextVar,errVar)  
   
   
                 (* Load the Kernel and Header Files *)  
                 val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.E_Var sourcesVar,CL.E_Int(1,CL.intTy)),  
                                                                           CL.mkApply(RN.clLoaderFN, [CL.E_Var clFNVar])),  
            CL.mkAssign(CL.mkSubscript(CL.E_Var sourcesVar,CL.E_Int(0,CL.intTy)),  
                                                                           CL.mkApply(RN.clLoaderFN, [CL.E_Var headerFNVar]))]  
   
                 (* val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.E_Var sourcesVar,CL.E_Int(1,CL.intTy)),  
                                                                           CL.mkApply(RN.clLoaderFN, [CL.E_Var clFNVar]))] *)  
   
   
                 (* Created Enqueue Statements *)  
 (* FIXME: simplify this code by function abstraction *)  
         val enqueueStm = if nDims = 1  
                         then [CL.mkAssign(CL.E_Var errVar,  
                                                           CL.mkApply("clEnqueueNDRangeKernel",  
                                                                                                 [CL.E_Var cmdVar,  
                                                                                                  CL.E_Var kernelVar,  
                                                                                                  CL.E_Int(1,CL.intTy),  
                                                                                                  CL.E_Var "NULL",  
                                                                                                  CL.E_Var globalVar,  
                                                                                                  CL.E_Var localVar,  
                                                                                                  CL.E_Int(0,CL.intTy),  
                                                                                                  CL.E_Var "NULL",  
                                                                                                  CL.E_Var "NULL"])),CL.mkCall("clFinish",[CL.E_Var cmdVar])]  
                         else if nDims = 2  then  
                          [CL.mkAssign(CL.E_Var errVar,  
                                                         CL.mkApply("clEnqueueNDRangeKernel",  
                                                                                                 [CL.E_Var cmdVar,  
                                                                                                  CL.E_Var kernelVar,  
                                                                                                  CL.E_Int(2,CL.intTy),  
                                                                                                  CL.E_Var "NULL",  
                                                                                                  CL.E_Var globalVar,  
                                                                                                  CL.E_Var localVar,  
                                                                                                  CL.E_Int(0,CL.intTy),  
                                                                                                  CL.E_Var "NULL",  
                                                                                                  CL.E_Var "NULL"])),CL.mkCall("clFinish",[CL.E_Var cmdVar])]  
                         else  
                           [CL.mkAssign(CL.E_Var errVar,  
                                                         CL.mkApply("clEnqueueNDRangeKernel",  
                                                                                                 [CL.E_Var cmdVar,  
                                                                                                  CL.E_Var kernelVar,  
                                                                                                  CL.E_Int(3,CL.intTy),  
                                                                                                  CL.E_Var "NULL",  
                                                                                                  CL.E_Var globalVar,  
                                                                                                  CL.E_Var localVar,  
                                                                                                  CL.E_Int(0,CL.intTy),  
                                                                                                  CL.E_Var "NULL",  
                                                                                                  CL.E_Var "NULL"])),CL.mkCall("clFinish",[CL.E_Var cmdVar])]  
   
   
   
                 (* Setup Global and Local variables *)  
   
                 val globalAndlocalStms = if nDims = 1 then  
                         [CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.E_Var "size", CL.E_Int(0,CL.intTy))),  
                          CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(0,CL.intTy)),  
                                                                   CL.E_Var "16")]  
   
   
                 else if nDims = 2 then  
                         [CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.E_Var "size", CL.E_Int(0,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(1,CL.intTy)),  
                                                                    CL.mkSubscript(CL.E_Var "size", CL.E_Int(1,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(0,CL.intTy)),  
                                                                   CL.E_Var "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(1,CL.intTy)),  
                                                                   CL.E_Var "16")]  
   
                 else  
                         [CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.E_Var "size", CL.E_Int(0,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(1,CL.intTy)),  
                                                                    CL.mkSubscript(CL.E_Var "size", CL.E_Int(1,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(2,CL.intTy)),  
                                                                    CL.mkSubscript(CL.E_Var "size", CL.E_Int(2,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(0,CL.intTy)),  
                                                                   CL.E_Var "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(1,CL.intTy)),  
                                                                   CL.E_Var "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(2,CL.intTy)),  
                                                                   CL.E_Var "16")]  
   
   
   
                 (* Setup Kernel arguments *)  
                 val kernelArguments = [CL.mkAssign(CL.E_Var errVar,CL.mkApply("clSetKernelArg",  
                                                                 [CL.E_Var kernelVar,  
                                                                  CL.E_Int(0,CL.intTy),  
                                                                  CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),  
                                                                  CL.E_UnOp(CL.%&,CL.E_Var clInstateVar)])),  
                                                             CL.mkExpStm(CL.mkAssignOp(CL.E_Var errVar, CL.|=,CL.mkApply("clSetKernelArg",  
                                                                 [CL.E_Var kernelVar,  
                                                                  CL.E_Int(1,CL.intTy),  
                                                                  CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),  
                                                                  CL.E_UnOp(CL.%&,CL.E_Var clOutStateVar)]))),  
                                                                   CL.mkExpStm(CL.mkAssignOp(CL.E_Var errVar, CL.|=,CL.mkApply("clSetKernelArg",  
                                                                 [CL.E_Var kernelVar,  
                                                                  CL.E_Int(2,CL.intTy),  
                                                                  CL.mkApply("sizeof",[CL.E_Var "int"]),  
                                                                  CL.E_UnOp(CL.%&,CL.E_Var "width")])))]  
   
            val clGlobalArguments = genGlobalArguments(!imgGlobals,3,kernelVar,errVar) @ [assertStm]  
   
                 (* Retrieve output *)  
                 val outputStm = CL.mkAssign(CL.E_Var errVar,  
                                                         CL.mkApply("clEnqueueReadBuffer",  
                                                                                                 [CL.E_Var cmdVar,  
                                                                                                  CL.E_Var clOutStateVar,  
                                                                                                  CL.E_Var "CL_TRUE",  
                                                                                                  CL.E_Int(0,CL.intTy),  
                                                                                                  CL.E_Var stateSizeVar,  
                                                                                                  CL.E_Var outStateVar,  
                                                                                                  CL.E_Int(0,CL.intTy),  
                                                                                                  CL.E_Var "NULL",  
                                                                                                  CL.E_Var "NULL"]))  
   
                 (* Free all the objects *)  
                 val freeStms = [CL.mkCall("clReleaseKernel",[CL.E_Var kernelVar]),  
                                                 CL.mkCall("clReleaseProgram",[CL.E_Var programVar ]),  
                                                 CL.mkCall("clReleaseCommandQueue",[CL.E_Var cmdVar]),  
                                                 CL.mkCall("clReleaseContext",[CL.E_Var contextVar]),  
                                                 CL.mkCall("clReleaseMemObject",[CL.E_Var clInstateVar]),  
                                                 CL.mkCall("clReleaseMemObject",[CL.E_Var clOutStateVar])]  
   
   
                 (*Setup Strand Print Function *)  
                 val outputData = [CL.mkDecl(CL.T_Ptr(CL.T_Named("FILE")), "outS", SOME(CL.I_Exp(CL.mkApply("fopen",  
                                                 [CL.E_Str "mip.txt",  
                                                 CL.E_Str "w"])))),  
                                                 CL.mkCall(concat[name, "_print"],  
                                                                         [CL.E_Var "outS",  
                                                                          CL.E_Var "size",  
                                                                          CL.E_Var "width",  
                                                                          CL.E_Var outStateVar])]  
   
   
   
463                  (* Body put all the statments together *)                  (* Body put all the statments together *)
464                  val body =  declarations @ [globalsDecl,initGlobalsCall] (*@ platformStm @ devicesStm *) @ contextStm @ commandStm @ !initially @ [strandSize] @                val body = CL.mkDecl(clIntTy, errVar, SOME(CL.I_Exp(CL.mkInt 0)))
465                                     strandsArrays @ globalAndlocalStms @ [widthDel,strands_init]  @ clStrandObjects @ clGlobalBuffers @ sourceStms  @ create_build_stms  (*@                      :: clGlobalBuffers @ clGlobalArguments
                                    kernelArguments @ clGlobalArguments @ enqueueStm @  [outputStm] @ freeStms @ outputData *)  
   
466                  in                  in
467                    CL.D_Func([],CL.voidTy,RN.globalsSetupName,params,CL.mkBlock(body))
                 CL.D_Func([],CL.voidTy,RN.setupFName,params,CL.mkBlock(body))  
   
468                  end                  end
469    
470  (* generate the data and global parameters *)  (* generate the data and global parameters *)
471          fun genKeneralGlobalParams ((name,tyname)::rest) =          fun genKeneralGlobalParams ((name,tyname)::rest) =
472                  CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), concat[RN.globalsVarName]) ::                  CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), concat[RN.globalsVarName]) ::
473                  CL.PARAM([], CL.T_Ptr(CL.T_Named (RN.imageTy tyname)),RN.addBufferSuffix name) ::                  CL.PARAM([], CL.T_Ptr(CL.T_Named (RN.imageTy tyname)),RN.addBufferSuffix name) ::
474                  CL.PARAM([], CL.T_Ptr(CL.voidTy),RN.addBufferSuffixData name) ::                  CL.PARAM([], CL.T_Ptr(CL.voidTy),RN.addBufferSuffixData name) ::
475                  genKeneralGlobalParams(rest)                genKeneralGlobalParams rest
476              | genKeneralGlobalParams [] = []
           | genKeneralGlobalParams ([]) = []  
477    
478          (*generate code for intilizing kernel global data *)          (*generate code for intilizing kernel global data *)
479          fun initKernelGlobals (globals,imgGlobals) = let  (* FIXME: should use List.map here *)
                 fun initGlobalStruct (CL.D_Var(_, _ , name, _)::rest) =  
                                 CL.mkAssign(CL.E_Var name, CL.mkIndirect(CL.E_Var RN.globalsVarName, name)) ::  
                                 initGlobalStruct(rest)  
                   | initGlobalStruct ( _::rest) = initGlobalStruct(rest)  
                   | initGlobalStruct([]) = []  
   
480                  fun initGlobalImages((name,tyname)::rest) =                  fun initGlobalImages((name,tyname)::rest) =
481                                  CL.mkAssign(CL.E_Var name, CL.E_Var (RN.addBufferSuffix name)) ::               CL.mkAssign(CL.mkIndirect(CL.E_Var RN.globalsVarName, name), CL.mkVar (RN.addBufferSuffix name)) ::
482                                  CL.mkAssign(CL.mkIndirect(CL.E_Var name,"data"),CL.E_Var (RN.addBufferSuffixData name)) ::               CL.mkAssign(CL.mkIndirect(CL.E_Var RN.globalsVarName,concat[name,"->","data"]),CL.mkVar (RN.addBufferSuffixData name)) ::
483                                  initGlobalImages(rest)               initGlobalImages rest
484                    | initGlobalImages([]) = []            | initGlobalImages [] = []
                 in  
                   initGlobalStruct(globals) @ initGlobalImages(imgGlobals)  
                 end  
485    
486          (* generate the main kernel function for the .cl file *)          (* generate the main kernel function for the .cl file *)
487          fun genKernelFun(Strand{name, tyName, state, output, code,...},nDims,globals,imgGlobals) = let          fun genKernelFun (strand, nDims, globals, imgGlobals) = let
488                  val Strand{name, tyName, state, output, code,...} = strand
489                   val fName = RN.kernelFuncName;                   val fName = RN.kernelFuncName;
490                   val inState = "strand_in"                   val inState = "strand_in"
491                   val outState = "strand_out"                   val outState = "strand_out"
# Line 830  Line 495 
495                        CL.PARAM(["__global"], CL.intTy, "width")                        CL.PARAM(["__global"], CL.intTy, "width")
496                      ] @ genKeneralGlobalParams(!imgGlobals)                      ] @ genKeneralGlobalParams(!imgGlobals)
497                    val thread_ids = if nDims = 1                    val thread_ids = if nDims = 1
498                          then [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy)))),                      then [
499                                    CL.mkAssign(CL.E_Var "x",CL.mkApply(RN.getGlobalThreadId,[CL.E_Int(0,CL.intTy)]))]                          CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt 0))),
500                          else                          CL.mkAssign(CL.mkVar "x",CL.mkApply(RN.getGlobalThreadId,[CL.mkInt 0]))
501                                  [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy)))),                        ]
502                                   CL.mkDecl(CL.intTy, "y", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy)))),                      else [
503                                    CL.mkAssign(CL.E_Var "x",  CL.mkApply(RN.getGlobalThreadId,[CL.E_Int(0,CL.intTy)])),                          CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt 0))),
504                                    CL.mkAssign(CL.E_Var "y",CL.mkApply(RN.getGlobalThreadId,[CL.E_Int(1,CL.intTy)]))]                          CL.mkDecl(CL.intTy, "y", SOME(CL.I_Exp(CL.mkInt 0))),
505                            CL.mkAssign(CL.mkVar "x",  CL.mkApply(RN.getGlobalThreadId,[CL.mkInt 0])),
506                    val strandDecl = [CL.mkDecl(CL.T_Named tyName, inState, NONE),                          CL.mkAssign(CL.mkVar "y",CL.mkApply(RN.getGlobalThreadId,[CL.mkInt 1]))
507                          ]
508                  val strandDecl = [
509                        CL.mkDecl(CL.T_Named tyName, inState, NONE),
510                                                          CL.mkDecl(CL.T_Named tyName, outState,NONE)]                                                          CL.mkDecl(CL.T_Named tyName, outState,NONE)]
511                    val strandObjects  = if nDims = 1                    val strandObjects  = if nDims = 1
512                          then [CL.mkAssign(CL.mkSubscript(CL.E_Var "selfIn",CL.E_Str "x"),                      then [
513                                                                           CL.E_Var inState),                          CL.mkAssign( CL.mkVar inState, CL.mkSubscript(CL.mkVar "selfIn", CL.mkStr "x")),
514                                    CL.mkAssign(CL.mkSubscript(CL.E_Var "selfOut",CL.E_Str "x"),                          CL.mkAssign(CL.mkVar outState,CL.mkSubscript(CL.mkVar "selfOut", CL.mkStr "x"))
515                                                                           CL.E_Var outState)]                        ]
516                          else let                          else let
517                                  val index = CL.mkBinOp(CL.mkBinOp(CL.E_Var "x",CL.#*,CL.E_Var "width"),CL.#+,CL.E_Var "y")                        val index = CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")
518                                  in                        in [
519                                          [CL.mkAssign(CL.mkSubscript(CL.E_Var "selfIn",index),                          CL.mkAssign(CL.mkVar inState, CL.mkSubscript(CL.mkVar "selfIn",index)),
520                                                                          CL.E_Var inState),                          CL.mkAssign(CL.mkVar outState,CL.mkSubscript(CL.mkVar "selfOut",index))
521                                           CL.mkAssign(CL.mkSubscript(CL.E_Var "selfOut",index),                        ] end
522                                                                          CL.E_Var outState)]                val status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.mkInt 0)))
523                                  end                val strand_Init_Stm = CL.mkCall(RN.strandInit name, [CL.E_Var RN.globalsVarName,CL.mkUnOp(CL.%&,CL.E_Var inState), CL.E_Var "x", CL.E_Var "y"])
524                  val local_vars = thread_ids @ initGlobalImages(!imgGlobals)  @ strandDecl @ strandObjects @ [strand_Init_Stm,status]
525                                      val while_exp = CL.mkBinOp(
526                    val status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy))))                      CL.mkBinOp(CL.mkVar "status",CL.#!=, CL.mkVar RN.kStabilize),
527                    val local_vars =  thread_ids @ initKernelGlobals(!globals,!imgGlobals)  @ strandDecl @ strandObjects @ [status]                      CL.#||,
528                    val while_exp = CL.mkBinOp(CL.mkBinOp(CL.E_Var "status",CL.#!=, CL.E_Var RN.kStabilize),CL.#||,CL.mkBinOp(CL.E_Var "status", CL.#!=, CL.E_Var RN.kDie))                      CL.mkBinOp(CL.mkVar "status", CL.#!=, CL.mkVar RN.kDie))
529                    val while_body = [CL.mkAssign(CL.E_Var "status", CL.mkApply(RN.strandUpdate name,[ CL.E_UnOp(CL.%&,CL.E_Var inState), CL.E_UnOp(CL.%&,CL.E_Var outState)])),                val whileBody = CL.mkBlock [
530                                                          CL.mkCall(RN.strandStabilize name,[ CL.E_UnOp(CL.%&,CL.E_Var inState),  CL.E_UnOp(CL.%&,CL.E_Var outState)])]                        CL.mkAssign(CL.mkVar "status",
531                            CL.mkApply(RN.strandUpdate name,
532                    val whileBlock = [CL.mkWhile(while_exp,CL.mkBlock while_body)]                            [CL.mkUnOp(CL.%&,CL.mkVar inState), CL.mkUnOp(CL.%&,CL.mkVar outState),CL.E_Var RN.globalsVarName])),
533                          CL.mkCall(RN.strandStabilize name,
534                            [CL.mkUnOp(CL.%&,CL.mkVar inState), CL.mkUnOp(CL.%&,CL.mkVar outState),CL.E_Var RN.globalsVarName])
535                        ]
536                  val whileBlock = [CL.mkWhile(while_exp, whileBody)]
537                    val body = CL.mkBlock(local_vars  @ whileBlock)                    val body = CL.mkBlock(local_vars  @ whileBlock)
538                  in                  in
539                     CL.D_Func(["__kernel"], CL.voidTy, fName, params, body)                     CL.D_Func(["__kernel"], CL.voidTy, fName, params, body)
540                  end                  end
541          (* generate a global structure from the globals *)          (* generate a global structure from the globals *)
542          fun genGlobalStruct(globals) = let          fun genGlobalStruct (targetTy, globals) = let
543                   fun getGlobals(CL.D_Var(_,ty,globalVar,_)::rest) = (ty,globalVar)::getGlobals(rest)                val globs = List.map (fn (x : mirror_var) => (targetTy x, #var x)) globals
                    | getGlobals([]) = []  
                    | getGlobals(_::rest) = getGlobals(rest)  
544                   in                   in
545                          CL.D_StructDef(getGlobals(globals),RN.globalsTy)                  CL.D_StructDef(globs, RN.globalsTy)
546                  end
547            fun genGlobals (declFn, targetTy, globals) = let
548                  fun doVar (x : mirror_var) = declFn (CL.D_Var([], targetTy x, #var x, NONE))
549                  in
550                    List.app doVar globals
551                  end
552    
553            fun genStrandDesc (Strand{name, output, ...}) = let
554                (* the strand's descriptor object *)
555                  val descI = let
556                        fun fnPtr (ty, f) = CL.I_Exp(CL.mkCast(CL.T_Named ty, CL.mkVar f))
557                        val SOME(outTy, _) = !output
558                        in
559                          CL.I_Struct[
560                              ("name", CL.I_Exp(CL.mkStr name)),
561                              ("stateSzb", CL.I_Exp(CL.mkSizeof(CL.T_Named(N.strandTy name)))),
562    (*
563                              ("outputSzb", CL.I_Exp(CL.mkSizeof(ToC.trTy outTy))),
564    *)
565                              ("update", fnPtr("update_method_t", "0")),
566                              ("print", fnPtr("print_method_t", name ^ "_print"))
567                            ]
568                        end
569                  val desc = CL.D_Var([], CL.T_Named N.strandDescTy, N.strandDesc name, SOME descI)
570                  in
571                    desc
572                    end                    end
573    
574        (* generate the table of strand descriptors *)        (* generate the table of strand descriptors *)
575          fun genStrandTable (ppStrm, strands) = let          fun genStrandTable (declFn, strands) = let
576                val nStrands = length strands                val nStrands = length strands
577                fun genInit (Strand{name, ...}) = CL.I_Exp(CL.mkUnOp(CL.%&, CL.E_Var(RN.strandDesc name)))                fun genInit (Strand{name, ...}) = CL.I_Exp(CL.mkUnOp(CL.%&, CL.E_Var(N.strandDesc name)))
578                fun genInits (_, []) = []                fun genInits (_, []) = []
579                  | genInits (i, s::ss) = (i, genInit s) :: genInits(i+1, ss)                  | genInits (i, s::ss) = (i, genInit s) :: genInits(i+1, ss)
               fun ppDecl dcl = PrintAsC.output(ppStrm, dcl)  
580                in                in
581                  ppDecl (CL.D_Var([], CL.int32, RN.numStrands,                  declFn (CL.D_Var([], CL.int32, N.numStrands,
582                    SOME(CL.I_Exp(CL.E_Int(IntInf.fromInt nStrands, CL.int32)))));                    SOME(CL.I_Exp(CL.E_Int(IntInf.fromInt nStrands, CL.int32)))));
583                  ppDecl (CL.D_Var([],                  declFn (CL.D_Var([],
584                    CL.T_Array(CL.T_Ptr(CL.T_Named RN.strandDescTy), SOME nStrands),                    CL.T_Array(CL.T_Ptr(CL.T_Named N.strandDescTy), SOME nStrands),
585                    RN.strands,                    N.strands,
586                    SOME(CL.I_Array(genInits (0, strands)))))                    SOME(CL.I_Array(genInits (0, strands)))))
587                end                end
588    
589            fun genSrc (baseName, prog) = let
590          fun genSrc (baseName, Prog{double,globals, topDecls, strands, initially,imgGlobals,numDims,...}) = let                val Prog{name,double, globals, topDecls, strands, initially, imgGlobals, numDims, ...} = prog
591                val clFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "cl"}                val clFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "cl"}
592                val cFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "c"}                val cFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "c"}
593                val clOutS = TextIO.openOut clFileName                val clOutS = TextIO.openOut clFileName
594                val cOutS = TextIO.openOut cFileName                val cOutS = TextIO.openOut cFileName
595  (* FIXME: need to use PrintAsC and PrintAsCL *)                val clppStrm = PrintAsCL.new clOutS
               val clppStrm = PrintAsC.new clOutS  
596                val cppStrm = PrintAsC.new cOutS                val cppStrm = PrintAsC.new cOutS
597                  val progName = name
598                fun cppDecl dcl = PrintAsC.output(cppStrm, dcl)                fun cppDecl dcl = PrintAsC.output(cppStrm, dcl)
599                fun clppDecl dcl = PrintAsC.output(clppStrm, dcl)                fun clppDecl dcl = PrintAsCL.output(clppStrm, dcl)
600                val strands = AtomTable.listItems strands                val strands = AtomTable.listItems strands
601                val [strand as Strand{name, tyName, code,init_code, ...}] = strands                val [strand as Strand{name, tyName, code,init_code, ...}] = strands
602                in                in
# Line 912  Line 606 
606                        then "#define DIDEROT_DOUBLE_PRECISION"                        then "#define DIDEROT_DOUBLE_PRECISION"
607                        else "#define DIDEROT_SINGLE_PRECISION",                        else "#define DIDEROT_SINGLE_PRECISION",
608                      "#define DIDEROT_TARGET_CL",                      "#define DIDEROT_TARGET_CL",
609                      "#include \"Diderot/cl-types.h\""                      "#include \"Diderot/cl-diderot.h\""
610                    ]));                    ]));
611                  List.app clppDecl (List.rev (!globals));                  clppDecl (genGlobalStruct (#gpuTy, !globals));
612                  clppDecl (genGlobalStruct (!globals));                  clppDecl (genStrandTyDef(#gpuTy, strand));
613                  clppDecl (genStrandTyDef strand);                  clppDecl  (!init_code);
614                  List.app clppDecl (!code);                  List.app clppDecl (!code);
615                  clppDecl (genKernelFun (strand,!numDims,globals,imgGlobals));                  clppDecl (genKernelFun (strand,!numDims,globals,imgGlobals));
616                (* Generate the Host file .c *)                (* Generate the Host C file *)
617                  cppDecl (CL.D_Verbatim([                  cppDecl (CL.D_Verbatim([
618                      if double                      if double
619                        then "#define DIDEROT_DOUBLE_PRECISION"                        then "#define DIDEROT_DOUBLE_PRECISION"
# Line 927  Line 621 
621                      "#define DIDEROT_TARGET_CL",                      "#define DIDEROT_TARGET_CL",
622                      "#include \"Diderot/diderot.h\""                      "#include \"Diderot/diderot.h\""
623                    ]));                    ]));
624                  List.app cppDecl (List.rev (!globals));                  cppDecl (CL.D_Var(["static"], CL.charPtr, "ProgramName",
625                  cppDecl (genGlobalStruct (!globals));                    SOME(CL.I_Exp(CL.mkStr progName))));
626                  cppDecl (genStrandTyDef strand);                  cppDecl (genGlobalStruct (#hostTy, !globals));
627                  cppDecl  (!init_code);                  cppDecl (CL.D_Var(["static"], CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName, NONE));
628                  cppDecl (genStrandInit(strand,!numDims));                  cppDecl (genStrandTyDef (#hostTy, strand));
629                  cppDecl (genStrandPrint(strand,!numDims));                  cppDecl (genStrandPrint strand);
                 (* cppDecl (genKernelLoader());*)  
630                  List.app cppDecl (List.rev (!topDecls));                  List.app cppDecl (List.rev (!topDecls));
631                  cppDecl (genHostSetupFunc (strand, clFileName, !numDims, initially, imgGlobals));                  cppDecl (genGlobalBuffersArgs imgGlobals);
632                    List.app (fn strand => cppDecl (genStrandDesc strand)) strands;
633                    genStrandTable (cppDecl, strands);
634                    cppDecl (!initially);
635                  PrintAsC.close cppStrm;                  PrintAsC.close cppStrm;
636                  PrintAsC.close clppStrm;                  PrintAsCL.close clppStrm;
637                  TextIO.closeOut cOutS;                  TextIO.closeOut cOutS;
638                  TextIO.closeOut clOutS                  TextIO.closeOut clOutS
639                end                end
# Line 999  Line 695 
695          fun init (Strand{name, tyName, code,init_code, ...}, params, init) = let          fun init (Strand{name, tyName, code,init_code, ...}, params, init) = let
696                val fName = RN.strandInit name                val fName = RN.strandInit name
697                val params =                val params =
698                        CL.PARAM([], globPtrTy, RN.globalsVarName) ::
699                      CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfOut") ::                      CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfOut") ::
700                        List.map (fn (ToCL.V(ty, x)) => CL.PARAM([], ty, x)) params                        List.map (fn (ToCL.V(ty, x)) => CL.PARAM([], ty, x)) params
701                val initFn = CL.D_Func([], CL.voidTy, fName, params, init)                val initFn = CL.D_Func([], CL.voidTy, fName, params, init)
# Line 1011  Line 708 
708                val fName = concat[name, "_", methName]                val fName = concat[name, "_", methName]
709                val params = [                val params = [
710                        CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfIn"),                        CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfIn"),
711                        CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfOut")                        CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfOut"),
712                                                             CL.PARAM([], CL.T_Ptr(CL.T_Named (RN.globalsTy)), RN.globalsVarName)
713                      ]                      ]
714                val methFn = CL.D_Func([], CL.int32, fName, params, body)                val methFn = CL.D_Func([], CL.int32, fName, params, body)
715                in                in

Legend:
Removed from v.1273  
changed lines
  Added in v.1322

root@smlnj-gforge.cs.uchicago.edu
ViewVC Help
Powered by ViewVC 1.0.0