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 1278, Mon Jun 6 16:27:28 2011 UTC revision 1315, Sat Jun 11 21:10:15 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
# Line 32  Line 75 
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 70  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                          | InitiallyScope => ToC.trFragment (vMap, blk)
119                          | _ => ToCL.trFragment (vMap, blk)
120                        (* end case *))
121                in                in
122                  (ENV{info=info, vMap=vMap, scope=scope}, stms)                  (ENV{info=info, vMap=vMap, scope=scope}, stms)
123                end                end
124          fun saveState cxt stateVars (env, args, stm) = (          fun block (ENV{vMap, scope, ...}, blk) = let
125                  fun saveState cxt stateVars trAssign (env, args, stm) = (
126                ListPair.foldrEq                ListPair.foldrEq
127                  (fn (x, e, stms) => ToCL.trAssign(env, x, e)@stms)                        (fn (x, e, stms) => trAssign(env, x, e)@stms)
128                    [stm]                    [stm]
129                      (stateVars, args)                      (stateVars, args)
130                ) handle ListPair.UnequalLengths => (                ) handle ListPair.UnequalLengths => (
131                  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"]);
132                  raise Fail(concat["saveState ", cxt, ": length mismatch"]))                  raise Fail(concat["saveState ", cxt, ": length mismatch"]))
133          fun block (ENV{vMap, scope, ...}, blk) = (case scope                in
134                 of StrandScope stateVars => ToCL.trBlock (vMap, saveState "StrandScope" stateVars, blk)                  case scope
135                  | 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! *)
136                  | _ => ToCL.trBlock (vMap, fn (_, _, stm) => [stm], blk)                   of StrandScope stateVars =>
137                (* end case *))                        ToC.trBlock (vMap, saveState "StrandScope" stateVars ToC.trAssign, blk)
138                      | MethodScope stateVars =>
139                          ToCL.trBlock (vMap, saveState "MethodScope" stateVars ToCL.trAssign, blk)
140                      | InitiallyScope => ToC.trBlock (vMap, fn (_, _, stm) => [stm], blk)
141                      | _ => ToC.trBlock (vMap, fn (_, _, stm) => [stm], blk)
142                    (* end case *)
143                  end
144          fun exp (ENV{vMap, ...}, e) = ToCL.trExp(vMap, e)          fun exp (ENV{vMap, ...}, e) = ToCL.trExp(vMap, e)
145        end        end
146    
# Line 95  Line 149 
149        struct        struct
150          fun name (ToCL.V(_, name)) = name          fun name (ToCL.V(_, name)) = name
151          fun global (Prog{globals, imgGlobals, ...}, name, ty) = let          fun global (Prog{globals, imgGlobals, ...}, name, ty) = let
152                val ty' = ToCL.trType ty                val x = {hostTy = ToC.trType ty, gpuTy = ToCL.trType ty, var = name}
153                fun isImgGlobal (imgGlobals, Ty.ImageTy(ImageInfo.ImgInfo{dim, ...}), name) =  imgGlobals  := (name,dim):: !imgGlobals                fun isImgGlobal (Ty.ImageTy(ImageInfo.ImgInfo{dim, ...}), name) =
154                  | isImgGlobal (imgGlobals, _, _) =  ()                      imgGlobals  := (name,dim) :: !imgGlobals
155                in                  | isImgGlobal _ =  ()
156                  globals := CL.D_Var([], ty', name, NONE) :: !globals;                in
157                  isImgGlobal(imgGlobals,ty,name);                  globals := x :: !globals;
158               ToCL.V(ty', name)                  isImgGlobal (ty, name);
159                    ToCL.V(#gpuTy x, name)
160                end                end
161          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)
162          fun state (Strand{state, ...}, x) = let          fun state (Strand{state, ...}, x) = let
163                val ty' = ToCL.trType(V.ty x)                val ty = V.ty x
164                val x' = ToCL.V(ty', V.name x)                val x' = {hostTy = ToC.trType ty, gpuTy = ToCL.trType ty, var = V.name x}
165                in                in
166                  state := x' :: !state;                  state := x' :: !state;
167                  x'                  ToCL.V(#gpuTy x', #var x')
168                end                end
169        end        end
170    
# Line 141  Line 196 
196        struct        struct
197          fun new {name, double, parallel, debug} = (          fun new {name, double, parallel, debug} = (
198                RN.initTargetSpec double;                RN.initTargetSpec double;
199                  CNames.initTargetSpec double;
200                Prog{                Prog{
201                    name = name,                    name = name,
202                    double = double, parallel = parallel, debug = debug,                    double = double, parallel = parallel, debug = debug,
203                    globals = ref [],                    globals = ref [],
204                    topDecls = ref [],                    topDecls = ref [],
205                    strands = AtomTable.mkTable (16, Fail "strand table"),                    strands = AtomTable.mkTable (16, Fail "strand table"),
206                    initially = ref([CL.S_Comment["missing initially"]]),                    initially = ref(CL.D_Comment["missing initially"]),
207                                    numDims = ref(0),                                    numDims = ref(0),
208                                    imgGlobals = ref[],                                    imgGlobals = ref[],
209                                    prFn = ref(CL.D_Comment(["No Print Function"]))                                    prFn = ref(CL.D_Comment(["No Print Function"]))
210                  })                  })
211        (* register the global initialization part of a program *)        (* register the global initialization part of a program *)
212    (* FIXME: unused code; can this be removed??
213            fun globalIndirects (globals,stms) = let            fun globalIndirects (globals,stms) = let
214                  fun getGlobals(CL.D_Var(_,_,globalVar,_)::rest) = CL.mkAssign(CL.mkIndirect(CL.mkVar RN.globalsVarName,globalVar),CL.mkVar globalVar)::getGlobals(rest)                  fun getGlobals ({name,target as TargetUtil.TARGET_CL}::rest) =
215                    | getGlobals([]) = []                        CL.mkAssign(CL.mkIndirect(CL.mkVar RN.globalsVarName,name),CL.mkVar name)
216                    | getGlobals(_::rest) = getGlobals(rest)                          ::getGlobals rest
217                      | getGlobals [] = []
218                      | getGlobals (_::rest) = getGlobals rest
219                  in                  in
220                    stms @ getGlobals(globals)                    stms @ getGlobals globals
221                  end                  end
222    *)
223        (* 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 *)
224          fun inputs (Prog{topDecls, ...}, stm) = let          fun inputs (Prog{topDecls, ...}, stm) = let
225                val inputsFn = CL.D_Func(                val inputsFn = CL.D_Func(
# Line 171  Line 230 
230                  topDecls := inputsFn :: !topDecls                  topDecls := inputsFn :: !topDecls
231                end                end
232    
233          fun init (Prog{globals, topDecls,...}, CL.S_Block(init)) = let        (* register the global initialization part of a program *)
234                val params = [          fun init (Prog{topDecls, ...}, init) = let
235                        CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName)                val initFn = CL.D_Func(
236                      ]                      [], CL.voidTy, RN.initGlobals, [CL.PARAM([], globPtrTy, RN.globalsVarName)],
237                val body = CL.S_Block(globalIndirects(!globals,init))                      init)
238                val initFn = CL.D_Func([], CL.voidTy, RN.initGlobals, params, body)                val shutdownFn = CL.D_Func(
239                in                      [], CL.voidTy, RN.shutdown,
240                  topDecls := initFn :: !topDecls                      [CL.PARAM([], CL.T_Ptr(CL.T_Named RN.worldTy), "wrld")],
241                end                      CL.S_Block[])
           | init (Prog{globals,topDecls,...}, init) = let  
               val params = [  
                       CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName)  
                     ]  
               val initFn = CL.D_Func([], CL.voidTy, RN.initGlobals, params, init)  
242                in                in
243                  topDecls := initFn :: !topDecls                  topDecls := shutdownFn :: initFn :: !topDecls
244                end                end
   
245        (* create and register the initially function for a program *)        (* create and register the initially function for a program *)
246          fun initially {          fun initially {
247                prog = Prog{strands, initially, numDims,...},                prog = Prog{name=progName, strands, initially, ...},
248                isArray : bool,                isArray : bool,
249                iterPrefix : stm list,                iterPrefix : stm list,
250                iters : (var * exp * exp) list,                iters : (var * exp * exp) list,
# Line 201  Line 254 
254              } = let              } = let
255                val name = Atom.toString strand                val name = Atom.toString strand
256                val nDims = List.length iters                val nDims = List.length iters
257                  val worldTy = CL.T_Ptr(CL.T_Named N.worldTy)
258                fun mapi f xs = let                fun mapi f xs = let
259                      fun mapf (_, []) = []                      fun mapf (_, []) = []
260                        | mapf (i, x::xs) = f(i, x) :: mapf(i+1, xs)                        | mapf (i, x::xs) = f(i, x) :: mapf(i+1, xs)
# Line 209  Line 263 
263                      end                      end
264                val baseInit = mapi (fn (i, (_, e, _)) => (i, CL.I_Exp e)) iters                val baseInit = mapi (fn (i, (_, e, _)) => (i, CL.I_Exp e)) iters
265                val sizeInit = mapi                val sizeInit = mapi
266                      (fn (i, (ToCL.V(ty, _), lo, hi)) =>                      (fn (i, (CL.V(ty, _), lo, hi)) =>
267                          (i, CL.I_Exp(CL.mkBinOp(CL.mkBinOp(hi, CL.#-, lo), CL.#+, CL.mkInt(1, ty))))                          (i, CL.I_Exp(CL.mkBinOp(CL.mkBinOp(hi, CL.#-, lo), CL.#+, CL.E_Int(1, ty))))
268                      ) iters                      ) iters
269                    val numStrandsVar = "numStrandsVar"              (* code to allocate the world and initial strands *)
270                val allocCode = iterPrefix @ [                val wrld = "wrld"
271                  val allocCode = [
272                        CL.mkComment["allocate initial block of strands"],                        CL.mkComment["allocate initial block of strands"],
273                        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)),
274                        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)),
275                        CL.mkDecl(CL.int32,"numDims",SOME(CL.I_Exp(CL.mkInt(IntInf.fromInt nDims, CL.int32))))                        CL.mkDecl(worldTy, wrld,
276                      ]                          SOME(CL.I_Exp(CL.E_Apply(N.allocInitially, [
277                val numStrandsLoopBody = CL.mkExpStm(CL.mkAssignOp(CL.mkVar numStrandsVar, CL.*=,CL.mkSubscript(CL.mkVar "size",CL.mkVar "i")))                              CL.mkVar "ProgramName",
278                val numStrandsLoop =  CL.mkFor([(CL.intTy, "i", CL.mkInt(0,CL.intTy))],                              CL.mkUnOp(CL.%&, CL.E_Var(N.strandDesc name)),
279                      CL.mkBinOp(CL.mkVar "i", CL.#<, CL.mkVar "numDims"),                              CL.E_Bool isArray,
280                      [CL.mkPostOp(CL.mkVar "i", CL.^++)], numStrandsLoopBody)                              CL.E_Int(IntInf.fromInt nDims, CL.int32),
281                in                              CL.E_Var "base",
282                  numDims := nDims;                              CL.E_Var "size"
283                  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")  
284                      ]                      ]
285                val body = let              (* create the loop nest for the initially iterations *)
286                      fun loopParams 3 = ["x", "y", "k"]                val indexVar = "ix"
287                        | loopParams 2 = ["x", "y"]                val strandTy = CL.T_Ptr(CL.T_Named(N.strandTy name))
288                        | loopParams 1 = ["x"]                fun mkLoopNest [] = CL.mkBlock(createPrefix @ [
289                        | loopParams _ = raise Fail "genStrandInit: missing size dim"                        CL.mkDecl(strandTy, "sp",
290                      fun mkLoopNest ([], _, nDims) = if nDims = 1                          SOME(CL.I_Exp(
291                            then CL.mkBlock [                            CL.E_Cast(strandTy,
292                                CL.mkCall(RN.strandInit name, [                            CL.E_Apply(N.inState, [CL.E_Var "wrld", CL.E_Var indexVar]))))),
293                                  CL.mkUnOp(CL.%&,CL.mkSubscript(CL.mkVar "strands",CL.mkStr "x")),                        CL.mkCall(N.strandInit name,
294                                                  CL.mkVar "x"])                          CL.E_Var RN.globalsVarName :: CL.E_Var "sp" :: args),
295                              ]                        CL.mkAssign(CL.E_Var indexVar, CL.mkBinOp(CL.E_Var indexVar, CL.#+, CL.E_Int(1, CL.uint32)))
296                                          else let                      ])
297                                                  val index = CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")                  | mkLoopNest ((CL.V(ty, param), lo, hi)::iters) = let
298                                          in                      val body = mkLoopNest iters
                                                 CL.mkBlock([CL.mkCall(RN.strandInit name, [CL.mkUnOp(CL.%&,CL.mkSubscript(CL.mkVar "strands",index)),  
                                                 CL.mkVar "x", CL.mkVar"y"])])  
                                         end  
   
                                 | mkLoopNest (param::rest,count,nDims) = let  
                                         val body = mkLoopNest (rest, count + 1,nDims)  
299                                     in                                     in
300                                          CL.mkFor(                                          CL.mkFor(
301                                                          [(CL.intTy, param, CL.mkInt(0,CL.intTy))],                          [(ty, param, lo)],
302                                                  CL.mkBinOp(CL.mkVar param, CL.#<, CL.mkSubscript(CL.mkVar "sizes",CL.mkInt(count,CL.intTy))),                          CL.mkBinOp(CL.E_Var param, CL.#<=, hi),
303                                                  [CL.mkPostOp(CL.mkVar param, CL.^++)],                          [CL.mkPostOp(CL.E_Var param, CL.^++)],
304                                                  body)                                                  body)
305                                     end                                     end
306                  val iterCode = [
307                          CL.mkComment["initially"],
308                          CL.mkDecl(CL.uint32, indexVar, SOME(CL.I_Exp(CL.E_Int(0, CL.uint32)))),
309                          mkLoopNest iters
310                        ]
311                  val body = CL.mkBlock(
312                        iterPrefix @
313                        allocCode @
314                        iterCode @
315                        [CL.mkReturn(SOME(CL.E_Var "wrld"))])
316                  val initFn = CL.D_Func([], worldTy, N.initially, [CL.PARAM([], globPtrTy, RN.globalsVarName)], body)
317                          in                          in
318                                  [mkLoopNest ((loopParams nDims),0,nDims)]                  initially := initFn
                         end  
                 in  
                         CL.D_Func(["static"], CL.voidTy, RN.strandInitSetup, params,CL.mkBlock(body))  
319                  end                  end
320          fun genStrandPrint (Strand{name, tyName, state, output, code,...},nDims) = let  
321          (***** OUTPUT *****)
322            fun genStrandPrint (Strand{name, tyName, state, output, code,...}) = let
323              (* the print function *)              (* the print function *)
324                val prFnName = concat[name, "_print"]                val prFnName = concat[name, "_print"]
   
325                val prFn = let                val prFn = let
326                      val params = [                      val params = [
327                            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"),  
328                            CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "self")                            CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "self")
329                          ]                          ]
   
330                     val SOME(ty, x) = !output                     val SOME(ty, x) = !output
331                     val outState = if nDims = 1 then                      val outState = CL.mkIndirect(CL.mkVar "self", x)
                           CL.mkSelect(CL.mkSubscript(CL.mkVar "self",CL.mkVar "x"), x)  
                         else if nDims = 2 then  
                                 CL.mkSelect(CL.mkSubscript(CL.mkVar "self",  
                                    CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")), x)  
   
                         else CL.mkSelect(CL.mkVar "self",x)  
   
332                      val prArgs = (case ty                      val prArgs = (case ty
333                             of Ty.IVecTy 1 => [CL.mkStr(!RN.gIntFormat ^ "\n"), outState]                             of Ty.IVecTy 1 => [CL.E_Str(!N.gIntFormat ^ "\n"), outState]
334                              | Ty.IVecTy d => let                              | Ty.IVecTy d => let
335                                  val fmt = CL.mkStr(                                  val fmt = CL.mkStr(
336                                        String.concatWith " " (List.tabulate(d, fn _ => !RN.gIntFormat))                                        String.concatWith " " (List.tabulate(d, fn _ => !N.gIntFormat))
337                                        ^ "\n")                                        ^ "\n")
338                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))                                  val args = List.tabulate (d, fn i => ToC.ivecIndex(outState, d, i))
339                                  in                                  in
340                                    fmt :: args                                    fmt :: args
341                                  end                                  end
# Line 305  Line 344 
344                                  val fmt = CL.mkStr(                                  val fmt = CL.mkStr(
345                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))
346                                        ^ "\n")                                        ^ "\n")
347                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))                                  val args = List.tabulate (d, fn i => ToC.vecIndex(outState, d, i))
348                                  in                                  in
349                                    fmt :: args                                    fmt :: args
350                                  end                                  end
351                              | _ => raise Fail("genStrand: unsupported output type " ^ Ty.toString ty)                              | _ => raise Fail("genStrand: unsupported output type " ^ Ty.toString ty)
352                            (* 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.mkInt(0,CL.intTy))],  
                                                 CL.mkBinOp(CL.mkVar param, CL.#<, CL.mkSubscript(CL.mkVar "sizes",CL.mkInt(count,CL.intTy))),  
                                                 [CL.mkPostOp(CL.mkVar param, CL.^++)],  
                                                 body)  
                                    end  
353                          in                          in
354                                  [mkLoopNest ((loopParams nDims),0)]                        CL.D_Func(["static"], CL.voidTy, prFnName, params,
355                          end                          CL.mkCall("fprintf", CL.mkVar "outS" :: prArgs))
   
                     in  
                       CL.D_Func(["static"], CL.voidTy, prFnName, params,CL.mkBlock(body))  
356                      end                      end
357                in                in
358                                   prFn                                   prFn
359                end                end
360          fun genStrandTyDef (Strand{tyName, state,...}) =  
361            fun genStrandTyDef (targetTy, Strand{tyName, state,...}) =
362              (* the type declaration for the strand's state struct *)              (* the type declaration for the strand's state struct *)
363                CL.D_StructDef(                CL.D_StructDef(
364                        List.rev (List.map (fn ToCL.V(ty, x) => (ty, x)) (!state)),                  List.rev (List.map (fn x => (targetTy x, #var x)) (!state)),
365                        tyName)                        tyName)
366    
   
367          (* generates the load kernel function *)          (* generates the load kernel function *)
368  (* 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;",  
                                                 "}"])  
369  (* generates the opencl buffers for the image data *)  (* generates the opencl buffers for the image data *)
370          fun getGlobalDataBuffers(globals,count,contextVar,errVar) = let          fun getGlobalDataBuffers (globals,contextVar,errVar) = let
371                  val globalBufferDecl =  CL.mkDecl(CL.clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)                val globalBufferDecl =  CL.mkDecl(clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)
372                  val globalBuffer = CL.mkAssign(CL.mkVar(concat[RN.globalsVarName,"_cl"]), CL.mkApply("clCreateBuffer",                val globalBuffer = CL.mkAssign(CL.mkVar(concat[RN.globalsVarName,"_cl"]),
373                                                                  [CL.mkVar contextVar,                      CL.mkApply("clCreateBuffer", [
374                            CL.mkVar contextVar,
375                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",
376                                                                  CL.mkApply("sizeof",[CL.mkVar RN.globalsTy]),                                                                  CL.mkApply("sizeof",[CL.mkVar RN.globalsTy]),
377                                                                  CL.mkVar RN.globalsVarName,                                                                  CL.mkVar RN.globalsVarName,
378                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))                          CL.mkUnOp(CL.%&,CL.mkVar errVar)
379                          ]))
380          fun genDataBuffers([],_,_,_) = []                fun genDataBuffers ([],_,_) = []
381            | genDataBuffers((var,nDims)::globals,count,contextVar,errVar) = let                  | genDataBuffers ((var,nDims)::globals, contextVar, errVar) = let
382                        val hostVar = CL.mkIndirect(CL.mkVar RN.globalsVarName, var)
383          (* FIXME: use CL constructors to  build expressions (not strings) *)          (* FIXME: use CL constructors to  build expressions (not strings) *)
384                    val size = if nDims = 1 then                      fun sizeExp i = CL.mkSubscript(CL.mkIndirect(hostVar, "size"), CL.mkInt i)
385                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,                      val size = CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*, sizeExp 0)
386                                           CL.mkIndirect(CL.mkVar var, "size[0]"))                      val size = if (nDims > 1)
387                                          else if nDims = 2 then                            then CL.mkBinOp(size, CL.#*, sizeExp 1)
388                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,                            else size
389                                            CL.mkIndirect(CL.mkVar var, concat["size[0]", " * ", var, "->size[1]"]))                      val size = if (nDims > 2)
390                                          else                            then CL.mkBinOp(size, CL.#*, sizeExp 2)
391                                           CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,                            else size
392                                            CL.mkIndirect(CL.mkVar var,concat["size[0]", " * ", var, "->size[1] * ", var, "->size[2]"]))                      in
393                          CL.mkDecl(clMemoryTy, RN.addBufferSuffix var ,NONE)::
394                   in                        CL.mkDecl(clMemoryTy, RN.addBufferSuffixData var ,NONE)::
395                     CL.mkDecl(CL.clMemoryTy,RN.addBufferSuffix var ,NONE)::                        CL.mkAssign(CL.mkVar(RN.addBufferSuffix var),
396                     CL.mkDecl(CL.clMemoryTy,RN.addBufferSuffixData var ,NONE)::                          CL.mkApply("clCreateBuffer", [
397                     CL.mkAssign(CL.mkVar(RN.addBufferSuffix var), CL.mkApply("clCreateBuffer",                              CL.mkVar contextVar,
                                                                 [CL.mkVar contextVar,  
398                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",
399                                                                  CL.mkApply("sizeof",[CL.mkVar (RN.imageTy nDims)]),                                                                  CL.mkApply("sizeof",[CL.mkVar (RN.imageTy nDims)]),
400                                                                  CL.mkVar var,                              hostVar,
401                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)])) ::                              CL.mkUnOp(CL.%&,CL.mkVar errVar)
402                          CL.mkAssign(CL.mkVar(RN.addBufferSuffixData var), CL.mkApply("clCreateBuffer",                            ])) ::
403                                                                  [CL.mkVar contextVar,                        CL.mkAssign(CL.mkVar(RN.addBufferSuffixData var),
404                            CL.mkApply("clCreateBuffer", [
405                                CL.mkVar contextVar,
406                                                                   CL.mkVar "CL_MEM_COPY_HOST_PTR",                                                                   CL.mkVar "CL_MEM_COPY_HOST_PTR",
407                                                                  size,                                                                  size,
408                                                                  CL.mkIndirect(CL.mkVar var,"data"),                              CL.mkIndirect(hostVar, "data"),
409                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)])):: genDataBuffers(globals,count + 2,contextVar,errVar)                              CL.mkUnOp(CL.%&,CL.mkVar errVar)
410                              ])) :: genDataBuffers(globals,contextVar,errVar)
411                  end                  end
412          in          in
413                  [globalBufferDecl] @ [globalBuffer] @ genDataBuffers(globals,count + 2,contextVar,errVar)                  globalBufferDecl :: globalBuffer :: genDataBuffers(globals,contextVar,errVar)
414          end          end
415    
   
416  (* generates the kernel arguments for the image data *)  (* generates the kernel arguments for the image data *)
417          fun genGlobalArguments(globals,count,kernelVar,errVar) = let          fun genGlobalArguments(globals,count,kernelVar,errVar) = let
418          val globalArgument = CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,CL.mkApply("clSetKernelArg",                val globalArgument = CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,
419                        CL.mkApply("clSetKernelArg",
420                                                                  [CL.mkVar kernelVar,                                                                  [CL.mkVar kernelVar,
421                                                                   CL.mkInt(count,CL.intTy),                         CL.mkPostOp(CL.E_Var count, CL.^++),
422                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
423                                                                   CL.mkUnOp(CL.%&,CL.mkVar(concat[RN.globalsVarName,"_cl"]))])))                                                                   CL.mkUnOp(CL.%&,CL.mkVar(concat[RN.globalsVarName,"_cl"]))])))
   
424          fun genDataArguments([],_,_,_) = []          fun genDataArguments([],_,_,_) = []
425            | genDataArguments((var,nDims)::globals,count,kernelVar,errVar) =            | genDataArguments((var,nDims)::globals,count,kernelVar,errVar) =
426                        CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,
427                  CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=, CL.mkApply("clSetKernelArg",                        CL.mkApply("clSetKernelArg",
428                                                                  [CL.mkVar kernelVar,                                                                  [CL.mkVar kernelVar,
429                                                                   CL.mkInt(count,CL.intTy),                           CL.mkPostOp(CL.E_Var count, CL.^++),
430                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
431                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffix var))])))::                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffix var))])))::
432                        CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,
433                          CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,CL.mkApply("clSetKernelArg",                        CL.mkApply("clSetKernelArg",
434                                                                  [CL.mkVar kernelVar,                                                                  [CL.mkVar kernelVar,
435                                                                   CL.mkInt((count + 1),CL.intTy),                           CL.mkPostOp(CL.E_Var count, CL.^++),
436                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
437                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffixData var))]))):: genDataArguments (globals, count + 2,kernelVar,errVar)                           CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffixData var))]))) ::
438                        genDataArguments (globals,count,kernelVar,errVar)
439          in          in
440                    globalArgument :: genDataArguments(globals, count, kernelVar, errVar)
                 [globalArgument] @ genDataArguments(globals,count + 1,kernelVar,errVar)  
   
441          end          end
442    
443          (* generates the main function of host code *)        (* generates the globals buffers and arguments function *)
444          fun genHostMain() = let          fun genGlobalBuffersArgs (imgGlobals) = let
               val setupCall = [CL.mkCall(RN.setupFName,[CL.mkVar 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.mkVar RN.globalsTy])]))))  
               val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.mkVar RN.globalsVarName])  
               val returnStm = [CL.mkReturn(SOME(CL.mkInt(0,CL.intTy)))]  
               val params = [  
                      CL.PARAM([],CL.intTy, "argc"),  
                      CL.PARAM([],CL.charArrayPtr,"argv")  
                    ]  
               val body = CL.mkBlock([globalsDecl] @ [initGlobalsCall]  @ setupCall @ returnStm)  
               in  
                 CL.D_Func([],CL.intTy,"main",params,body)  
               end  
   
       (* generates the host-side setup function *)  
         fun genHostSetupFunc (strand as Strand{name,tyName,...}, filename, nDims, initially, imgGlobals) = let  
445              (* 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"  
446                val errVar = "err"                val errVar = "err"
447                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.mkVar errVar, CL.#==, CL.mkVar "CL_SUCCESS")])  
448                val params = [                val params = [
449                        CL.PARAM([],CL.T_Named("cl_device_id"), deviceVar)                        CL.PARAM([], globPtrTy, RN.globalsVarName),
450                      ]                        CL.PARAM([],CL.T_Named("cl_context"), "context"),
451                val declarations = [                        CL.PARAM([],CL.T_Named("cl_kernel"), "kernel"),
452                      CL.mkDecl(CL.clProgramTy, programVar, NONE),                        CL.PARAM([],CL.T_Named("int"), "argStart")
                     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.mkInt(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.mkStr filename))),  
 (* FIXME:  use Paths.diderotInclude *)  
                     CL.mkDecl(CL.charPtr, headerFNVar,SOME(CL.I_Exp(CL.mkStr "../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.mkInt(~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.mkInt(~1,CL.intTy))))  
453                  ]                  ]
454              (* Setup Global Variables *)                val clGlobalBuffers = getGlobalDataBuffers(!imgGlobals, "context", errVar)
455                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.mkVar RN.globalsTy])]))))  
               val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.mkVar RN.globalsVarName])  
   
                 (* Retrieve the platforms  
                 val platformStm = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetPlatformIDs",  
                                                   [CL.mkInt(10,CL.intTy),  
                                                    CL.mkVar platformsVar,  
                                                    CL.mkUnOp(CL.%&,CL.mkVar numPlatformsVar)])),  
                                                    assertStm]  
   
                 val devicesStm = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetDeviceIDs",  
                                                   [CL.mkSubscript(CL.mkVar platformsVar,CL.mkInt(0,CL.intTy)),  
                                                    CL.mkVar "CL_DEVICE_TYPE_GPU",  
                                                    CL.mkInt(1,CL.intTy),  
                                                    CL.mkUnOp(CL.%&,CL.mkVar deviceVar),  
                                                    CL.mkUnOp(CL.%&,CL.mkVar numDevicesVar)])),  
                                                    assertStm] *)  
   
                 (* Create Context *)  
                 val contextStm = [CL.mkAssign(CL.mkVar contextVar, CL.mkApply("clCreateContext",  
                                                   [CL.mkInt(0,CL.intTy),  
                                                   CL.mkInt(1,CL.intTy),  
                                                   CL.mkUnOp(CL.%&,CL.mkVar deviceVar),  
                                                   CL.mkVar "NULL",  
                                                   CL.mkVar "NULL",  
                                                   CL.mkUnOp(CL.%&,CL.mkVar errVar)])),  
                                                   assertStm]  
   
                 (* Create Command Queue *)  
                 val commandStm = [CL.mkAssign(CL.mkVar cmdVar, CL.mkApply("clCreateCommandQueue",  
                                                   [CL.mkVar contextVar,  
                                                   CL.mkVar deviceVar,  
                                                   CL.mkInt(0,CL.intTy),  
                                                   CL.mkUnOp(CL.%&,CL.mkVar errVar)])),  
                                                   assertStm]  
   
   
                 (*Create Program/Build/Kernel with Source statement *)  
                 val createProgStm = CL.mkAssign(CL.mkVar programVar, CL.mkApply("clCreateProgramWithSource",  
                                                                                                                 [CL.mkVar contextVar,  
                                                                                                                  CL.mkInt(2,CL.intTy),  
                                                                                                                  CL.mkCast(CL.T_Ptr(CL.T_Named("const char *")),CL.mkUnOp(CL.%&,CL.mkVar sourcesVar)),  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))  
   
                 (* FIXME: Remove after testing purposes, Build Log for OpenCL*)  
                 val buildLog = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clBuildProgram",  
                                                                                                                 [CL.mkVar programVar,  
                                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkVar "NULL"])),  
                                           CL.mkDecl(CL.charPtr, "build", NONE),  
                                           CL.mkDecl(CL.T_Named("size_t"),"ret_val_size",NONE),  
                                            CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetProgramBuildInfo",  
                                                                                                                 [CL.mkVar programVar,  
                                                                                                                 CL.mkVar deviceVar,  
                                                                                                                  CL.mkVar "CL_PROGRAM_BUILD_LOG",  
                                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkUnOp(CL.%&,CL.mkVar "ret_val_size")])),  
                                           CL.mkAssign(CL.mkVar "build", CL.mkApply("malloc", [CL.mkVar "ret_val_size"])),  
                                                 CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetProgramBuildInfo",  
                                                                                                                 [CL.mkVar programVar,  
                                                                                                                 CL.mkVar deviceVar,  
                                                                                                                  CL.mkVar "CL_PROGRAM_BUILD_LOG",  
                                                                                                                  CL.mkVar "ret_val_size",  
                                                                                                                  CL.mkVar "build",  
                                                                                                                  CL.mkVar "NULL"])),  
                                                 CL.mkAssign(CL.mkSubscript(CL.mkVar "build",CL.mkVar "ret_val_size"),CL.mkVar ("'\\" ^ "0'")),  
                                                 CL.mkCall("printf",[CL.mkStr ( "Build Log:" ^ "\n" ^ "%s" ^ "\n"), CL.mkVar "build"])]  
   
   
   
   
                 val createKernel = CL.mkAssign(CL.mkVar kernelVar, CL.mkApply("clCreateKernel",  
                                                                                                                 [CL.mkVar programVar,  
                                                                                                                  CL.mkStr RN.kernelFuncName,  
                                                                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))  
   
   
                 val create_build_stms = [createProgStm,assertStm] @ buildLog @ [assertStm,createKernel,assertStm]  
   
   
   
                 (* Create Memory Buffers for Strand States and Globals *)  
                 val strandSize = CL.mkAssign(CL.mkVar stateSizeVar,CL.mkBinOp(CL.mkApply("sizeof",  
                                                                         [CL.mkVar tyName]), CL.#*,CL.mkVar numStrandsVar))  
   
                 val clStrandObjects = [CL.mkAssign(CL.mkVar clInstateVar, CL.mkApply("clCreateBuffer",  
                                                                 [CL.mkVar contextVar,  
                                                                 CL.mkVar "CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR",  
                                                                 CL.mkVar stateSizeVar,  
                                                                 CL.mkVar "NULL",  
                                                                 CL.mkUnOp(CL.%&,CL.mkVar errVar)])),  
                                                          CL.mkAssign(CL.mkVar clOutStateVar, CL.mkApply("clCreateBuffer",  
                                                                 [CL.mkVar contextVar,  
                                                                 CL.mkVar "CL_MEM_READ_WRITE",  
                                                                 CL.mkVar stateSizeVar,  
                                                                 CL.mkVar "NULL",  
                                                                 CL.mkUnOp(CL.%&,CL.mkVar errVar)]))]  
   
   
                 (* Setup up selfOut variable *)  
                 val strandsArrays = [CL.mkAssign(CL.mkVar outStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.mkVar numStrandsVar,  
                                                                         CL.#*, CL.mkApply("sizeof",[CL.mkVar tyName]))])),  
                                                                 CL.mkAssign(CL.mkVar inStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.mkVar numStrandsVar,  
                                                                         CL.#*, CL.mkApply("sizeof",[CL.mkVar tyName]))]))]  
   
   
                 (* Initialize Width Parameter *)  
                 val widthDel = if nDims = 2 then  
                           CL.mkAssign(CL.mkVar "width",CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1, CL.intTy)))  
                    else  
                           CL.mkAssign(CL.mkVar "width",CL.mkInt(0,CL.intTy))  
   
   
                 val strands_init = CL.mkCall(RN.strandInitSetup, [  
                         CL.mkVar "size", CL.mkVar "width", CL.mkVar inStateVar  
                       ])  
   
             val clGlobalBuffers = getGlobalDataBuffers(!imgGlobals,3,contextVar,errVar)  
   
   
                 (* Load the Kernel and Header Files *)  
                 val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(1,CL.intTy)),  
                                                                           CL.mkApply(RN.clLoaderFN, [CL.mkVar clFNVar])),  
            CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(0,CL.intTy)),  
                                                                           CL.mkApply(RN.clLoaderFN, [CL.mkVar headerFNVar]))]  
   
                 (* val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(1,CL.intTy)),  
                                                                           CL.mkApply(RN.clLoaderFN, [CL.mkVar clFNVar]))] *)  
   
   
                 (* Created Enqueue Statements *)  
 (* FIXME: simplify this code by function abstraction *)  
         val enqueueStm = if nDims = 1  
                         then [CL.mkAssign(CL.mkVar errVar,  
                                                           CL.mkApply("clEnqueueNDRangeKernel",  
                                                                                                 [CL.mkVar cmdVar,  
                                                                                                  CL.mkVar kernelVar,  
                                                                                                  CL.mkInt(1,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar globalVar,  
                                                                                                  CL.mkVar localVar,  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]  
                         else if nDims = 2  then  
                          [CL.mkAssign(CL.mkVar errVar,  
                                                         CL.mkApply("clEnqueueNDRangeKernel",  
                                                                                                 [CL.mkVar cmdVar,  
                                                                                                  CL.mkVar kernelVar,  
                                                                                                  CL.mkInt(2,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar globalVar,  
                                                                                                  CL.mkVar localVar,  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]  
                         else  
                           [CL.mkAssign(CL.mkVar errVar,  
                                                         CL.mkApply("clEnqueueNDRangeKernel",  
                                                                                                 [CL.mkVar cmdVar,  
                                                                                                  CL.mkVar kernelVar,  
                                                                                                  CL.mkInt(3,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar globalVar,  
                                                                                                  CL.mkVar localVar,  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]  
   
   
   
                 (* Setup Global and Local variables *)  
   
                 val globalAndlocalStms = if nDims = 1 then  
                         [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),  
                          CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),  
                                                                   CL.mkVar "16")]  
   
   
                 else if nDims = 2 then  
                         [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(1,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),  
                                                                   CL.mkVar "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(1,CL.intTy)),  
                                                                   CL.mkVar "16")]  
   
                 else  
                         [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(1,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(2,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(2,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),  
                                                                   CL.mkVar "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(1,CL.intTy)),  
                                                                   CL.mkVar "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(2,CL.intTy)),  
                                                                   CL.mkVar "16")]  
   
   
   
                 (* Setup Kernel arguments *)  
                 val kernelArguments = [CL.mkAssign(CL.mkVar errVar,CL.mkApply("clSetKernelArg",  
                                                                 [CL.mkVar kernelVar,  
                                                                  CL.mkInt(0,CL.intTy),  
                                                                  CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),  
                                                                  CL.mkUnOp(CL.%&,CL.mkVar clInstateVar)])),  
                                                             CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar, CL.|=,CL.mkApply("clSetKernelArg",  
                                                                 [CL.mkVar kernelVar,  
                                                                  CL.mkInt(1,CL.intTy),  
                                                                  CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),  
                                                                  CL.mkUnOp(CL.%&,CL.mkVar clOutStateVar)]))),  
                                                                   CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar, CL.|=,CL.mkApply("clSetKernelArg",  
                                                                 [CL.mkVar kernelVar,  
                                                                  CL.mkInt(2,CL.intTy),  
                                                                  CL.mkApply("sizeof",[CL.mkVar "int"]),  
                                                                  CL.mkUnOp(CL.%&,CL.mkVar "width")])))]  
   
            val clGlobalArguments = genGlobalArguments(!imgGlobals,3,kernelVar,errVar) @ [assertStm]  
   
                 (* Retrieve output *)  
                 val outputStm = CL.mkAssign(CL.mkVar errVar,  
                                                         CL.mkApply("clEnqueueReadBuffer",  
                                                                                                 [CL.mkVar cmdVar,  
                                                                                                  CL.mkVar clOutStateVar,  
                                                                                                  CL.mkVar "CL_TRUE",  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar stateSizeVar,  
                                                                                                  CL.mkVar outStateVar,  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar "NULL"]))  
   
                 (* Free all the objects *)  
                 val freeStms = [CL.mkCall("clReleaseKernel",[CL.mkVar kernelVar]),  
                                                 CL.mkCall("clReleaseProgram",[CL.mkVar programVar ]),  
                                                 CL.mkCall("clReleaseCommandQueue",[CL.mkVar cmdVar]),  
                                                 CL.mkCall("clReleaseContext",[CL.mkVar contextVar]),  
                                                 CL.mkCall("clReleaseMemObject",[CL.mkVar clInstateVar]),  
                                                 CL.mkCall("clReleaseMemObject",[CL.mkVar 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.mkStr "mip.txt",  
                                                 CL.mkStr "w"])))),  
                                                 CL.mkCall(concat[name, "_print"],  
                                                                         [CL.mkVar "outS",  
                                                                          CL.mkVar "size",  
                                                                          CL.mkVar "width",  
                                                                          CL.mkVar outStateVar])]  
   
   
   
456                  (* Body put all the statments together *)                  (* Body put all the statments together *)
457                  val body =  declarations @ [globalsDecl,initGlobalsCall] (*@ platformStm @ devicesStm *) @ contextStm @ commandStm @ !initially @ [strandSize] @                val body = CL.mkDecl(clIntTy, errVar, SOME(CL.I_Exp(CL.mkInt 0)))
458                                     strandsArrays @ globalAndlocalStms @ [widthDel,strands_init]  @ clStrandObjects @ clGlobalBuffers @ sourceStms  @ create_build_stms  (*@                      :: clGlobalBuffers @ clGlobalArguments
                                    kernelArguments @ clGlobalArguments @ enqueueStm @  [outputStm] @ freeStms @ outputData *)  
   
459                  in                  in
460    (* FIXME: we ought to check the error condition! *)
461                  CL.D_Func([],CL.voidTy,RN.setupFName,params,CL.mkBlock(body))                  CL.D_Func([],CL.voidTy,RN.globalsSetupName,params,CL.mkBlock(body))
   
462                  end                  end
463    
464  (* generate the data and global parameters *)  (* generate the data and global parameters *)
465          fun genKeneralGlobalParams ((name,tyname)::rest) =          fun genKeneralGlobalParams ((name,tyname)::rest) =
466                  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]) ::
467                  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) ::
468                  CL.PARAM([], CL.T_Ptr(CL.voidTy),RN.addBufferSuffixData name) ::                  CL.PARAM([], CL.T_Ptr(CL.voidTy),RN.addBufferSuffixData name) ::
469                  genKeneralGlobalParams(rest)                genKeneralGlobalParams rest
470              | genKeneralGlobalParams [] = []
           | genKeneralGlobalParams ([]) = []  
471    
472          (*generate code for intilizing kernel global data *)          (*generate code for intilizing kernel global data *)
473          fun initKernelGlobals (globals,imgGlobals) = let          fun initKernelGlobals (globals,imgGlobals) = let
474                  fun initGlobalStruct (CL.D_Var(_, _ , name, _)::rest) =  (* FIXME: should use List.map here *)
475                                  CL.mkAssign(CL.mkVar name, CL.mkIndirect(CL.mkVar RN.globalsVarName, name)) ::                fun initGlobalStruct ({hostTy, gpuTy, var}::rest) =
476                                  initGlobalStruct(rest)                      CL.mkAssign(CL.mkVar var, CL.mkIndirect(CL.mkVar RN.globalsVarName, var)) ::
477                    | initGlobalStruct ( _::rest) = initGlobalStruct(rest)                      initGlobalStruct rest
478                    | initGlobalStruct([]) = []                  | initGlobalStruct [] = []
   
479                  fun initGlobalImages((name,tyname)::rest) =                  fun initGlobalImages((name,tyname)::rest) =
480                                  CL.mkAssign(CL.mkVar name, CL.mkVar (RN.addBufferSuffix name)) ::                                  CL.mkAssign(CL.mkVar name, CL.mkVar (RN.addBufferSuffix name)) ::
481                                  CL.mkAssign(CL.mkIndirect(CL.mkVar name,"data"),CL.mkVar (RN.addBufferSuffixData name)) ::                                  CL.mkAssign(CL.mkIndirect(CL.mkVar name,"data"),CL.mkVar (RN.addBufferSuffixData name)) ::
482                                  initGlobalImages(rest)                      initGlobalImages rest
483                    | initGlobalImages([]) = []                    | initGlobalImages [] = []
484                  in                  in
485                    initGlobalStruct(globals) @ initGlobalImages(imgGlobals)                  initGlobalStruct globals @ initGlobalImages(imgGlobals)
486                  end                  end
487    
488          (* generate the main kernel function for the .cl file *)          (* generate the main kernel function for the .cl file *)
489          fun genKernelFun(Strand{name, tyName, state, output, code,...},nDims,globals,imgGlobals) = let          fun genKernelFun (strand, nDims, globals, imgGlobals) = let
490                  val Strand{name, tyName, state, output, code,...} = strand
491                   val fName = RN.kernelFuncName;                   val fName = RN.kernelFuncName;
492                   val inState = "strand_in"                   val inState = "strand_in"
493                   val outState = "strand_out"                   val outState = "strand_out"
# Line 828  Line 497 
497                        CL.PARAM(["__global"], CL.intTy, "width")                        CL.PARAM(["__global"], CL.intTy, "width")
498                      ] @ genKeneralGlobalParams(!imgGlobals)                      ] @ genKeneralGlobalParams(!imgGlobals)
499                    val thread_ids = if nDims = 1                    val thread_ids = if nDims = 1
500                          then [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),                      then [
501                                    CL.mkAssign(CL.mkVar "x",CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(0,CL.intTy)]))]                          CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt 0))),
502                          else                          CL.mkAssign(CL.mkVar "x",CL.mkApply(RN.getGlobalThreadId,[CL.mkInt 0]))
503                                  [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),                        ]
504                                   CL.mkDecl(CL.intTy, "y", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),                      else [
505                                    CL.mkAssign(CL.mkVar "x",  CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(0,CL.intTy)])),                          CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt 0))),
506                                    CL.mkAssign(CL.mkVar "y",CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(1,CL.intTy)]))]                          CL.mkDecl(CL.intTy, "y", SOME(CL.I_Exp(CL.mkInt 0))),
507                            CL.mkAssign(CL.mkVar "x",  CL.mkApply(RN.getGlobalThreadId,[CL.mkInt 0])),
508                    val strandDecl = [CL.mkDecl(CL.T_Named tyName, inState, NONE),                          CL.mkAssign(CL.mkVar "y",CL.mkApply(RN.getGlobalThreadId,[CL.mkInt 1]))
509                          ]
510                  val strandDecl = [
511                        CL.mkDecl(CL.T_Named tyName, inState, NONE),
512                                                          CL.mkDecl(CL.T_Named tyName, outState,NONE)]                                                          CL.mkDecl(CL.T_Named tyName, outState,NONE)]
513                    val strandObjects  = if nDims = 1                    val strandObjects  = if nDims = 1
514                          then [CL.mkAssign(CL.mkSubscript(CL.mkVar "selfIn",CL.mkStr "x"),                      then [
515                                                                           CL.mkVar inState),                          CL.mkAssign( CL.mkVar inState, CL.mkSubscript(CL.mkVar "selfIn", CL.mkStr "x")),
516                                    CL.mkAssign(CL.mkSubscript(CL.mkVar "selfOut",CL.mkStr "x"),                          CL.mkAssign(CL.mkVar outState,CL.mkSubscript(CL.mkVar "selfOut", CL.mkStr "x"))
517                                                                           CL.mkVar outState)]                        ]
518                          else let                          else let
519                                  val index = CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")                                  val index = CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")
520                                  in                        in [
521                                          [CL.mkAssign(CL.mkSubscript(CL.mkVar "selfIn",index),                          CL.mkAssign(CL.mkVar inState, CL.mkSubscript(CL.mkVar "selfIn",index)),
522                                                                          CL.mkVar inState),                          CL.mkAssign(CL.mkVar outState,CL.mkSubscript(CL.mkVar "selfOut",index))
523                                           CL.mkAssign(CL.mkSubscript(CL.mkVar "selfOut",index),                        ] end
524                                                                          CL.mkVar outState)]                val status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.mkInt 0)))
                                 end  
   
   
                   val status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy))))  
525                    val local_vars =  thread_ids @ initKernelGlobals(!globals,!imgGlobals)  @ strandDecl @ strandObjects @ [status]                    val local_vars =  thread_ids @ initKernelGlobals(!globals,!imgGlobals)  @ strandDecl @ strandObjects @ [status]
526                    val while_exp = CL.mkBinOp(CL.mkBinOp(CL.mkVar "status",CL.#!=, CL.mkVar RN.kStabilize),CL.#||,CL.mkBinOp(CL.mkVar "status", CL.#!=, CL.mkVar RN.kDie))                val while_exp = CL.mkBinOp(
527                    val while_body = [CL.mkAssign(CL.mkVar "status", CL.mkApply(RN.strandUpdate name,[ CL.mkUnOp(CL.%&,CL.mkVar inState), CL.mkUnOp(CL.%&,CL.mkVar outState)])),                      CL.mkBinOp(CL.mkVar "status",CL.#!=, CL.mkVar RN.kStabilize),
528                                                          CL.mkCall(RN.strandStabilize name,[ CL.mkUnOp(CL.%&,CL.mkVar inState),  CL.mkUnOp(CL.%&,CL.mkVar outState)])]                      CL.#||,
529                        CL.mkBinOp(CL.mkVar "status", CL.#!=, CL.mkVar RN.kDie))
530                    val whileBlock = [CL.mkWhile(while_exp,CL.mkBlock while_body)]                val whileBody = CL.mkBlock [
531                          CL.mkAssign(CL.mkVar "status",
532                            CL.mkApply(RN.strandUpdate name,
533                              [CL.mkUnOp(CL.%&,CL.mkVar inState), CL.mkUnOp(CL.%&,CL.mkVar outState)])),
534                          CL.mkCall(RN.strandStabilize name,
535                            [CL.mkUnOp(CL.%&,CL.mkVar inState), CL.mkUnOp(CL.%&,CL.mkVar outState)])
536                        ]
537                  val whileBlock = [CL.mkWhile(while_exp, whileBody)]
538                    val body = CL.mkBlock(local_vars  @ whileBlock)                    val body = CL.mkBlock(local_vars  @ whileBlock)
539                  in                  in
540                     CL.D_Func(["__kernel"], CL.voidTy, fName, params, body)                     CL.D_Func(["__kernel"], CL.voidTy, fName, params, body)
541                  end                  end
542          (* generate a global structure from the globals *)          (* generate a global structure from the globals *)
543          fun genGlobalStruct(globals) = let          fun genGlobalStruct (targetTy, globals) = let
544                   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
545                     | getGlobals([]) = []                in
546                     | getGlobals(_::rest) = getGlobals(rest)                  CL.D_StructDef(globs, RN.globalsTy)
547                  end
548            fun genGlobals (declFn, targetTy, globals) = let
549                  fun doVar (x : mirror_var) = declFn (CL.D_Var([], targetTy x, #var x, NONE))
550                  in
551                    List.app doVar globals
552                  end
553    
554            fun genStrandDesc (Strand{name, output, ...}) = let
555                (* the strand's descriptor object *)
556                  val descI = let
557                        fun fnPtr (ty, f) = CL.I_Exp(CL.mkCast(CL.T_Named ty, CL.mkVar f))
558                        val SOME(outTy, _) = !output
559                   in                   in
560                          CL.D_StructDef(getGlobals(globals),RN.globalsTy)                        CL.I_Struct[
561                              ("name", CL.I_Exp(CL.mkStr name)),
562                              ("stateSzb", CL.I_Exp(CL.mkSizeof(CL.T_Named(N.strandTy name)))),
563    (*
564                              ("outputSzb", CL.I_Exp(CL.mkSizeof(ToC.trTy outTy))),
565    *)
566                              ("update", fnPtr("update_method_t", "0")),
567                              ("print", fnPtr("print_method_t", name ^ "_print"))
568                            ]
569                        end
570                  val desc = CL.D_Var([], CL.T_Named N.strandDescTy, N.strandDesc name, SOME descI)
571                  in
572                    desc
573                    end                    end
574    
575        (* generate the table of strand descriptors *)        (* generate the table of strand descriptors *)
576          fun genStrandTable (ppStrm, strands) = let          fun genStrandTable (declFn, strands) = let
577                val nStrands = length strands                val nStrands = length strands
578                fun genInit (Strand{name, ...}) = CL.I_Exp(CL.mkUnOp(CL.%&, CL.mkVar(RN.strandDesc name)))                fun genInit (Strand{name, ...}) = CL.I_Exp(CL.mkUnOp(CL.%&, CL.E_Var(N.strandDesc name)))
579                fun genInits (_, []) = []                fun genInits (_, []) = []
580                  | 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)  
581                in                in
582                  ppDecl (CL.D_Var([], CL.int32, RN.numStrands,                  declFn (CL.D_Var([], CL.int32, N.numStrands,
583                    SOME(CL.I_Exp(CL.mkInt(IntInf.fromInt nStrands, CL.int32)))));                    SOME(CL.I_Exp(CL.E_Int(IntInf.fromInt nStrands, CL.int32)))));
584                  ppDecl (CL.D_Var([],                  declFn (CL.D_Var([],
585                    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),
586                    RN.strands,                    N.strands,
587                    SOME(CL.I_Array(genInits (0, strands)))))                    SOME(CL.I_Array(genInits (0, strands)))))
588                end                end
589    
590            fun genSrc (baseName, prog) = let
591          fun genSrc (baseName, Prog{double,globals, topDecls, strands, initially,imgGlobals,numDims,...}) = let                val Prog{double, globals, topDecls, strands, initially, imgGlobals, numDims, ...} = prog
592                val clFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "cl"}                val clFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "cl"}
593                val cFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "c"}                val cFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "c"}
594                val clOutS = TextIO.openOut clFileName                val clOutS = TextIO.openOut clFileName
595                val cOutS = TextIO.openOut cFileName                val cOutS = TextIO.openOut cFileName
596  (* FIXME: need to use PrintAsC and PrintAsCL *)                val clppStrm = PrintAsCL.new clOutS
               val clppStrm = PrintAsC.new clOutS  
597                val cppStrm = PrintAsC.new cOutS                val cppStrm = PrintAsC.new cOutS
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 910  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));                  genGlobals (clppDecl, #gpuTy, !globals);
612                  clppDecl (genGlobalStruct (!globals));                  clppDecl (genGlobalStruct (#gpuTy, !globals));
613                  clppDecl (genStrandTyDef strand);                  clppDecl (genStrandTyDef(#gpuTy, strand));
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 925  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 name))));
626                  cppDecl (genStrandTyDef strand);  (* FIXME: I don't think that the following is necessary, since we have the global struct. [jhr]
627                    genGlobals (cppDecl, #hostTy, !globals);
628    *)
629                    cppDecl (genGlobalStruct (#hostTy, !globals));
630                    cppDecl (genStrandTyDef (#hostTy, strand));
631                  cppDecl  (!init_code);                  cppDecl  (!init_code);
632                  cppDecl (genStrandInit(strand,!numDims));                  cppDecl (genStrandPrint strand);
                 cppDecl (genStrandPrint(strand,!numDims));  
                 (* cppDecl (genKernelLoader());*)  
633                  List.app cppDecl (List.rev (!topDecls));                  List.app cppDecl (List.rev (!topDecls));
634                  cppDecl (genHostSetupFunc (strand, clFileName, !numDims, initially, imgGlobals));                  cppDecl (genGlobalBuffersArgs imgGlobals);
635                    List.app (fn strand => cppDecl (genStrandDesc strand)) strands;
636                    genStrandTable (cppDecl, strands);
637                    cppDecl (!initially);
638                  PrintAsC.close cppStrm;                  PrintAsC.close cppStrm;
639                  PrintAsC.close clppStrm;                  PrintAsCL.close clppStrm;
640                  TextIO.closeOut cOutS;                  TextIO.closeOut cOutS;
641                  TextIO.closeOut clOutS                  TextIO.closeOut clOutS
642                end                end
# Line 997  Line 698 
698          fun init (Strand{name, tyName, code,init_code, ...}, params, init) = let          fun init (Strand{name, tyName, code,init_code, ...}, params, init) = let
699                val fName = RN.strandInit name                val fName = RN.strandInit name
700                val params =                val params =
701                        CL.PARAM([], globPtrTy, RN.globalsVarName) ::
702                      CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfOut") ::                      CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfOut") ::
703                        List.map (fn (ToCL.V(ty, x)) => CL.PARAM([], ty, x)) params                        List.map (fn (ToCL.V(ty, x)) => CL.PARAM([], ty, x)) params
704                val initFn = CL.D_Func([], CL.voidTy, fName, params, init)                val initFn = CL.D_Func([], CL.voidTy, fName, params, init)

Legend:
Removed from v.1278  
changed lines
  Added in v.1315

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