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 1286, Tue Jun 7 10:54:18 2011 UTC
# Line 14  Line 14 
14      structure RN = RuntimeNames      structure RN = RuntimeNames
15      structure ToCL = TreeToCL      structure ToCL = TreeToCL
16    
17      type var = ToCL.var    (* C variable translation *)
18        structure TrCVar =
19          struct
20            type env = CL.typed_var TreeIL.Var.Map.map
21            fun lookup (env, x) = (case V.Map.find (env, x)
22                   of SOME(CL.V(_, x')) => x'
23                    | NONE => raise Fail(concat["TrCVar.lookup(_, ", V.name x, ")"])
24                  (* end case *))
25          (* translate a variable that occurs in an l-value context (i.e., as the target of an assignment) *)
26            fun lvalueVar (env, x) = (case V.kind x
27                   of IL.VK_Global => CL.mkIndirect(CL.mkVar RN.globalsVarName, lookup(env, x))
28                    | IL.VK_State strand => raise Fail "unexpected strand context"
29                    | IL.VK_Local => CL.mkVar(lookup(env, x))
30                  (* end case *))
31          (* translate a variable that occurs in an r-value context *)
32            val rvalueVar = lvalueVar
33          end
34    
35        structure ToC = TreeToCFn (TrCVar)
36    
37        type var = CL.typed_var
38      type exp = CL.exp      type exp = CL.exp
39      type stm = CL.stm      type stm = CL.stm
40    
41      (* OpenCL specific types *)
42        val clProgramTy = CL.T_Named "cl_program"
43        val clKernelTy  = CL.T_Named "cl_kernel"
44        val clCmdQueueTy = CL.T_Named "cl_command_queue"
45        val clContextTy = CL.T_Named "cl_context"
46        val clDeviceIdTy = CL.T_Named "cl_device_id"
47        val clPlatformIdTy = CL.T_Named "cl_platform_id"
48        val clMemoryTy = CL.T_Named "cl_mem"
49    
50      datatype strand = Strand of {      datatype strand = Strand of {
51          name : string,          name : string,
52          tyName : string,          tyName : string,
# Line 28  Line 57 
57        }        }
58    
59      datatype program = Prog of {      datatype program = Prog of {
60            name : string,                  (* stem of source file *)
61          double : bool,                  (* true for double-precision support *)          double : bool,                  (* true for double-precision support *)
62          parallel : bool,                (* true for multithreaded (or multi-GPU) target *)          parallel : bool,                (* true for multithreaded (or multi-GPU) target *)
63          debug : bool,                   (* true for debug support in executable *)          debug : bool,                   (* true for debug support in executable *)
# Line 68  Line 98 
98    (* TreeIL to target translations *)    (* TreeIL to target translations *)
99      structure Tr =      structure Tr =
100        struct        struct
101          (* this function is used for the initially clause, so it generates OpenCL *)
102          fun fragment (ENV{info, vMap, scope}, blk) = let          fun fragment (ENV{info, vMap, scope}, blk) = let
103                val (vMap, stms) = ToCL.trFragment (vMap, blk)                val (vMap, stms) = ToCL.trFragment (vMap, blk)
104                in                in
# Line 84  Line 115 
115          fun block (ENV{vMap, scope, ...}, blk) = (case scope          fun block (ENV{vMap, scope, ...}, blk) = (case scope
116                 of StrandScope stateVars => ToCL.trBlock (vMap, saveState "StrandScope" stateVars, blk)                 of StrandScope stateVars => ToCL.trBlock (vMap, saveState "StrandScope" stateVars, blk)
117                  | MethodScope stateVars => ToCL.trBlock (vMap, saveState "MethodScope" stateVars, blk)                  | MethodScope stateVars => ToCL.trBlock (vMap, saveState "MethodScope" stateVars, blk)
118                  | _ => ToCL.trBlock (vMap, fn (_, _, stm) => [stm], blk)                  | InitiallyScope => ToCL.trBlock (vMap, fn (_, _, stm) => [stm], blk)
119                    | _ => ToC.trBlock (vMap, fn (_, _, stm) => [stm], blk)
120                (* end case *))                (* end case *))
121          fun exp (ENV{vMap, ...}, e) = ToCL.trExp(vMap, e)          fun exp (ENV{vMap, ...}, e) = ToCL.trExp(vMap, e)
122        end        end
# Line 138  Line 170 
170    (* programs *)    (* programs *)
171      structure Program =      structure Program =
172        struct        struct
173          fun new {double, parallel, debug} = (          fun new {name, double, parallel, debug} = (
174                RN.initTargetSpec double;                RN.initTargetSpec double;
175                  CNames.initTargetSpec double;
176                Prog{                Prog{
177                      name = name,
178                    double = double, parallel = parallel, debug = debug,                    double = double, parallel = parallel, debug = debug,
179                    globals = ref [],                    globals = ref [],
180                    topDecls = ref [],                    topDecls = ref [],
# Line 152  Line 186 
186                  })                  })
187        (* register the global initialization part of a program *)        (* register the global initialization part of a program *)
188            fun globalIndirects (globals,stms) = let            fun globalIndirects (globals,stms) = let
189                  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 (CL.D_Var(_,_,globalVar,_)::rest) =
190                    | getGlobals([]) = []                        CL.mkAssign(CL.mkIndirect(CL.mkVar RN.globalsVarName,globalVar),CL.mkVar globalVar)
191                    | getGlobals(_::rest) = getGlobals(rest)                          ::getGlobals rest
192                      | getGlobals [] = []
193                      | getGlobals (_::rest) = getGlobals rest
194                  in                  in
195                    stms @ getGlobals(globals)                    stms @ getGlobals globals
196                  end                  end
197    
198        (* 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 *)
# Line 169  Line 205 
205                  topDecls := inputsFn :: !topDecls                  topDecls := inputsFn :: !topDecls
206                end                end
207    
208          fun init (Prog{globals,topDecls,...}, CL.S_Block(init)) = let        (* register the global initialization part of a program *)
209                val params = [          fun init (Prog{topDecls, ...}, init) = let
210                        CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName)                val globPtrTy = CL.T_Ptr(CL.T_Named RN.globalsTy)
211                      ]                val initFn = CL.D_Func(
212                val body = CL.S_Block(globalIndirects(!globals,init))                      [], CL.voidTy, RN.initGlobals, [CL.PARAM([], globPtrTy, RN.globalsVarName)],
213                val initFn = CL.D_Func([], CL.voidTy, RN.initGlobals, params, body)                      init)
214                in                val shutdownFn = CL.D_Func(
215                  topDecls := initFn :: !topDecls                      [], CL.voidTy, RN.shutdown,
216                end                      [CL.PARAM([], CL.T_Ptr(CL.T_Named RN.worldTy), "wrld")],
217            | init (Prog{globals,topDecls,...}, init) = let                      CL.S_Block[])
               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)  
218                in                in
219                  topDecls := initFn :: !topDecls                  topDecls := shutdownFn :: initFn :: !topDecls
220                end                end
221    
222        (* create and register the initially function for a program *)        (* create and register the initially function for a program *)
# Line 208  Line 240 
240                val baseInit = mapi (fn (i, (_, e, _)) => (i, CL.I_Exp e)) iters                val baseInit = mapi (fn (i, (_, e, _)) => (i, CL.I_Exp e)) iters
241                val sizeInit = mapi                val sizeInit = mapi
242                      (fn (i, (ToCL.V(ty, _), lo, hi)) =>                      (fn (i, (ToCL.V(ty, _), lo, hi)) =>
243                          (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.mkInt(1, ty))))
244                      ) iters                      ) iters
245                    val numStrandsVar = "numStrandsVar"                    val numStrandsVar = "numStrandsVar"
246                val allocCode = iterPrefix @ [                val allocCode = iterPrefix @ [
247                        CL.mkComment["allocate initial block of strands"],                        CL.mkComment["allocate initial block of strands"],
248                        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)),
249                        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)),
250                        CL.mkDecl(CL.int32,"numDims",SOME(CL.I_Exp(CL.E_Int(IntInf.fromInt nDims, CL.int32))))                        CL.mkDecl(CL.int32,"numDims",SOME(CL.I_Exp(CL.mkInt(IntInf.fromInt nDims, CL.int32))))
251                      ]                      ]
252                val numStrandsLoopBody = CL.mkExpStm(CL.mkAssignOp(CL.E_Var numStrandsVar, CL.*=,CL.mkSubscript(CL.E_Var "size",CL.E_Var "i")))                val numStrandsLoopBody =
253                val numStrandsLoop =  CL.mkFor([(CL.intTy, "i", CL.E_Int(0,CL.intTy))],                      CL.mkExpStm(CL.mkAssignOp(CL.mkVar numStrandsVar, CL.*=,CL.mkSubscript(CL.mkVar "size",CL.mkVar "i")))
254                      CL.mkBinOp(CL.E_Var "i", CL.#<, CL.E_Var "numDims"),                val numStrandsLoop =  CL.mkFor([(CL.intTy, "i", CL.mkInt(0,CL.intTy))],
255                      [CL.mkPostOp(CL.E_Var "i", CL.^++)], numStrandsLoopBody)                      CL.mkBinOp(CL.mkVar "i", CL.#<, CL.mkVar "numDims"),
256                        [CL.mkPostOp(CL.mkVar "i", CL.^++)], numStrandsLoopBody)
257                in                in
258                  numDims := nDims;                  numDims := nDims;
259                  initially := allocCode @ [numStrandsLoop]                  initially := allocCode @ [numStrandsLoop]
# Line 234  Line 267 
267                            CL.PARAM([], CL.intTy, "width"),                            CL.PARAM([], CL.intTy, "width"),
268                            CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "strands")                            CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "strands")
269                          ]                          ]
   
270                  val body = let                  val body = let
271                              fun loopParams (3) =                      fun loopParams 3 = ["x", "y", "k"]
272                                   "x"::"y"::"k"::[]                        | loopParams 2 = ["x", "y"]
273                                | loopParams (2) =                        | loopParams 1 = ["x"]
274                                   "x"::"y"::[]                        | loopParams _ = raise Fail "genStrandInit: missing size dim"
275                                | loopParams (1) =                      fun mkLoopNest ([], _, nDims) = if nDims = 1
276                                   "x"::[]                            then CL.mkBlock [
277                                | loopParams (_) =                                CL.mkCall(RN.strandInit name, [
278                                  raise Fail("genStrandInit: missing size dim")                                  CL.mkUnOp(CL.%&,CL.mkSubscript(CL.mkVar "strands",CL.mkStr "x")),
279                             fun mkLoopNest ([],_,nDims) =  if nDims = 1 then                                                  CL.mkVar "x"])
280                                ]
                                         CL.mkBlock ([CL.mkCall(RN.strandInit name, [CL.E_UnOp(CL.%&,CL.mkSubscript(CL.E_Var "strands",CL.E_Str "x")),  
                                                 CL.E_Var "x"])])  
281                                          else let                                          else let
282                                                  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")
283                                          in                                          in
284                                                  CL.mkBlock([CL.mkCall(RN.strandInit name, [CL.E_UnOp(CL.%&,CL.mkSubscript(CL.E_Var "strands",index)),                                CL.mkBlock([CL.mkCall(RN.strandInit name, [CL.mkUnOp(CL.%&,CL.mkSubscript(CL.mkVar "strands",index)),
285                                                  CL.E_Var "x", CL.E_Var"y"])])                                CL.mkVar "x", CL.mkVar"y"])])
286                                          end                                          end
   
287                                  | mkLoopNest (param::rest,count,nDims) = let                                  | mkLoopNest (param::rest,count,nDims) = let
288                                          val body = mkLoopNest (rest, count + 1,nDims)                                          val body = mkLoopNest (rest, count + 1,nDims)
289                                     in                                     in
290                                          CL.mkFor(                                          CL.mkFor(
291                                                          [(CL.intTy, param, CL.E_Int(0,CL.intTy))],                                  [(CL.intTy, param, CL.mkInt(0,CL.intTy))],
292                                                  CL.mkBinOp(CL.E_Var param, CL.#<, CL.mkSubscript(CL.E_Var "sizes",CL.E_Int(count,CL.intTy))),                                  CL.mkBinOp(CL.mkVar param, CL.#<, CL.mkSubscript(CL.mkVar "sizes",CL.mkInt(count,CL.intTy))),
293                                                  [CL.mkPostOp(CL.E_Var param, CL.^++)],                                  [CL.mkPostOp(CL.mkVar param, CL.^++)],
294                                                  body)                                                  body)
295                                     end                                     end
296                          in                          in
# Line 270  Line 299 
299                  in                  in
300                          CL.D_Func(["static"], CL.voidTy, RN.strandInitSetup, params,CL.mkBlock(body))                          CL.D_Func(["static"], CL.voidTy, RN.strandInitSetup, params,CL.mkBlock(body))
301                  end                  end
302    
303          fun genStrandPrint (Strand{name, tyName, state, output, code,...},nDims) = let          fun genStrandPrint (Strand{name, tyName, state, output, code,...},nDims) = let
304              (* the print function *)              (* the print function *)
305                val prFnName = concat[name, "_print"]                val prFnName = concat[name, "_print"]
   
306                val prFn = let                val prFn = let
307                      val params = [                      val params = [
308                            CL.PARAM([], CL.T_Ptr(CL.T_Named "FILE"), "outS"),                            CL.PARAM([], CL.T_Ptr(CL.T_Named "FILE"), "outS"),
# Line 284  Line 313 
313    
314                     val SOME(ty, x) = !output                     val SOME(ty, x) = !output
315                     val outState = if nDims = 1 then                     val outState = if nDims = 1 then
316                            CL.mkSelect(CL.mkSubscript(CL.mkVar "self",CL.E_Var "x"), x)                            CL.mkSelect(CL.mkSubscript(CL.mkVar "self",CL.mkVar "x"), x)
317                          else if nDims = 2 then                          else if nDims = 2 then
318                                  CL.mkSelect(CL.mkSubscript(CL.mkVar "self",                                  CL.mkSelect(CL.mkSubscript(CL.mkVar "self",
319                                     CL.mkBinOp(CL.mkBinOp(CL.E_Var "x",CL.#*,CL.E_Var "width"),CL.#+,CL.E_Var "y")), x)                                     CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")), x)
320    
321                          else CL.mkSelect(CL.mkVar "self",x)                          else CL.mkSelect(CL.mkVar "self",x)
322    
323                      val prArgs = (case ty                      val prArgs = (case ty
324                             of Ty.IVecTy 1 => [CL.E_Str(!RN.gIntFormat ^ "\n"), outState]                             of Ty.IVecTy 1 => [CL.mkStr(!RN.gIntFormat ^ "\n"), outState]
325                              | Ty.IVecTy d => let                              | Ty.IVecTy d => let
326                                  val fmt = CL.E_Str(                                  val fmt = CL.mkStr(
327                                        String.concatWith " " (List.tabulate(d, fn _ => !RN.gIntFormat))                                        String.concatWith " " (List.tabulate(d, fn _ => !RN.gIntFormat))
328                                        ^ "\n")                                        ^ "\n")
329                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))
330                                  in                                  in
331                                    fmt :: args                                    fmt :: args
332                                  end                                  end
333                              | Ty.TensorTy[] => [CL.E_Str "%f\n", outState]                              | Ty.TensorTy[] => [CL.mkStr "%f\n", outState]
334                              | Ty.TensorTy[d] => let                              | Ty.TensorTy[d] => let
335                                  val fmt = CL.E_Str(                                  val fmt = CL.mkStr(
336                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))
337                                        ^ "\n")                                        ^ "\n")
338                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))
# Line 330  Line 359 
359                                          val body = mkLoopNest (rest, count + 1)                                          val body = mkLoopNest (rest, count + 1)
360                                     in                                     in
361                                                  CL.mkFor(                                                  CL.mkFor(
362                                                          [(CL.intTy, param, CL.E_Int(0,CL.intTy))],                                                          [(CL.intTy, param, CL.mkInt(0,CL.intTy))],
363                                                  CL.mkBinOp(CL.E_Var param, CL.#<, CL.mkSubscript(CL.E_Var "sizes",CL.E_Int(count,CL.intTy))),                                                  CL.mkBinOp(CL.mkVar param, CL.#<, CL.mkSubscript(CL.mkVar "sizes",CL.mkInt(count,CL.intTy))),
364                                                  [CL.mkPostOp(CL.E_Var param, CL.^++)],                                                  [CL.mkPostOp(CL.mkVar param, CL.^++)],
365                                                  body)                                                  body)
366                                     end                                     end
367                          in                          in
# Line 371  Line 400 
400                                                  "}"])                                                  "}"])
401  (* generates the opencl buffers for the image data *)  (* generates the opencl buffers for the image data *)
402          fun getGlobalDataBuffers(globals,count,contextVar,errVar) = let          fun getGlobalDataBuffers(globals,count,contextVar,errVar) = let
403                  val globalBufferDecl =  CL.mkDecl(CL.clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)                  val globalBufferDecl =  CL.mkDecl(clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)
404                  val globalBuffer = CL.mkAssign(CL.E_Var(concat[RN.globalsVarName,"_cl"]), CL.mkApply("clCreateBuffer",                  val globalBuffer = CL.mkAssign(CL.mkVar(concat[RN.globalsVarName,"_cl"]), CL.mkApply("clCreateBuffer",
405                                                                  [CL.E_Var contextVar,                                                                  [CL.mkVar contextVar,
406                                                                  CL.E_Var "CL_MEM_COPY_HOST_PTR",                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",
407                                                                  CL.mkApply("sizeof",[CL.E_Var RN.globalsTy]),                                                                  CL.mkApply("sizeof",[CL.mkVar RN.globalsTy]),
408                                                                  CL.E_Var RN.globalsVarName,                                                                  CL.mkVar RN.globalsVarName,
409                                                                  CL.E_UnOp(CL.%&,CL.E_Var errVar)]))                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))
410    
411          fun genDataBuffers([],_,_,_) = []          fun genDataBuffers([],_,_,_) = []
412            | genDataBuffers((var,nDims)::globals,count,contextVar,errVar) = let            | genDataBuffers((var,nDims)::globals,count,contextVar,errVar) = let
413          (* FIXME: use CL constructors to  build expressions (not strings) *)          (* FIXME: use CL constructors to  build expressions (not strings) *)
414                    val size = if nDims = 1 then                    val size = if nDims = 1 then
415                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.E_Var "float"]), CL.#*,                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,
416                                           CL.mkIndirect(CL.E_Var var, "size[0]"))                                           CL.mkIndirect(CL.mkVar var, "size[0]"))
417                                          else if nDims = 2 then                                          else if nDims = 2 then
418                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.E_Var "float"]), CL.#*,                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,
419                                            CL.mkIndirect(CL.E_Var var, concat["size[0]", " * ", var, "->size[1]"]))                                            CL.mkIndirect(CL.mkVar var, concat["size[0]", " * ", var, "->size[1]"]))
420                                          else                                          else
421                                           CL.mkBinOp(CL.mkApply("sizeof",[CL.E_Var "float"]), CL.#*,                                           CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,
422                                            CL.mkIndirect(CL.E_Var var,concat["size[0]", " * ", var, "->size[1] * ", var, "->size[2]"]))                                            CL.mkIndirect(CL.mkVar var,concat["size[0]", " * ", var, "->size[1] * ", var, "->size[2]"]))
423    
424                   in                   in
425                     CL.mkDecl(CL.clMemoryTy,RN.addBufferSuffix var ,NONE)::                     CL.mkDecl(clMemoryTy, RN.addBufferSuffix var ,NONE)::
426                     CL.mkDecl(CL.clMemoryTy,RN.addBufferSuffixData var ,NONE)::                     CL.mkDecl(clMemoryTy, RN.addBufferSuffixData var ,NONE)::
427                     CL.mkAssign(CL.E_Var(RN.addBufferSuffix var), CL.mkApply("clCreateBuffer",                     CL.mkAssign(CL.mkVar(RN.addBufferSuffix var), CL.mkApply("clCreateBuffer",
428                                                                  [CL.E_Var contextVar,                                                                  [CL.mkVar contextVar,
429                                                                  CL.E_Var "CL_MEM_COPY_HOST_PTR",                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",
430                                                                  CL.mkApply("sizeof",[CL.E_Var (RN.imageTy nDims)]),                                                                  CL.mkApply("sizeof",[CL.mkVar (RN.imageTy nDims)]),
431                                                                  CL.E_Var var,                                                                  CL.mkVar var,
432                                                                  CL.E_UnOp(CL.%&,CL.E_Var errVar)])) ::                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)])) ::
433                          CL.mkAssign(CL.E_Var(RN.addBufferSuffixData var), CL.mkApply("clCreateBuffer",                          CL.mkAssign(CL.mkVar(RN.addBufferSuffixData var), CL.mkApply("clCreateBuffer",
434                                                                  [CL.E_Var contextVar,                                                                  [CL.mkVar contextVar,
435                                                                   CL.E_Var "CL_MEM_COPY_HOST_PTR",                                                                   CL.mkVar "CL_MEM_COPY_HOST_PTR",
436                                                                  size,                                                                  size,
437                                                                  CL.mkIndirect(CL.E_Var var,"data"),                                                                  CL.mkIndirect(CL.mkVar var,"data"),
438                                                                  CL.E_UnOp(CL.%&,CL.E_Var errVar)])):: genDataBuffers(globals,count + 2,contextVar,errVar)                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)])):: genDataBuffers(globals,count + 2,contextVar,errVar)
439                  end                  end
440          in          in
441                  [globalBufferDecl] @ [globalBuffer] @ genDataBuffers(globals,count + 2,contextVar,errVar)                  [globalBufferDecl] @ [globalBuffer] @ genDataBuffers(globals,count + 2,contextVar,errVar)
# Line 415  Line 444 
444    
445  (* generates the kernel arguments for the image data *)  (* generates the kernel arguments for the image data *)
446          fun genGlobalArguments(globals,count,kernelVar,errVar) = let          fun genGlobalArguments(globals,count,kernelVar,errVar) = let
447          val globalArgument = CL.mkExpStm(CL.mkAssignOp(CL.E_Var errVar,CL.|=,CL.mkApply("clSetKernelArg",          val globalArgument = CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,CL.mkApply("clSetKernelArg",
448                                                                  [CL.E_Var kernelVar,                                                                  [CL.mkVar kernelVar,
449                                                                   CL.E_Int(count,CL.intTy),                                                                   CL.mkInt(count,CL.intTy),
450                                                                   CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
451                                                                   CL.E_UnOp(CL.%&,CL.E_Var(concat[RN.globalsVarName,"_cl"]))])))                                                                   CL.mkUnOp(CL.%&,CL.mkVar(concat[RN.globalsVarName,"_cl"]))])))
452    
453          fun genDataArguments([],_,_,_) = []          fun genDataArguments([],_,_,_) = []
454            | genDataArguments((var,nDims)::globals,count,kernelVar,errVar) =            | genDataArguments((var,nDims)::globals,count,kernelVar,errVar) =
455    
456                  CL.mkExpStm(CL.mkAssignOp(CL.E_Var errVar,CL.|=, CL.mkApply("clSetKernelArg",                  CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=, CL.mkApply("clSetKernelArg",
457                                                                  [CL.E_Var kernelVar,                                                                  [CL.mkVar kernelVar,
458                                                                   CL.E_Int(count,CL.intTy),                                                                   CL.mkInt(count,CL.intTy),
459                                                                   CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
460                                                                   CL.E_UnOp(CL.%&,CL.E_Var(RN.addBufferSuffix var))])))::                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffix var))])))::
461    
462                          CL.mkExpStm(CL.mkAssignOp(CL.E_Var errVar,CL.|=,CL.mkApply("clSetKernelArg",                          CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,CL.mkApply("clSetKernelArg",
463                                                                  [CL.E_Var kernelVar,                                                                  [CL.mkVar kernelVar,
464                                                                   CL.E_Int((count + 1),CL.intTy),                                                                   CL.mkInt((count + 1),CL.intTy),
465                                                                   CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
466                                                                   CL.E_UnOp(CL.%&,CL.E_Var(RN.addBufferSuffixData var))]))):: genDataArguments (globals, count + 2,kernelVar,errVar)                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffixData var))]))):: genDataArguments (globals, count + 2,kernelVar,errVar)
467    
468          in          in
469    
# Line 444  Line 473 
473    
474          (* generates the main function of host code *)          (* generates the main function of host code *)
475          fun genHostMain() = let          fun genHostMain() = let
476                val setupCall = [CL.mkCall(RN.setupFName,[CL.E_Var RN.globalsVarName])]                val setupCall = [CL.mkCall(RN.setupFName,[CL.mkVar RN.globalsVarName])]
477                val globalsDecl = CL.mkDecl(                val globalsDecl = CL.mkDecl(
478                      CL.T_Ptr(CL.T_Named RN.globalsTy),                      CL.T_Ptr(CL.T_Named RN.globalsTy),
479                      RN.globalsVarName,                      RN.globalsVarName,
480                      SOME(CL.I_Exp(CL.mkApply("malloc", [CL.mkApply("sizeof",[CL.E_Var RN.globalsTy])]))))                      SOME(CL.I_Exp(CL.mkApply("malloc", [CL.mkApply("sizeof",[CL.mkVar RN.globalsTy])]))))
481                val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.E_Var RN.globalsVarName])                val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.mkVar RN.globalsVarName])
482                val returnStm = [CL.mkReturn(SOME(CL.E_Int(0,CL.intTy)))]                val returnStm = [CL.mkReturn(SOME(CL.mkInt(0,CL.intTy)))]
483                val params = [                val params = [
484                       CL.PARAM([],CL.intTy, "argc"),                       CL.PARAM([],CL.intTy, "argc"),
485                       CL.PARAM([],CL.charArrayPtr,"argv")                       CL.PARAM([],CL.charArrayPtr,"argv")
# Line 485  Line 514 
514                val platformsVar = "platforms"                val platformsVar = "platforms"
515                val numPlatformsVar = "num_platforms"                val numPlatformsVar = "num_platforms"
516                val numDevicesVar = "num_devices"                val numDevicesVar = "num_devices"
517                val assertStm = CL.mkCall("assert",[CL.mkBinOp(CL.E_Var errVar, CL.#==, CL.E_Var "CL_SUCCESS")])                val assertStm = CL.mkCall("assert",[CL.mkBinOp(CL.mkVar errVar, CL.#==, CL.mkVar "CL_SUCCESS")])
518                val params = [                val params = [
519                        CL.PARAM([],CL.T_Named("cl_device_id"), deviceVar)                        CL.PARAM([],CL.T_Named("cl_device_id"), deviceVar)
520                      ]                      ]
521                val declarations = [                val declarations = [
522                      CL.mkDecl(CL.clProgramTy, programVar, NONE),                      CL.mkDecl(clProgramTy, programVar, NONE),
523                      CL.mkDecl(CL.clKernelTy, kernelVar, NONE),                      CL.mkDecl(clKernelTy, kernelVar, NONE),
524                      CL.mkDecl(CL.clCmdQueueTy, cmdVar, NONE),                      CL.mkDecl(clCmdQueueTy, cmdVar, NONE),
525                      CL.mkDecl(CL.clContextTy, contextVar, NONE),                      CL.mkDecl(clContextTy, contextVar, NONE),
526                      CL.mkDecl(CL.intTy, errVar, NONE),                      CL.mkDecl(CL.intTy, errVar, NONE),
527                      CL.mkDecl(CL.intTy, numStrandsVar, SOME(CL.I_Exp(CL.E_Int(1,CL.intTy)))),                      CL.mkDecl(CL.intTy, numStrandsVar, SOME(CL.I_Exp(CL.mkInt(1,CL.intTy)))),
528                      CL.mkDecl(CL.intTy, stateSizeVar, NONE),                      CL.mkDecl(CL.intTy, stateSizeVar, NONE),
529                      CL.mkDecl(CL.intTy, "width", NONE),                      CL.mkDecl(CL.intTy, "width", NONE),
530                      CL.mkDecl(CL.intTy, imgDataSizeVar, NONE),                      CL.mkDecl(CL.intTy, imgDataSizeVar, NONE),
531                      (*CL.mkDecl(CL.clDeviceIdTy, deviceVar, NONE), *)                      (*CL.mkDecl(clDeviceIdTy, deviceVar, NONE), *)
532                      CL.mkDecl(CL.T_Ptr(CL.T_Named tyName), inStateVar,NONE),                      CL.mkDecl(CL.T_Ptr(CL.T_Named tyName), inStateVar,NONE),
533                      CL.mkDecl(CL.clMemoryTy,clInstateVar,NONE),                      CL.mkDecl(clMemoryTy,clInstateVar,NONE),
534                      CL.mkDecl(CL.clMemoryTy,clOutStateVar,NONE),                      CL.mkDecl(clMemoryTy,clOutStateVar,NONE),
535                      CL.mkDecl(CL.T_Ptr(CL.T_Named tyName), outStateVar,NONE),                      CL.mkDecl(CL.T_Ptr(CL.T_Named tyName), outStateVar,NONE),
536                      CL.mkDecl(CL.charPtr, clFNVar,SOME(CL.I_Exp(CL.E_Str filename))),                      CL.mkDecl(CL.charPtr, clFNVar,SOME(CL.I_Exp(CL.mkStr filename))),
537  (* FIXME:  use Paths.diderotInclude *)  (* FIXME:  use Paths.diderotInclude *)
538                      CL.mkDecl(CL.charPtr, headerFNVar,SOME(CL.I_Exp(CL.E_Str "../src/include/Diderot/cl-types.h"))),                      CL.mkDecl(CL.charPtr, headerFNVar,SOME(CL.I_Exp(CL.mkStr "../src/include/Diderot/cl-types.h"))),
539                      CL.mkDecl(CL.T_Array(CL.charPtr,SOME(2)),sourcesVar,NONE),                      CL.mkDecl(CL.T_Array(CL.charPtr,SOME(2)),sourcesVar,NONE),
540                      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)),globalVar,NONE),
541                      CL.mkDecl(CL.T_Array(CL.T_Named "size_t",SOME(nDims)),localVar,NONE),                      CL.mkDecl(CL.T_Array(CL.T_Named "size_t",SOME(nDims)),localVar,NONE),
542                      CL.mkDecl(CL.intTy,numDevicesVar,SOME(CL.I_Exp(CL.E_Int(~1,CL.intTy)))),                      CL.mkDecl(CL.intTy,numDevicesVar,SOME(CL.I_Exp(CL.mkInt(~1,CL.intTy)))),
543                      CL.mkDecl(CL.T_Array(CL.T_Named "cl_platform_id", SOME(1)), platformsVar, NONE),                      CL.mkDecl(CL.T_Array(CL.T_Named "cl_platform_id", SOME(1)), platformsVar, NONE),
544                      CL.mkDecl(CL.intTy,"num_platforms",SOME(CL.I_Exp(CL.E_Int(~1,CL.intTy))))                      CL.mkDecl(CL.intTy,"num_platforms",SOME(CL.I_Exp(CL.mkInt(~1,CL.intTy))))
545                  ]                  ]
546              (* Setup Global Variables *)              (* Setup Global Variables *)
547                val globalsDecl = CL.mkDecl(                val globalsDecl = CL.mkDecl(
548                      CL.T_Ptr(CL.T_Named RN.globalsTy),                      CL.T_Ptr(CL.T_Named RN.globalsTy),
549                      RN.globalsVarName,                      RN.globalsVarName,
550                      SOME(CL.I_Exp(CL.mkApply("malloc", [CL.mkApply("sizeof",[CL.E_Var RN.globalsTy])]))))                      SOME(CL.I_Exp(CL.mkApply("malloc", [CL.mkApply("sizeof",[CL.mkVar RN.globalsTy])]))))
551                val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.E_Var RN.globalsVarName])                val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.mkVar RN.globalsVarName])
552    
553                  (* Retrieve the platforms                  (* Retrieve the platforms
554                  val platformStm = [CL.mkAssign(CL.E_Var errVar, CL.mkApply("clGetPlatformIDs",                  val platformStm = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetPlatformIDs",
555                                                    [CL.E_Int(10,CL.intTy),                                                    [CL.mkInt(10,CL.intTy),
556                                                     CL.E_Var platformsVar,                                                     CL.mkVar platformsVar,
557                                                     CL.E_UnOp(CL.%&,CL.E_Var numPlatformsVar)])),                                                     CL.mkUnOp(CL.%&,CL.mkVar numPlatformsVar)])),
558                                                     assertStm]                                                     assertStm]
559    
560                  val devicesStm = [CL.mkAssign(CL.E_Var errVar, CL.mkApply("clGetDeviceIDs",                  val devicesStm = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetDeviceIDs",
561                                                    [CL.mkSubscript(CL.E_Var platformsVar,CL.E_Int(0,CL.intTy)),                                                    [CL.mkSubscript(CL.mkVar platformsVar,CL.mkInt(0,CL.intTy)),
562                                                     CL.E_Var "CL_DEVICE_TYPE_GPU",                                                     CL.mkVar "CL_DEVICE_TYPE_GPU",
563                                                     CL.E_Int(1,CL.intTy),                                                     CL.mkInt(1,CL.intTy),
564                                                     CL.E_UnOp(CL.%&,CL.E_Var deviceVar),                                                     CL.mkUnOp(CL.%&,CL.mkVar deviceVar),
565                                                     CL.E_UnOp(CL.%&,CL.E_Var numDevicesVar)])),                                                     CL.mkUnOp(CL.%&,CL.mkVar numDevicesVar)])),
566                                                     assertStm] *)                                                     assertStm] *)
567    
568                  (* Create Context *)                  (* Create Context *)
569                  val contextStm = [CL.mkAssign(CL.E_Var contextVar, CL.mkApply("clCreateContext",                  val contextStm = [CL.mkAssign(CL.mkVar contextVar, CL.mkApply("clCreateContext",
570                                                    [CL.E_Int(0,CL.intTy),                                                    [CL.mkInt(0,CL.intTy),
571                                                    CL.E_Int(1,CL.intTy),                                                    CL.mkInt(1,CL.intTy),
572                                                    CL.E_UnOp(CL.%&,CL.E_Var deviceVar),                                                    CL.mkUnOp(CL.%&,CL.mkVar deviceVar),
573                                                    CL.E_Var "NULL",                                                    CL.mkVar "NULL",
574                                                    CL.E_Var "NULL",                                                    CL.mkVar "NULL",
575                                                    CL.E_UnOp(CL.%&,CL.E_Var errVar)])),                                                    CL.mkUnOp(CL.%&,CL.mkVar errVar)])),
576                                                    assertStm]                                                    assertStm]
577    
578                  (* Create Command Queue *)                  (* Create Command Queue *)
579                  val commandStm = [CL.mkAssign(CL.E_Var cmdVar, CL.mkApply("clCreateCommandQueue",                  val commandStm = [CL.mkAssign(CL.mkVar cmdVar, CL.mkApply("clCreateCommandQueue",
580                                                    [CL.E_Var contextVar,                                                    [CL.mkVar contextVar,
581                                                    CL.E_Var deviceVar,                                                    CL.mkVar deviceVar,
582                                                    CL.E_Int(0,CL.intTy),                                                    CL.mkInt(0,CL.intTy),
583                                                    CL.E_UnOp(CL.%&,CL.E_Var errVar)])),                                                    CL.mkUnOp(CL.%&,CL.mkVar errVar)])),
584                                                    assertStm]                                                    assertStm]
585    
586    
587                  (*Create Program/Build/Kernel with Source statement *)                  (*Create Program/Build/Kernel with Source statement *)
588                  val createProgStm = CL.mkAssign(CL.E_Var programVar, CL.mkApply("clCreateProgramWithSource",                  val createProgStm = CL.mkAssign(CL.mkVar programVar, CL.mkApply("clCreateProgramWithSource",
589                                                                                                                  [CL.E_Var contextVar,                                                                                                                  [CL.mkVar contextVar,
590                                                                                                                   CL.E_Int(2,CL.intTy),                                                                                                                   CL.mkInt(2,CL.intTy),
591                                                                                                                   CL.E_Cast(CL.T_Ptr(CL.T_Named("const char *")),CL.E_UnOp(CL.%&,CL.E_Var sourcesVar)),                                                                                                                   CL.mkCast(CL.T_Ptr(CL.T_Named("const char *")),CL.mkUnOp(CL.%&,CL.mkVar sourcesVar)),
592                                                                                                                   CL.E_Var "NULL",                                                                                                                   CL.mkVar "NULL",
593                                                                                                                   CL.E_UnOp(CL.%&,CL.E_Var errVar)]))                                                                                                                   CL.mkUnOp(CL.%&,CL.mkVar errVar)]))
594    
595                  (* FIXME: Remove after testing purposes, Build Log for OpenCL*)                  (* FIXME: Remove after testing purposes, Build Log for OpenCL*)
596                  val buildLog = [CL.mkAssign(CL.E_Var errVar, CL.mkApply("clBuildProgram",                  val buildLog = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clBuildProgram",
597                                                                                                                  [CL.E_Var programVar,                                                                                                                  [CL.mkVar programVar,
598                                                                                                                   CL.E_Int(0,CL.intTy),                                                                                                                   CL.mkInt(0,CL.intTy),
599                                                                                                                   CL.E_Var "NULL",                                                                                                                   CL.mkVar "NULL",
600                                                                                                                   CL.E_Var "NULL",                                                                                                                   CL.mkVar "NULL",
601                                                                                                                   CL.E_Var "NULL",                                                                                                                   CL.mkVar "NULL",
602                                                                                                                   CL.E_Var "NULL"])),                                                                                                                   CL.mkVar "NULL"])),
603                                            CL.mkDecl(CL.charPtr, "build", NONE),                                            CL.mkDecl(CL.charPtr, "build", NONE),
604                                            CL.mkDecl(CL.T_Named("size_t"),"ret_val_size",NONE),                                            CL.mkDecl(CL.T_Named("size_t"),"ret_val_size",NONE),
605                                             CL.mkAssign(CL.E_Var errVar, CL.mkApply("clGetProgramBuildInfo",                                             CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetProgramBuildInfo",
606                                                                                                                  [CL.E_Var programVar,                                                                                                                  [CL.mkVar programVar,
607                                                                                                                  CL.E_Var deviceVar,                                                                                                                  CL.mkVar deviceVar,
608                                                                                                                   CL.E_Var "CL_PROGRAM_BUILD_LOG",                                                                                                                   CL.mkVar "CL_PROGRAM_BUILD_LOG",
609                                                                                                                   CL.E_Int(0,CL.intTy),                                                                                                                   CL.mkInt(0,CL.intTy),
610                                                                                                                   CL.E_Var "NULL",                                                                                                                   CL.mkVar "NULL",
611                                                                                                                   CL.E_UnOp(CL.%&,CL.E_Var "ret_val_size")])),                                                                                                                   CL.mkUnOp(CL.%&,CL.mkVar "ret_val_size")])),
612                                            CL.mkAssign(CL.E_Var "build", CL.mkApply("malloc", [CL.E_Var "ret_val_size"])),                                            CL.mkAssign(CL.mkVar "build", CL.mkApply("malloc", [CL.mkVar "ret_val_size"])),
613                                                  CL.mkAssign(CL.E_Var errVar, CL.mkApply("clGetProgramBuildInfo",                                                  CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetProgramBuildInfo",
614                                                                                                                  [CL.E_Var programVar,                                                                                                                  [CL.mkVar programVar,
615                                                                                                                  CL.E_Var deviceVar,                                                                                                                  CL.mkVar deviceVar,
616                                                                                                                   CL.E_Var "CL_PROGRAM_BUILD_LOG",                                                                                                                   CL.mkVar "CL_PROGRAM_BUILD_LOG",
617                                                                                                                   CL.E_Var "ret_val_size",                                                                                                                   CL.mkVar "ret_val_size",
618                                                                                                                   CL.E_Var "build",                                                                                                                   CL.mkVar "build",
619                                                                                                                   CL.E_Var "NULL"])),                                                                                                                   CL.mkVar "NULL"])),
620                                                  CL.mkAssign(CL.mkSubscript(CL.E_Var "build",CL.E_Var "ret_val_size"),CL.E_Var ("'\\" ^ "0'")),                                                  CL.mkAssign(CL.mkSubscript(CL.mkVar "build",CL.mkVar "ret_val_size"),CL.mkVar ("'\\" ^ "0'")),
621                                                  CL.mkCall("printf",[CL.E_Str ( "Build Log:" ^ "\n" ^ "%s" ^ "\n"), CL.E_Var "build"])]                                                  CL.mkCall("printf",[CL.mkStr ( "Build Log:" ^ "\n" ^ "%s" ^ "\n"), CL.mkVar "build"])]
622    
623    
624    
625    
626                  val createKernel = CL.mkAssign(CL.E_Var kernelVar, CL.mkApply("clCreateKernel",                  val createKernel = CL.mkAssign(CL.mkVar kernelVar, CL.mkApply("clCreateKernel",
627                                                                                                                  [CL.E_Var programVar,                                                                                                                  [CL.mkVar programVar,
628                                                                                                                   CL.E_Str RN.kernelFuncName,                                                                                                                   CL.mkStr RN.kernelFuncName,
629                                                                                                                   CL.E_UnOp(CL.%&,CL.E_Var errVar)]))                                                                                                                   CL.mkUnOp(CL.%&,CL.mkVar errVar)]))
630    
631    
632                  val create_build_stms = [createProgStm,assertStm] @ buildLog @ [assertStm,createKernel,assertStm]                  val create_build_stms = [createProgStm,assertStm] @ buildLog @ [assertStm,createKernel,assertStm]
# Line 605  Line 634 
634    
635    
636                  (* Create Memory Buffers for Strand States and Globals *)                  (* Create Memory Buffers for Strand States and Globals *)
637                  val strandSize = CL.mkAssign(CL.E_Var stateSizeVar,CL.mkBinOp(CL.mkApply("sizeof",                  val strandSize = CL.mkAssign(CL.mkVar stateSizeVar,CL.mkBinOp(CL.mkApply("sizeof",
638                                                                          [CL.E_Var tyName]), CL.#*,CL.E_Var numStrandsVar))                                                                          [CL.mkVar tyName]), CL.#*,CL.mkVar numStrandsVar))
639    
640                  val clStrandObjects = [CL.mkAssign(CL.E_Var clInstateVar, CL.mkApply("clCreateBuffer",                  val clStrandObjects = [CL.mkAssign(CL.mkVar clInstateVar, CL.mkApply("clCreateBuffer",
641                                                                  [CL.E_Var contextVar,                                                                  [CL.mkVar contextVar,
642                                                                  CL.E_Var "CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR",                                                                  CL.mkVar "CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR",
643                                                                  CL.E_Var stateSizeVar,                                                                  CL.mkVar stateSizeVar,
644                                                                  CL.E_Var "NULL",                                                                  CL.mkVar "NULL",
645                                                                  CL.E_UnOp(CL.%&,CL.E_Var errVar)])),                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)])),
646                                                           CL.mkAssign(CL.E_Var clOutStateVar, CL.mkApply("clCreateBuffer",                                                           CL.mkAssign(CL.mkVar clOutStateVar, CL.mkApply("clCreateBuffer",
647                                                                  [CL.E_Var contextVar,                                                                  [CL.mkVar contextVar,
648                                                                  CL.E_Var "CL_MEM_READ_WRITE",                                                                  CL.mkVar "CL_MEM_READ_WRITE",
649                                                                  CL.E_Var stateSizeVar,                                                                  CL.mkVar stateSizeVar,
650                                                                  CL.E_Var "NULL",                                                                  CL.mkVar "NULL",
651                                                                  CL.E_UnOp(CL.%&,CL.E_Var errVar)]))]                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))]
652    
653    
654                  (* Setup up selfOut variable *)                  (* Setup up selfOut variable *)
655                  val strandsArrays = [CL.mkAssign(CL.E_Var outStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.E_Var numStrandsVar,                  val strandsArrays = [CL.mkAssign(CL.mkVar outStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.mkVar numStrandsVar,
656                                                                          CL.#*, CL.mkApply("sizeof",[CL.E_Var tyName]))])),                                                                          CL.#*, CL.mkApply("sizeof",[CL.mkVar tyName]))])),
657                                                                  CL.mkAssign(CL.E_Var inStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.E_Var numStrandsVar,                                                                  CL.mkAssign(CL.mkVar inStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.mkVar numStrandsVar,
658                                                                          CL.#*, CL.mkApply("sizeof",[CL.E_Var tyName]))]))]                                                                          CL.#*, CL.mkApply("sizeof",[CL.mkVar tyName]))]))]
659    
660    
661                  (* Initialize Width Parameter *)                  (* Initialize Width Parameter *)
662                  val widthDel = if nDims = 2 then                  val widthDel = if nDims = 2 then
663                            CL.mkAssign(CL.E_Var "width",CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(1, CL.intTy)))                            CL.mkAssign(CL.mkVar "width",CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1, CL.intTy)))
664                     else                     else
665                            CL.mkAssign(CL.E_Var "width",CL.E_Int(0,CL.intTy))                            CL.mkAssign(CL.mkVar "width",CL.mkInt(0,CL.intTy))
666    
667    
668                  val strands_init =      CL.mkCall(RN.strandInitSetup,                  val strands_init = CL.mkCall(RN.strandInitSetup, [
669                                                                           [CL.E_Var "size",                          CL.mkVar "size", CL.mkVar "width", CL.mkVar inStateVar
670                                                                           CL.E_Var "width",                        ])
                                                                          CL.E_Var inStateVar])  
671    
672              val clGlobalBuffers = getGlobalDataBuffers(!imgGlobals,3,contextVar,errVar)              val clGlobalBuffers = getGlobalDataBuffers(!imgGlobals,3,contextVar,errVar)
673    
674    
675                  (* Load the Kernel and Header Files *)                  (* Load the Kernel and Header Files *)
676                  val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.E_Var sourcesVar,CL.E_Int(1,CL.intTy)),                  val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(1,CL.intTy)),
677                                                                            CL.mkApply(RN.clLoaderFN, [CL.E_Var clFNVar])),                                                                            CL.mkApply(RN.clLoaderFN, [CL.mkVar clFNVar])),
678             CL.mkAssign(CL.mkSubscript(CL.E_Var sourcesVar,CL.E_Int(0,CL.intTy)),             CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(0,CL.intTy)),
679                                                                            CL.mkApply(RN.clLoaderFN, [CL.E_Var headerFNVar]))]                                                                            CL.mkApply(RN.clLoaderFN, [CL.mkVar headerFNVar]))]
680    
681                  (* val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.E_Var sourcesVar,CL.E_Int(1,CL.intTy)),                  (* val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(1,CL.intTy)),
682                                                                            CL.mkApply(RN.clLoaderFN, [CL.E_Var clFNVar]))] *)                                                                            CL.mkApply(RN.clLoaderFN, [CL.mkVar clFNVar]))] *)
683    
684    
685                  (* Created Enqueue Statements *)                  (* Created Enqueue Statements *)
686  (* FIXME: simplify this code by function abstraction *)  (* FIXME: simplify this code by function abstraction *)
687          val enqueueStm = if nDims = 1          val enqueueStm = if nDims = 1
688                          then [CL.mkAssign(CL.E_Var errVar,                          then [CL.mkAssign(CL.mkVar errVar,
689                                                            CL.mkApply("clEnqueueNDRangeKernel",                                                            CL.mkApply("clEnqueueNDRangeKernel",
690                                                                                                  [CL.E_Var cmdVar,                                                                                                  [CL.mkVar cmdVar,
691                                                                                                   CL.E_Var kernelVar,                                                                                                   CL.mkVar kernelVar,
692                                                                                                   CL.E_Int(1,CL.intTy),                                                                                                   CL.mkInt(1,CL.intTy),
693                                                                                                   CL.E_Var "NULL",                                                                                                   CL.mkVar "NULL",
694                                                                                                   CL.E_Var globalVar,                                                                                                   CL.mkVar globalVar,
695                                                                                                   CL.E_Var localVar,                                                                                                   CL.mkVar localVar,
696                                                                                                   CL.E_Int(0,CL.intTy),                                                                                                   CL.mkInt(0,CL.intTy),
697                                                                                                   CL.E_Var "NULL",                                                                                                   CL.mkVar "NULL",
698                                                                                                   CL.E_Var "NULL"])),CL.mkCall("clFinish",[CL.E_Var cmdVar])]                                                                                                   CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]
699                          else if nDims = 2  then                          else if nDims = 2  then
700                           [CL.mkAssign(CL.E_Var errVar,                           [CL.mkAssign(CL.mkVar errVar,
701                                                          CL.mkApply("clEnqueueNDRangeKernel",                                                          CL.mkApply("clEnqueueNDRangeKernel",
702                                                                                                  [CL.E_Var cmdVar,                                                                                                  [CL.mkVar cmdVar,
703                                                                                                   CL.E_Var kernelVar,                                                                                                   CL.mkVar kernelVar,
704                                                                                                   CL.E_Int(2,CL.intTy),                                                                                                   CL.mkInt(2,CL.intTy),
705                                                                                                   CL.E_Var "NULL",                                                                                                   CL.mkVar "NULL",
706                                                                                                   CL.E_Var globalVar,                                                                                                   CL.mkVar globalVar,
707                                                                                                   CL.E_Var localVar,                                                                                                   CL.mkVar localVar,
708                                                                                                   CL.E_Int(0,CL.intTy),                                                                                                   CL.mkInt(0,CL.intTy),
709                                                                                                   CL.E_Var "NULL",                                                                                                   CL.mkVar "NULL",
710                                                                                                   CL.E_Var "NULL"])),CL.mkCall("clFinish",[CL.E_Var cmdVar])]                                                                                                   CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]
711                          else                          else
712                            [CL.mkAssign(CL.E_Var errVar,                            [CL.mkAssign(CL.mkVar errVar,
713                                                          CL.mkApply("clEnqueueNDRangeKernel",                                                          CL.mkApply("clEnqueueNDRangeKernel",
714                                                                                                  [CL.E_Var cmdVar,                                                                                                  [CL.mkVar cmdVar,
715                                                                                                   CL.E_Var kernelVar,                                                                                                   CL.mkVar kernelVar,
716                                                                                                   CL.E_Int(3,CL.intTy),                                                                                                   CL.mkInt(3,CL.intTy),
717                                                                                                   CL.E_Var "NULL",                                                                                                   CL.mkVar "NULL",
718                                                                                                   CL.E_Var globalVar,                                                                                                   CL.mkVar globalVar,
719                                                                                                   CL.E_Var localVar,                                                                                                   CL.mkVar localVar,
720                                                                                                   CL.E_Int(0,CL.intTy),                                                                                                   CL.mkInt(0,CL.intTy),
721                                                                                                   CL.E_Var "NULL",                                                                                                   CL.mkVar "NULL",
722                                                                                                   CL.E_Var "NULL"])),CL.mkCall("clFinish",[CL.E_Var cmdVar])]                                                                                                   CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]
723    
724    
725    
726                  (* Setup Global and Local variables *)                  (* Setup Global and Local variables *)
727    
728                  val globalAndlocalStms = if nDims = 1 then                  val globalAndlocalStms = if nDims = 1 then
729                          [CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(0,CL.intTy)),                          [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),
730                                                                     CL.mkSubscript(CL.E_Var "size", CL.E_Int(0,CL.intTy))),                                                                     CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),
731                           CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(0,CL.intTy)),                           CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),
732                                                                    CL.E_Var "16")]                                                                    CL.mkVar "16")]
733    
734    
735                  else if nDims = 2 then                  else if nDims = 2 then
736                          [CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(0,CL.intTy)),                          [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),
737                                                                     CL.mkSubscript(CL.E_Var "size", CL.E_Int(0,CL.intTy))),                                                                     CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),
738                          CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(1,CL.intTy)),                          CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1,CL.intTy)),
739                                                                     CL.mkSubscript(CL.E_Var "size", CL.E_Int(1,CL.intTy))),                                                                     CL.mkSubscript(CL.mkVar "size", CL.mkInt(1,CL.intTy))),
740                          CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(0,CL.intTy)),                          CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),
741                                                                    CL.E_Var "16"),                                                                    CL.mkVar "16"),
742                          CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(1,CL.intTy)),                          CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(1,CL.intTy)),
743                                                                    CL.E_Var "16")]                                                                    CL.mkVar "16")]
744    
745                  else                  else
746                          [CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(0,CL.intTy)),                          [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),
747                                                                     CL.mkSubscript(CL.E_Var "size", CL.E_Int(0,CL.intTy))),                                                                     CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),
748                          CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(1,CL.intTy)),                          CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1,CL.intTy)),
749                                                                     CL.mkSubscript(CL.E_Var "size", CL.E_Int(1,CL.intTy))),                                                                     CL.mkSubscript(CL.mkVar "size", CL.mkInt(1,CL.intTy))),
750                          CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(2,CL.intTy)),                          CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(2,CL.intTy)),
751                                                                     CL.mkSubscript(CL.E_Var "size", CL.E_Int(2,CL.intTy))),                                                                     CL.mkSubscript(CL.mkVar "size", CL.mkInt(2,CL.intTy))),
752                          CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(0,CL.intTy)),                          CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),
753                                                                    CL.E_Var "16"),                                                                    CL.mkVar "16"),
754                          CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(1,CL.intTy)),                          CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(1,CL.intTy)),
755                                                                    CL.E_Var "16"),                                                                    CL.mkVar "16"),
756                          CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(2,CL.intTy)),                          CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(2,CL.intTy)),
757                                                                    CL.E_Var "16")]                                                                    CL.mkVar "16")]
758    
759    
760    
761                  (* Setup Kernel arguments *)                  (* Setup Kernel arguments *)
762                  val kernelArguments = [CL.mkAssign(CL.E_Var errVar,CL.mkApply("clSetKernelArg",                  val kernelArguments = [CL.mkAssign(CL.mkVar errVar,CL.mkApply("clSetKernelArg",
763                                                                  [CL.E_Var kernelVar,                                                                  [CL.mkVar kernelVar,
764                                                                   CL.E_Int(0,CL.intTy),                                                                   CL.mkInt(0,CL.intTy),
765                                                                   CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
766                                                                   CL.E_UnOp(CL.%&,CL.E_Var clInstateVar)])),                                                                   CL.mkUnOp(CL.%&,CL.mkVar clInstateVar)])),
767                                                              CL.mkExpStm(CL.mkAssignOp(CL.E_Var errVar, CL.|=,CL.mkApply("clSetKernelArg",                                                              CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar, CL.|=,CL.mkApply("clSetKernelArg",
768                                                                  [CL.E_Var kernelVar,                                                                  [CL.mkVar kernelVar,
769                                                                   CL.E_Int(1,CL.intTy),                                                                   CL.mkInt(1,CL.intTy),
770                                                                   CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
771                                                                   CL.E_UnOp(CL.%&,CL.E_Var clOutStateVar)]))),                                                                   CL.mkUnOp(CL.%&,CL.mkVar clOutStateVar)]))),
772                                                                    CL.mkExpStm(CL.mkAssignOp(CL.E_Var errVar, CL.|=,CL.mkApply("clSetKernelArg",                                                                    CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar, CL.|=,CL.mkApply("clSetKernelArg",
773                                                                  [CL.E_Var kernelVar,                                                                  [CL.mkVar kernelVar,
774                                                                   CL.E_Int(2,CL.intTy),                                                                   CL.mkInt(2,CL.intTy),
775                                                                   CL.mkApply("sizeof",[CL.E_Var "int"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "int"]),
776                                                                   CL.E_UnOp(CL.%&,CL.E_Var "width")])))]                                                                   CL.mkUnOp(CL.%&,CL.mkVar "width")])))]
777    
778             val clGlobalArguments = genGlobalArguments(!imgGlobals,3,kernelVar,errVar) @ [assertStm]             val clGlobalArguments = genGlobalArguments(!imgGlobals,3,kernelVar,errVar) @ [assertStm]
779    
780                  (* Retrieve output *)                  (* Retrieve output *)
781                  val outputStm = CL.mkAssign(CL.E_Var errVar,                  val outputStm = CL.mkAssign(CL.mkVar errVar,
782                                                          CL.mkApply("clEnqueueReadBuffer",                                                          CL.mkApply("clEnqueueReadBuffer",
783                                                                                                  [CL.E_Var cmdVar,                                                                                                  [CL.mkVar cmdVar,
784                                                                                                   CL.E_Var clOutStateVar,                                                                                                   CL.mkVar clOutStateVar,
785                                                                                                   CL.E_Var "CL_TRUE",                                                                                                   CL.mkVar "CL_TRUE",
786                                                                                                   CL.E_Int(0,CL.intTy),                                                                                                   CL.mkInt(0,CL.intTy),
787                                                                                                   CL.E_Var stateSizeVar,                                                                                                   CL.mkVar stateSizeVar,
788                                                                                                   CL.E_Var outStateVar,                                                                                                   CL.mkVar outStateVar,
789                                                                                                   CL.E_Int(0,CL.intTy),                                                                                                   CL.mkInt(0,CL.intTy),
790                                                                                                   CL.E_Var "NULL",                                                                                                   CL.mkVar "NULL",
791                                                                                                   CL.E_Var "NULL"]))                                                                                                   CL.mkVar "NULL"]))
792    
793                  (* Free all the objects *)                  (* Free all the objects *)
794                  val freeStms = [CL.mkCall("clReleaseKernel",[CL.E_Var kernelVar]),                  val freeStms = [CL.mkCall("clReleaseKernel",[CL.mkVar kernelVar]),
795                                                  CL.mkCall("clReleaseProgram",[CL.E_Var programVar ]),                                                  CL.mkCall("clReleaseProgram",[CL.mkVar programVar ]),
796                                                  CL.mkCall("clReleaseCommandQueue",[CL.E_Var cmdVar]),                                                  CL.mkCall("clReleaseCommandQueue",[CL.mkVar cmdVar]),
797                                                  CL.mkCall("clReleaseContext",[CL.E_Var contextVar]),                                                  CL.mkCall("clReleaseContext",[CL.mkVar contextVar]),
798                                                  CL.mkCall("clReleaseMemObject",[CL.E_Var clInstateVar]),                                                  CL.mkCall("clReleaseMemObject",[CL.mkVar clInstateVar]),
799                                                  CL.mkCall("clReleaseMemObject",[CL.E_Var clOutStateVar])]                                                  CL.mkCall("clReleaseMemObject",[CL.mkVar clOutStateVar])]
800    
801    
802                  (*Setup Strand Print Function *)                  (*Setup Strand Print Function *)
803                  val outputData = [CL.mkDecl(CL.T_Ptr(CL.T_Named("FILE")), "outS", SOME(CL.I_Exp(CL.mkApply("fopen",                  val outputData = [CL.mkDecl(CL.T_Ptr(CL.T_Named("FILE")), "outS", SOME(CL.I_Exp(CL.mkApply("fopen",
804                                                  [CL.E_Str "mip.txt",                                                  [CL.mkStr "mip.txt",
805                                                  CL.E_Str "w"])))),                                                  CL.mkStr "w"])))),
806                                                  CL.mkCall(concat[name, "_print"],                                                  CL.mkCall(concat[name, "_print"],
807                                                                          [CL.E_Var "outS",                                                                          [CL.mkVar "outS",
808                                                                           CL.E_Var "size",                                                                           CL.mkVar "size",
809                                                                           CL.E_Var "width",                                                                           CL.mkVar "width",
810                                                                           CL.E_Var outStateVar])]                                                                           CL.mkVar outStateVar])]
811    
812    
813    
# Line 805  Line 833 
833          (*generate code for intilizing kernel global data *)          (*generate code for intilizing kernel global data *)
834          fun initKernelGlobals (globals,imgGlobals) = let          fun initKernelGlobals (globals,imgGlobals) = let
835                  fun initGlobalStruct (CL.D_Var(_, _ , name, _)::rest) =                  fun initGlobalStruct (CL.D_Var(_, _ , name, _)::rest) =
836                                  CL.mkAssign(CL.E_Var name, CL.mkIndirect(CL.E_Var RN.globalsVarName, name)) ::                                  CL.mkAssign(CL.mkVar name, CL.mkIndirect(CL.mkVar RN.globalsVarName, name)) ::
837                                  initGlobalStruct(rest)                                  initGlobalStruct(rest)
838                    | initGlobalStruct ( _::rest) = initGlobalStruct(rest)                    | initGlobalStruct ( _::rest) = initGlobalStruct(rest)
839                    | initGlobalStruct([]) = []                    | initGlobalStruct([]) = []
840    
841                  fun initGlobalImages((name,tyname)::rest) =                  fun initGlobalImages((name,tyname)::rest) =
842                                  CL.mkAssign(CL.E_Var name, CL.E_Var (RN.addBufferSuffix name)) ::                                  CL.mkAssign(CL.mkVar name, CL.mkVar (RN.addBufferSuffix name)) ::
843                                  CL.mkAssign(CL.mkIndirect(CL.E_Var name,"data"),CL.E_Var (RN.addBufferSuffixData name)) ::                                  CL.mkAssign(CL.mkIndirect(CL.mkVar name,"data"),CL.mkVar (RN.addBufferSuffixData name)) ::
844                                  initGlobalImages(rest)                                  initGlobalImages(rest)
845                    | initGlobalImages([]) = []                    | initGlobalImages([]) = []
846                  in                  in
# Line 830  Line 858 
858                        CL.PARAM(["__global"], CL.intTy, "width")                        CL.PARAM(["__global"], CL.intTy, "width")
859                      ] @ genKeneralGlobalParams(!imgGlobals)                      ] @ genKeneralGlobalParams(!imgGlobals)
860                    val thread_ids = if nDims = 1                    val thread_ids = if nDims = 1
861                          then [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy)))),                          then [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),
862                                    CL.mkAssign(CL.E_Var "x",CL.mkApply(RN.getGlobalThreadId,[CL.E_Int(0,CL.intTy)]))]                                    CL.mkAssign(CL.mkVar "x",CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(0,CL.intTy)]))]
863                          else                          else
864                                  [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy)))),                                  [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),
865                                   CL.mkDecl(CL.intTy, "y", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy)))),                                   CL.mkDecl(CL.intTy, "y", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),
866                                    CL.mkAssign(CL.E_Var "x",  CL.mkApply(RN.getGlobalThreadId,[CL.E_Int(0,CL.intTy)])),                                    CL.mkAssign(CL.mkVar "x",  CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(0,CL.intTy)])),
867                                    CL.mkAssign(CL.E_Var "y",CL.mkApply(RN.getGlobalThreadId,[CL.E_Int(1,CL.intTy)]))]                                    CL.mkAssign(CL.mkVar "y",CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(1,CL.intTy)]))]
868    
869                    val strandDecl = [CL.mkDecl(CL.T_Named tyName, inState, NONE),                    val strandDecl = [CL.mkDecl(CL.T_Named tyName, inState, NONE),
870                                                          CL.mkDecl(CL.T_Named tyName, outState,NONE)]                                                          CL.mkDecl(CL.T_Named tyName, outState,NONE)]
871                    val strandObjects  = if nDims = 1                    val strandObjects  = if nDims = 1
872                          then [CL.mkAssign(CL.mkSubscript(CL.E_Var "selfIn",CL.E_Str "x"),                          then [CL.mkAssign(CL.mkSubscript(CL.mkVar "selfIn",CL.mkStr "x"),
873                                                                           CL.E_Var inState),                                                                           CL.mkVar inState),
874                                    CL.mkAssign(CL.mkSubscript(CL.E_Var "selfOut",CL.E_Str "x"),                                    CL.mkAssign(CL.mkSubscript(CL.mkVar "selfOut",CL.mkStr "x"),
875                                                                           CL.E_Var outState)]                                                                           CL.mkVar outState)]
876                          else let                          else let
877                                  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")
878                                  in                                  in
879                                          [CL.mkAssign(CL.mkSubscript(CL.E_Var "selfIn",index),                                          [CL.mkAssign(CL.mkSubscript(CL.mkVar "selfIn",index),
880                                                                          CL.E_Var inState),                                                                          CL.mkVar inState),
881                                           CL.mkAssign(CL.mkSubscript(CL.E_Var "selfOut",index),                                           CL.mkAssign(CL.mkSubscript(CL.mkVar "selfOut",index),
882                                                                          CL.E_Var outState)]                                                                          CL.mkVar outState)]
883                                  end                                  end
884    
885    
886                    val status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy))))                    val status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy))))
887                    val local_vars =  thread_ids @ initKernelGlobals(!globals,!imgGlobals)  @ strandDecl @ strandObjects @ [status]                    val local_vars =  thread_ids @ initKernelGlobals(!globals,!imgGlobals)  @ strandDecl @ strandObjects @ [status]
888                    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))                    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))
889                    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 while_body = [CL.mkAssign(CL.mkVar "status", CL.mkApply(RN.strandUpdate name,[ CL.mkUnOp(CL.%&,CL.mkVar inState), CL.mkUnOp(CL.%&,CL.mkVar outState)])),
890                                                          CL.mkCall(RN.strandStabilize name,[ CL.E_UnOp(CL.%&,CL.E_Var inState),  CL.E_UnOp(CL.%&,CL.E_Var outState)])]                                                          CL.mkCall(RN.strandStabilize name,[ CL.mkUnOp(CL.%&,CL.mkVar inState),  CL.mkUnOp(CL.%&,CL.mkVar outState)])]
891    
892                    val whileBlock = [CL.mkWhile(while_exp,CL.mkBlock while_body)]                    val whileBlock = [CL.mkWhile(while_exp,CL.mkBlock while_body)]
893    
# Line 879  Line 907 
907        (* generate the table of strand descriptors *)        (* generate the table of strand descriptors *)
908          fun genStrandTable (ppStrm, strands) = let          fun genStrandTable (ppStrm, strands) = let
909                val nStrands = length strands                val nStrands = length strands
910                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.mkVar(RN.strandDesc name)))
911                fun genInits (_, []) = []                fun genInits (_, []) = []
912                  | genInits (i, s::ss) = (i, genInit s) :: genInits(i+1, ss)                  | genInits (i, s::ss) = (i, genInit s) :: genInits(i+1, ss)
913                fun ppDecl dcl = PrintAsC.output(ppStrm, dcl)                fun ppDecl dcl = PrintAsC.output(ppStrm, dcl)
914                in                in
915                  ppDecl (CL.D_Var([], CL.int32, RN.numStrands,                  ppDecl (CL.D_Var([], CL.int32, RN.numStrands,
916                    SOME(CL.I_Exp(CL.E_Int(IntInf.fromInt nStrands, CL.int32)))));                    SOME(CL.I_Exp(CL.mkInt(IntInf.fromInt nStrands, CL.int32)))));
917                  ppDecl (CL.D_Var([],                  ppDecl (CL.D_Var([],
918                    CL.T_Array(CL.T_Ptr(CL.T_Named RN.strandDescTy), SOME nStrands),                    CL.T_Array(CL.T_Ptr(CL.T_Named RN.strandDescTy), SOME nStrands),
919                    RN.strands,                    RN.strands,
# Line 902  Line 930 
930                val clppStrm = PrintAsC.new clOutS                val clppStrm = PrintAsC.new clOutS
931                val cppStrm = PrintAsC.new cOutS                val cppStrm = PrintAsC.new cOutS
932                fun cppDecl dcl = PrintAsC.output(cppStrm, dcl)                fun cppDecl dcl = PrintAsC.output(cppStrm, dcl)
933                fun clppDecl dcl = PrintAsC.output(clppStrm, dcl)                fun clppDecl dcl = PrintAsCL.output(clppStrm, dcl)
934                val strands = AtomTable.listItems strands                val strands = AtomTable.listItems strands
935                val [strand as Strand{name, tyName, code,init_code, ...}] = strands                val [strand as Strand{name, tyName, code,init_code, ...}] = strands
936                in                in
# Line 937  Line 965 
965                  List.app cppDecl (List.rev (!topDecls));                  List.app cppDecl (List.rev (!topDecls));
966                  cppDecl (genHostSetupFunc (strand, clFileName, !numDims, initially, imgGlobals));                  cppDecl (genHostSetupFunc (strand, clFileName, !numDims, initially, imgGlobals));
967                  PrintAsC.close cppStrm;                  PrintAsC.close cppStrm;
968                  PrintAsC.close clppStrm;                  PrintAsCL.close clppStrm;
969                  TextIO.closeOut cOutS;                  TextIO.closeOut cOutS;
970                  TextIO.closeOut clOutS                  TextIO.closeOut clOutS
971                end                end

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

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