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 1286, Tue Jun 7 10:54:18 2011 UTC revision 1305, Sat Jun 11 00:29:00 2011 UTC
# Line 13  Line 13 
13      structure CL = CLang      structure CL = CLang
14      structure RN = RuntimeNames      structure RN = RuntimeNames
15      structure ToCL = TreeToCL      structure ToCL = TreeToCL
16        structure N = CNames
17    
18      (* variable translation *)
19        structure TrVar =
20          struct
21            type env = CL.typed_var TreeIL.Var.Map.map
22            fun lookup (env, x) = (case V.Map.find (env, x)
23                   of SOME(CL.V(_, x')) => x'
24                    | NONE => raise Fail(concat["lookup(_, ", V.name x, ")"])
25                  (* end case *))
26          (* translate a variable that occurs in an l-value context (i.e., as the target of an assignment) *)
27            fun lvalueVar (env, x) = (case V.kind x
28                   of IL.VK_Global => CL.mkVar(lookup(env, x))
29                    | IL.VK_State strand => CL.mkIndirect(CL.mkVar "selfOut", lookup(env, x))
30                    | IL.VK_Local => CL.mkVar(lookup(env, x))
31                  (* end case *))
32          (* translate a variable that occurs in an r-value context *)
33            fun rvalueVar (env, x) = (case V.kind x
34                   of IL.VK_Global => CL.mkVar(lookup(env, x))
35                    | IL.VK_State strand => CL.mkIndirect(CL.mkVar "selfIn", lookup(env, x))
36                    | IL.VK_Local => CL.mkVar(lookup(env, x))
37                  (* end case *))
38          end
39    
40            structure ToC = TreeToCFn (TrVar)
41    
42    (* C variable translation *)    (* C variable translation *)
43      structure TrCVar =      structure TrCVar =
# Line 61  Line 86 
86          double : bool,                  (* true for double-precision support *)          double : bool,                  (* true for double-precision support *)
87          parallel : bool,                (* true for multithreaded (or multi-GPU) target *)          parallel : bool,                (* true for multithreaded (or multi-GPU) target *)
88          debug : bool,                   (* true for debug support in executable *)          debug : bool,                   (* true for debug support in executable *)
89          globals : CL.decl list ref,          globals : {target:TargetUtil.target, globalTy:CL.ty, name:CLang.var} list ref,
90          topDecls : CL.decl list ref,          topDecls : CL.decl list ref,
91          strands : strand AtomTable.hash_table,          strands : strand AtomTable.hash_table,
92          initially : CL.stm list ref,          initially :  CL.decl ref,
93          numDims: int ref,          numDims: int ref,
94          imgGlobals: (string * int) list ref,          imgGlobals: (string * int) list ref,
95          prFn: CL.decl ref          prFn: CL.decl ref
# Line 125  Line 150 
150      structure Var =      structure Var =
151        struct        struct
152          fun name (ToCL.V(_, name)) = name          fun name (ToCL.V(_, name)) = name
153          fun global (Prog{globals, imgGlobals, ...}, name, ty) = let          fun global (Prog{globals, imgGlobals, ...}, global_name, ty) = let
154                val ty' = ToCL.trType ty                val cl_ty  = ToCL.trType ty
155                                      val c_ty = ToC.trType ty
156                fun isImgGlobal (imgGlobals, Ty.ImageTy(ImageInfo.ImgInfo{dim, ...}), name) =  imgGlobals  := (name,dim):: !imgGlobals                fun isImgGlobal (imgGlobals, Ty.ImageTy(ImageInfo.ImgInfo{dim, ...}), name) =  imgGlobals  := (name,dim):: !imgGlobals
157                  | isImgGlobal (imgGlobals, _, _) =  ()                  | isImgGlobal (imgGlobals, _, _) =  ()
158                in                in
159                  globals := CL.D_Var([], ty', name, NONE) :: !globals;                  globals := {target =TargetUtil.TARGET_CL,globalTy = cl_ty, name = global_name} :: !globals;
160                  isImgGlobal(imgGlobals,ty,name);                                           globals := {target =TargetUtil.TARGET_C, globalTy = c_ty, name = global_name} :: !globals;
161                  ToCL.V(ty', name)                  isImgGlobal(imgGlobals,ty,global_name);
162                    ToCL.V(cl_ty, global_name)
163                end                end
164          fun param x = ToCL.V(ToCL.trType(V.ty x), V.name x)          fun param x = ToCL.V(ToCL.trType(V.ty x), V.name x)
165          fun state (Strand{state, ...}, x) = let          fun state (Strand{state, ...}, x) = let
# Line 179  Line 206 
206                    globals = ref [],                    globals = ref [],
207                    topDecls = ref [],                    topDecls = ref [],
208                    strands = AtomTable.mkTable (16, Fail "strand table"),                    strands = AtomTable.mkTable (16, Fail "strand table"),
209                    initially = ref([CL.S_Comment["missing initially"]]),                initially = ref(CL.D_Comment["missing initially"]),
210                                    numDims = ref(0),                                    numDims = ref(0),
211                                    imgGlobals = ref[],                                    imgGlobals = ref[],
212                                    prFn = ref(CL.D_Comment(["No Print Function"]))                                    prFn = ref(CL.D_Comment(["No Print Function"]))
213                  })                  })
214        (* register the global initialization part of a program *)        (* register the global initialization part of a program *)
215            fun globalIndirects (globals,stms) = let            fun globalIndirects (globals,stms) = let
216                  fun getGlobals (CL.D_Var(_,_,globalVar,_)::rest) =                  fun getGlobals ({name,target as TargetUtil.TARGET_CL}::rest) =
217                        CL.mkAssign(CL.mkIndirect(CL.mkVar RN.globalsVarName,globalVar),CL.mkVar globalVar)                        CL.mkAssign(CL.mkIndirect(CL.mkVar RN.globalsVarName,name),CL.mkVar name)
218                          ::getGlobals rest                          ::getGlobals rest
219                    | getGlobals [] = []                    | getGlobals [] = []
220                    | getGlobals (_::rest) = getGlobals rest                    | getGlobals (_::rest) = getGlobals rest
# Line 218  Line 245 
245                in                in
246                  topDecls := shutdownFn :: initFn :: !topDecls                  topDecls := shutdownFn :: initFn :: !topDecls
247                end                end
   
248        (* create and register the initially function for a program *)        (* create and register the initially function for a program *)
249          fun initially {          fun initially {
250                prog = Prog{strands, initially, numDims,...},                prog = Prog{name=progName, strands, initially, ...},
251                isArray : bool,                isArray : bool,
252                iterPrefix : stm list,                iterPrefix : stm list,
253                iters : (var * exp * exp) list,                iters : (var * exp * exp) list,
# Line 231  Line 257 
257              } = let              } = let
258                val name = Atom.toString strand                val name = Atom.toString strand
259                val nDims = List.length iters                val nDims = List.length iters
260                  val worldTy = CL.T_Ptr(CL.T_Named N.worldTy)
261                fun mapi f xs = let                fun mapi f xs = let
262                      fun mapf (_, []) = []                      fun mapf (_, []) = []
263                        | mapf (i, x::xs) = f(i, x) :: mapf(i+1, xs)                        | mapf (i, x::xs) = f(i, x) :: mapf(i+1, xs)
# Line 239  Line 266 
266                      end                      end
267                val baseInit = mapi (fn (i, (_, e, _)) => (i, CL.I_Exp e)) iters                val baseInit = mapi (fn (i, (_, e, _)) => (i, CL.I_Exp e)) iters
268                val sizeInit = mapi                val sizeInit = mapi
269                      (fn (i, (ToCL.V(ty, _), lo, hi)) =>                      (fn (i, (CL.V(ty, _), lo, hi)) =>
270                          (i, CL.I_Exp(CL.mkBinOp(CL.mkBinOp(hi, CL.#-, lo), CL.#+, CL.mkInt(1, ty))))                          (i, CL.I_Exp(CL.mkBinOp(CL.mkBinOp(hi, CL.#-, lo), CL.#+, CL.E_Int(1, ty))))
271                      ) iters                      ) iters
272                    val numStrandsVar = "numStrandsVar"              (* code to allocate the world and initial strands *)
273                val allocCode = iterPrefix @ [                val wrld = "wrld"
274                  val allocCode = [
275                        CL.mkComment["allocate initial block of strands"],                        CL.mkComment["allocate initial block of strands"],
276                        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)),
277                        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)),
278                        CL.mkDecl(CL.int32,"numDims",SOME(CL.I_Exp(CL.mkInt(IntInf.fromInt nDims, CL.int32))))                        CL.mkDecl(worldTy, wrld,
279                            SOME(CL.I_Exp(CL.E_Apply(N.allocInitially, [
280                                CL.mkVar "ProgramName",
281                                CL.mkUnOp(CL.%&, CL.E_Var(N.strandDesc name)),
282                                CL.E_Bool isArray,
283                                CL.E_Int(IntInf.fromInt nDims, CL.int32),
284                                CL.E_Var "base",
285                                CL.E_Var "size"
286                              ]))))
287                      ]                      ]
288                val numStrandsLoopBody =              (* create the loop nest for the initially iterations *)
289                      CL.mkExpStm(CL.mkAssignOp(CL.mkVar numStrandsVar, CL.*=,CL.mkSubscript(CL.mkVar "size",CL.mkVar "i")))                val indexVar = "ix"
290                val numStrandsLoop =  CL.mkFor([(CL.intTy, "i", CL.mkInt(0,CL.intTy))],                val strandTy = CL.T_Ptr(CL.T_Named(N.strandTy name))
291                      CL.mkBinOp(CL.mkVar "i", CL.#<, CL.mkVar "numDims"),                fun mkLoopNest [] = CL.mkBlock(createPrefix @ [
292                      [CL.mkPostOp(CL.mkVar "i", CL.^++)], numStrandsLoopBody)                        CL.mkDecl(strandTy, "sp",
293                in                          SOME(CL.I_Exp(
294                  numDims := nDims;                            CL.E_Cast(strandTy,
295                  initially := allocCode @ [numStrandsLoop]                            CL.E_Apply(N.inState, [CL.E_Var "wrld", CL.E_Var indexVar]))))),
296                end                        CL.mkCall(N.strandInit name, CL.E_Var "sp" :: args),
297                          CL.mkAssign(CL.E_Var indexVar, CL.mkBinOp(CL.E_Var indexVar, CL.#+, CL.E_Int(1, CL.uint32)))
298                        ])
299        (***** OUTPUT *****)                  | mkLoopNest ((CL.V(ty, param), lo, hi)::iters) = let
300          fun genStrandInit (Strand{name,tyName,state,output,code,...}, nDims) = let                      val body = mkLoopNest iters
               val params = [  
                       CL.PARAM([], CL.T_Ptr(CL.uint32), "sizes" ),  
                       CL.PARAM([], CL.intTy, "width"),  
                       CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "strands")  
                     ]  
               val body = let  
                     fun loopParams 3 = ["x", "y", "k"]  
                       | loopParams 2 = ["x", "y"]  
                       | loopParams 1 = ["x"]  
                       | loopParams _ = raise Fail "genStrandInit: missing size dim"  
                     fun mkLoopNest ([], _, nDims) = if nDims = 1  
                           then CL.mkBlock [  
                               CL.mkCall(RN.strandInit name, [  
                                 CL.mkUnOp(CL.%&,CL.mkSubscript(CL.mkVar "strands",CL.mkStr "x")),  
                                                 CL.mkVar "x"])  
                             ]  
                           else let  
                             val index = CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")  
                             in  
                               CL.mkBlock([CL.mkCall(RN.strandInit name, [CL.mkUnOp(CL.%&,CL.mkSubscript(CL.mkVar "strands",index)),  
                               CL.mkVar "x", CL.mkVar"y"])])  
                             end  
                       | mkLoopNest (param::rest,count,nDims) = let  
                           val body = mkLoopNest (rest, count + 1,nDims)  
301                            in                            in
302                              CL.mkFor(                              CL.mkFor(
303                                  [(CL.intTy, param, CL.mkInt(0,CL.intTy))],                          [(ty, param, lo)],
304                                  CL.mkBinOp(CL.mkVar param, CL.#<, CL.mkSubscript(CL.mkVar "sizes",CL.mkInt(count,CL.intTy))),                          CL.mkBinOp(CL.E_Var param, CL.#<=, hi),
305                                  [CL.mkPostOp(CL.mkVar param, CL.^++)],                          [CL.mkPostOp(CL.E_Var param, CL.^++)],
306                                  body)                                  body)
307                            end                            end
308                  val iterCode = [
309                          CL.mkComment["initially"],
310                          CL.mkDecl(CL.uint32, indexVar, SOME(CL.I_Exp(CL.E_Int(0, CL.uint32)))),
311                          mkLoopNest iters
312                        ]
313                  val body = CL.mkBlock(
314                        iterPrefix @
315                        allocCode @
316                        iterCode @
317                        [CL.mkReturn(SOME(CL.E_Var "wrld"))])
318                  val initFn = CL.D_Func([], worldTy, N.initially, [], body)
319                      in                      in
320                        [mkLoopNest ((loopParams nDims),0,nDims)]                  initially := initFn
                     end  
                 in  
                   CL.D_Func(["static"], CL.voidTy, RN.strandInitSetup, params,CL.mkBlock(body))  
321                  end                  end
322    
323          fun genStrandPrint (Strand{name, tyName, state, output, code,...},nDims) = let        (***** OUTPUT *****)
324            fun genStrandPrint (Strand{name, tyName, state, output, code,...}) = let
325              (* the print function *)              (* the print function *)
326                val prFnName = concat[name, "_print"]                val prFnName = concat[name, "_print"]
327                val prFn = let                val prFn = let
328                      val params = [                      val params = [
329                            CL.PARAM([], CL.T_Ptr(CL.T_Named "FILE"), "outS"),                            CL.PARAM([], CL.T_Ptr(CL.T_Named "FILE"), "outS"),
                           CL.PARAM([], CL.T_Ptr(CL.uint32), "sizes" ),  
                           CL.PARAM([], CL.intTy, "width"),  
330                            CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "self")                            CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "self")
331                          ]                          ]
   
332                     val SOME(ty, x) = !output                     val SOME(ty, x) = !output
333                     val outState = if nDims = 1 then                      val outState = CL.mkIndirect(CL.mkVar "self", x)
                           CL.mkSelect(CL.mkSubscript(CL.mkVar "self",CL.mkVar "x"), x)  
                         else if nDims = 2 then  
                                 CL.mkSelect(CL.mkSubscript(CL.mkVar "self",  
                                    CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")), x)  
   
                         else CL.mkSelect(CL.mkVar "self",x)  
   
334                      val prArgs = (case ty                      val prArgs = (case ty
335                             of Ty.IVecTy 1 => [CL.mkStr(!RN.gIntFormat ^ "\n"), outState]                             of Ty.IVecTy 1 => [CL.E_Str(!N.gIntFormat ^ "\n"), outState]
336                              | Ty.IVecTy d => let                              | Ty.IVecTy d => let
337                                  val fmt = CL.mkStr(                                  val fmt = CL.E_Str(
338                                        String.concatWith " " (List.tabulate(d, fn _ => !RN.gIntFormat))                                        String.concatWith " " (List.tabulate(d, fn _ => !N.gIntFormat))
339                                        ^ "\n")                                        ^ "\n")
340                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))                                  val args = List.tabulate (d, fn i => ToC.ivecIndex(outState, d, i))
341                                  in                                  in
342                                    fmt :: args                                    fmt :: args
343                                  end                                  end
344                              | Ty.TensorTy[] => [CL.mkStr "%f\n", outState]                              | Ty.TensorTy[] => [CL.E_Str "%f\n", outState]
345                              | Ty.TensorTy[d] => let                              | Ty.TensorTy[d] => let
346                                  val fmt = CL.mkStr(                                  val fmt = CL.E_Str(
347                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))
348                                        ^ "\n")                                        ^ "\n")
349                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))                                  val args = List.tabulate (d, fn i => ToC.vecIndex(outState, d, i))
350                                  in                                  in
351                                    fmt :: args                                    fmt :: args
352                                  end                                  end
353                              | _ => raise Fail("genStrand: unsupported output type " ^ Ty.toString ty)                              | _ => raise Fail("genStrand: unsupported output type " ^ Ty.toString ty)
354                            (* end case *))                            (* end case *))
   
                           val body = let  
   
                             fun loopParams (3) =  
                                  "x"::"y"::"k"::[]  
                               | loopParams (2) =  
                                  "x"::"y"::[]  
                               | loopParams (1) =  
                                  "x"::[]  
                               | loopParams (_) =  
                                 raise Fail("genStrandPrint: unsupported output type " ^ Ty.toString ty)  
   
                            fun mkLoopNest ([],_) =  
                                                 CL.mkCall("fprintf", CL.mkVar "outS" :: prArgs)  
                                 | mkLoopNest (param::rest,count) = let  
                                         val body = mkLoopNest (rest, count + 1)  
355                                     in                                     in
356                                                  CL.mkFor(                        CL.D_Func(["static"], CL.voidTy, prFnName, params,
357                                                          [(CL.intTy, param, CL.mkInt(0,CL.intTy))],                          CL.mkCall("fprintf", CL.mkVar "outS" :: prArgs))
                                                 CL.mkBinOp(CL.mkVar param, CL.#<, CL.mkSubscript(CL.mkVar "sizes",CL.mkInt(count,CL.intTy))),  
                                                 [CL.mkPostOp(CL.mkVar param, CL.^++)],  
                                                 body)  
                                    end  
                         in  
                                 [mkLoopNest ((loopParams nDims),0)]  
                         end  
   
                     in  
                       CL.D_Func(["static"], CL.voidTy, prFnName, params,CL.mkBlock(body))  
358                      end                      end
359                in                in
360                                   prFn                                   prFn
361                end                end
362    
363          fun genStrandTyDef (Strand{tyName, state,...}) =          fun genStrandTyDef (Strand{tyName, state,...}) =
364              (* the type declaration for the strand's state struct *)              (* the type declaration for the strand's state struct *)
365                CL.D_StructDef(                CL.D_StructDef(
366                        List.rev (List.map (fn ToCL.V(ty, x) => (ty, x)) (!state)),                        List.rev (List.map (fn CL.V(ty, x) => (ty, x)) (!state)),
367                        tyName)                        tyName)
368    
369    
370          (* generates the load kernel function *)          (* generates the load kernel function *)
371  (* FIXME: this code might be part of the runtime system *)  
         fun genKernelLoader() =  
                 CL.D_Verbatim ( ["/* Loads the Kernel from a file */",  
                                                 "char * loadKernel (const char * filename) {",  
                                                 "struct stat statbuf;",  
                                                 "FILE *fh;",  
                                                 "char *source;",  
                                                 "fh = fopen(filename, \"r\");",  
                                                 "if (fh == 0)",  
                                                 "   return 0;",  
                                                 "stat(filename, &statbuf);",  
                                                 "source = (char *) malloc(statbuf.st_size + 1);",  
                                                 "fread(source, statbuf.st_size, 1, fh);",  
                                                 "fread(source, statbuf.st_size, 1, fh);",  
                                                 "return source;",  
                                                 "}"])  
372  (* generates the opencl buffers for the image data *)  (* generates the opencl buffers for the image data *)
373          fun getGlobalDataBuffers(globals,count,contextVar,errVar) = let          fun getGlobalDataBuffers(globals,contextVar,errVar) = let
374                  val globalBufferDecl =  CL.mkDecl(clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)                  val globalBufferDecl =  CL.mkDecl(clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)
375                  val globalBuffer = CL.mkAssign(CL.mkVar(concat[RN.globalsVarName,"_cl"]), CL.mkApply("clCreateBuffer",                  val globalBuffer = CL.mkAssign(CL.mkVar(concat[RN.globalsVarName,"_cl"]), CL.mkApply("clCreateBuffer",
376                                                                  [CL.mkVar contextVar,                                                                  [CL.mkVar contextVar,
# Line 408  Line 379 
379                                                                  CL.mkVar RN.globalsVarName,                                                                  CL.mkVar RN.globalsVarName,
380                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))
381    
382          fun genDataBuffers([],_,_,_) = []          fun genDataBuffers([],_,_) = []
383            | genDataBuffers((var,nDims)::globals,count,contextVar,errVar) = let            | genDataBuffers((var,nDims)::globals,contextVar,errVar) = let
384          (* FIXME: use CL constructors to  build expressions (not strings) *)          (* FIXME: use CL constructors to  build expressions (not strings) *)
385                    val size = if nDims = 1 then                    val size = if nDims = 1 then
386                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,
# Line 435  Line 406 
406                                                                   CL.mkVar "CL_MEM_COPY_HOST_PTR",                                                                   CL.mkVar "CL_MEM_COPY_HOST_PTR",
407                                                                  size,                                                                  size,
408                                                                  CL.mkIndirect(CL.mkVar var,"data"),                                                                  CL.mkIndirect(CL.mkVar var,"data"),
409                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)])):: genDataBuffers(globals,count + 2,contextVar,errVar)                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)])):: genDataBuffers(globals,contextVar,errVar)
410                  end                  end
411          in          in
412                  [globalBufferDecl] @ [globalBuffer] @ genDataBuffers(globals,count + 2,contextVar,errVar)                  [globalBufferDecl] @ [globalBuffer] @ genDataBuffers(globals,contextVar,errVar)
413          end          end
414    
415    
# Line 446  Line 417 
417          fun genGlobalArguments(globals,count,kernelVar,errVar) = let          fun genGlobalArguments(globals,count,kernelVar,errVar) = let
418          val globalArgument = CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,CL.mkApply("clSetKernelArg",          val globalArgument = CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,CL.mkApply("clSetKernelArg",
419                                                                  [CL.mkVar kernelVar,                                                                  [CL.mkVar kernelVar,
420                                                                   CL.mkInt(count,CL.intTy),                                                                   CL.mkPostOp(CL.E_Var count, CL.^++),
421                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
422                                                                   CL.mkUnOp(CL.%&,CL.mkVar(concat[RN.globalsVarName,"_cl"]))])))                                                                   CL.mkUnOp(CL.%&,CL.mkVar(concat[RN.globalsVarName,"_cl"]))])))
423    
# Line 455  Line 426 
426    
427                  CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=, CL.mkApply("clSetKernelArg",                  CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=, CL.mkApply("clSetKernelArg",
428                                                                  [CL.mkVar kernelVar,                                                                  [CL.mkVar kernelVar,
429                                                                   CL.mkInt(count,CL.intTy),                                                                   CL.mkPostOp(CL.E_Var count, CL.^++),
430                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
431                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffix var))])))::                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffix var))])))::
432    
433                          CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,CL.mkApply("clSetKernelArg",                          CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,CL.mkApply("clSetKernelArg",
434                                                                  [CL.mkVar kernelVar,                                                                  [CL.mkVar kernelVar,
435                                                                   CL.mkInt((count + 1),CL.intTy),                                                                   CL.mkPostOp(CL.E_Var count, CL.^++),
436                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
437                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffixData var))]))):: genDataArguments (globals, count + 2,kernelVar,errVar)                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffixData var))]))):: genDataArguments (globals,count,kernelVar,errVar)
438    
439          in          in
440    
441                  [globalArgument] @ genDataArguments(globals,count + 1,kernelVar,errVar)                  [globalArgument] @ genDataArguments(globals,count,kernelVar,errVar)
442    
443          end          end
444    
445          (* generates the main function of host code *)        (* generates the globals buffers and arguments function *)
446          fun genHostMain() = let          fun genGlobal_Buffers_Args (imgGlobals) = let
               val setupCall = [CL.mkCall(RN.setupFName,[CL.mkVar RN.globalsVarName])]  
               val globalsDecl = CL.mkDecl(  
                     CL.T_Ptr(CL.T_Named RN.globalsTy),  
                     RN.globalsVarName,  
                     SOME(CL.I_Exp(CL.mkApply("malloc", [CL.mkApply("sizeof",[CL.mkVar RN.globalsTy])]))))  
               val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.mkVar RN.globalsVarName])  
               val returnStm = [CL.mkReturn(SOME(CL.mkInt(0,CL.intTy)))]  
               val params = [  
                      CL.PARAM([],CL.intTy, "argc"),  
                      CL.PARAM([],CL.charArrayPtr,"argv")  
                    ]  
               val body = CL.mkBlock([globalsDecl] @ [initGlobalsCall]  @ setupCall @ returnStm)  
               in  
                 CL.D_Func([],CL.intTy,"main",params,body)  
               end  
   
       (* generates the host-side setup function *)  
         fun genHostSetupFunc (strand as Strand{name,tyName,...}, filename, nDims, initially, imgGlobals) = let  
447              (* Delcare opencl setup objects *)              (* Delcare opencl setup objects *)
               val programVar= "program"  
               val kernelVar = "kernel"  
               val cmdVar = "queue"  
               val inStateVar = "selfin"  
               val outStateVar = "selfout"  
               val stateSizeVar= "state_mem_size"  
               val clInstateVar = "clSelfIn"  
               val clOutStateVar = "clSelfOut"  
               val clGlobals = "clGlobals"  
               val sourcesVar = "sources"  
               val contextVar = "context"  
448                val errVar = "err"                val errVar = "err"
449                val imgDataSizeVar = "image_dataSize"                val imgDataSizeVar = "image_dataSize"
               val globalVar = "global_work_size"  
               val localVar = "local_work_size"  
               val clFNVar = "filename"  
               val numStrandsVar = "numStrandsVar"  
               val headerFNVar = "header"  
               val deviceVar = "device"  
               val platformsVar = "platforms"  
               val numPlatformsVar = "num_platforms"  
               val numDevicesVar = "num_devices"  
               val assertStm = CL.mkCall("assert",[CL.mkBinOp(CL.mkVar errVar, CL.#==, CL.mkVar "CL_SUCCESS")])  
450                val params = [                val params = [
451                        CL.PARAM([],CL.T_Named("cl_device_id"), deviceVar)                        CL.PARAM([],CL.T_Named("cl_context"), "context"),
452                      ]                                                           CL.PARAM([],CL.T_Named("cl_kernel"), "kernel"),
453                val declarations = [                                                           CL.PARAM([],CL.T_Named("int"), "argStart")
                     CL.mkDecl(clProgramTy, programVar, NONE),  
                     CL.mkDecl(clKernelTy, kernelVar, NONE),  
                     CL.mkDecl(clCmdQueueTy, cmdVar, NONE),  
                     CL.mkDecl(clContextTy, contextVar, NONE),  
                     CL.mkDecl(CL.intTy, errVar, NONE),  
                     CL.mkDecl(CL.intTy, numStrandsVar, SOME(CL.I_Exp(CL.mkInt(1,CL.intTy)))),  
                     CL.mkDecl(CL.intTy, stateSizeVar, NONE),  
                     CL.mkDecl(CL.intTy, "width", NONE),  
                     CL.mkDecl(CL.intTy, imgDataSizeVar, NONE),  
                     (*CL.mkDecl(clDeviceIdTy, deviceVar, NONE), *)  
                     CL.mkDecl(CL.T_Ptr(CL.T_Named tyName), inStateVar,NONE),  
                     CL.mkDecl(clMemoryTy,clInstateVar,NONE),  
                     CL.mkDecl(clMemoryTy,clOutStateVar,NONE),  
                     CL.mkDecl(CL.T_Ptr(CL.T_Named tyName), outStateVar,NONE),  
                     CL.mkDecl(CL.charPtr, clFNVar,SOME(CL.I_Exp(CL.mkStr filename))),  
 (* FIXME:  use Paths.diderotInclude *)  
                     CL.mkDecl(CL.charPtr, headerFNVar,SOME(CL.I_Exp(CL.mkStr "../src/include/Diderot/cl-types.h"))),  
                     CL.mkDecl(CL.T_Array(CL.charPtr,SOME(2)),sourcesVar,NONE),  
                     CL.mkDecl(CL.T_Array(CL.T_Named "size_t",SOME(nDims)),globalVar,NONE),  
                     CL.mkDecl(CL.T_Array(CL.T_Named "size_t",SOME(nDims)),localVar,NONE),  
                     CL.mkDecl(CL.intTy,numDevicesVar,SOME(CL.I_Exp(CL.mkInt(~1,CL.intTy)))),  
                     CL.mkDecl(CL.T_Array(CL.T_Named "cl_platform_id", SOME(1)), platformsVar, NONE),  
                     CL.mkDecl(CL.intTy,"num_platforms",SOME(CL.I_Exp(CL.mkInt(~1,CL.intTy))))  
454                  ]                  ]
             (* Setup Global Variables *)  
               val globalsDecl = CL.mkDecl(  
                     CL.T_Ptr(CL.T_Named RN.globalsTy),  
                     RN.globalsVarName,  
                     SOME(CL.I_Exp(CL.mkApply("malloc", [CL.mkApply("sizeof",[CL.mkVar RN.globalsTy])]))))  
               val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.mkVar RN.globalsVarName])  
   
                 (* Retrieve the platforms  
                 val platformStm = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetPlatformIDs",  
                                                   [CL.mkInt(10,CL.intTy),  
                                                    CL.mkVar platformsVar,  
                                                    CL.mkUnOp(CL.%&,CL.mkVar numPlatformsVar)])),  
                                                    assertStm]  
   
                 val devicesStm = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetDeviceIDs",  
                                                   [CL.mkSubscript(CL.mkVar platformsVar,CL.mkInt(0,CL.intTy)),  
                                                    CL.mkVar "CL_DEVICE_TYPE_GPU",  
                                                    CL.mkInt(1,CL.intTy),  
                                                    CL.mkUnOp(CL.%&,CL.mkVar deviceVar),  
                                                    CL.mkUnOp(CL.%&,CL.mkVar numDevicesVar)])),  
                                                    assertStm] *)  
   
                 (* Create Context *)  
                 val contextStm = [CL.mkAssign(CL.mkVar contextVar, CL.mkApply("clCreateContext",  
                                                   [CL.mkInt(0,CL.intTy),  
                                                   CL.mkInt(1,CL.intTy),  
                                                   CL.mkUnOp(CL.%&,CL.mkVar deviceVar),  
                                                   CL.mkVar "NULL",  
                                                   CL.mkVar "NULL",  
                                                   CL.mkUnOp(CL.%&,CL.mkVar errVar)])),  
                                                   assertStm]  
455    
                 (* Create Command Queue *)  
                 val commandStm = [CL.mkAssign(CL.mkVar cmdVar, CL.mkApply("clCreateCommandQueue",  
                                                   [CL.mkVar contextVar,  
                                                   CL.mkVar deviceVar,  
                                                   CL.mkInt(0,CL.intTy),  
                                                   CL.mkUnOp(CL.%&,CL.mkVar errVar)])),  
                                                   assertStm]  
   
   
                 (*Create Program/Build/Kernel with Source statement *)  
                 val createProgStm = CL.mkAssign(CL.mkVar programVar, CL.mkApply("clCreateProgramWithSource",  
                                                                                                                 [CL.mkVar contextVar,  
                                                                                                                  CL.mkInt(2,CL.intTy),  
                                                                                                                  CL.mkCast(CL.T_Ptr(CL.T_Named("const char *")),CL.mkUnOp(CL.%&,CL.mkVar sourcesVar)),  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))  
   
                 (* FIXME: Remove after testing purposes, Build Log for OpenCL*)  
                 val buildLog = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clBuildProgram",  
                                                                                                                 [CL.mkVar programVar,  
                                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkVar "NULL"])),  
                                           CL.mkDecl(CL.charPtr, "build", NONE),  
                                           CL.mkDecl(CL.T_Named("size_t"),"ret_val_size",NONE),  
                                            CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetProgramBuildInfo",  
                                                                                                                 [CL.mkVar programVar,  
                                                                                                                 CL.mkVar deviceVar,  
                                                                                                                  CL.mkVar "CL_PROGRAM_BUILD_LOG",  
                                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkUnOp(CL.%&,CL.mkVar "ret_val_size")])),  
                                           CL.mkAssign(CL.mkVar "build", CL.mkApply("malloc", [CL.mkVar "ret_val_size"])),  
                                                 CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetProgramBuildInfo",  
                                                                                                                 [CL.mkVar programVar,  
                                                                                                                 CL.mkVar deviceVar,  
                                                                                                                  CL.mkVar "CL_PROGRAM_BUILD_LOG",  
                                                                                                                  CL.mkVar "ret_val_size",  
                                                                                                                  CL.mkVar "build",  
                                                                                                                  CL.mkVar "NULL"])),  
                                                 CL.mkAssign(CL.mkSubscript(CL.mkVar "build",CL.mkVar "ret_val_size"),CL.mkVar ("'\\" ^ "0'")),  
                                                 CL.mkCall("printf",[CL.mkStr ( "Build Log:" ^ "\n" ^ "%s" ^ "\n"), CL.mkVar "build"])]  
   
   
   
   
                 val createKernel = CL.mkAssign(CL.mkVar kernelVar, CL.mkApply("clCreateKernel",  
                                                                                                                 [CL.mkVar programVar,  
                                                                                                                  CL.mkStr RN.kernelFuncName,  
                                                                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))  
   
   
                 val create_build_stms = [createProgStm,assertStm] @ buildLog @ [assertStm,createKernel,assertStm]  
   
   
   
                 (* Create Memory Buffers for Strand States and Globals *)  
                 val strandSize = CL.mkAssign(CL.mkVar stateSizeVar,CL.mkBinOp(CL.mkApply("sizeof",  
                                                                         [CL.mkVar tyName]), CL.#*,CL.mkVar numStrandsVar))  
   
                 val clStrandObjects = [CL.mkAssign(CL.mkVar clInstateVar, CL.mkApply("clCreateBuffer",  
                                                                 [CL.mkVar contextVar,  
                                                                 CL.mkVar "CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR",  
                                                                 CL.mkVar stateSizeVar,  
                                                                 CL.mkVar "NULL",  
                                                                 CL.mkUnOp(CL.%&,CL.mkVar errVar)])),  
                                                          CL.mkAssign(CL.mkVar clOutStateVar, CL.mkApply("clCreateBuffer",  
                                                                 [CL.mkVar contextVar,  
                                                                 CL.mkVar "CL_MEM_READ_WRITE",  
                                                                 CL.mkVar stateSizeVar,  
                                                                 CL.mkVar "NULL",  
                                                                 CL.mkUnOp(CL.%&,CL.mkVar errVar)]))]  
   
   
                 (* Setup up selfOut variable *)  
                 val strandsArrays = [CL.mkAssign(CL.mkVar outStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.mkVar numStrandsVar,  
                                                                         CL.#*, CL.mkApply("sizeof",[CL.mkVar tyName]))])),  
                                                                 CL.mkAssign(CL.mkVar inStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.mkVar numStrandsVar,  
                                                                         CL.#*, CL.mkApply("sizeof",[CL.mkVar tyName]))]))]  
   
   
                 (* Initialize Width Parameter *)  
                 val widthDel = if nDims = 2 then  
                           CL.mkAssign(CL.mkVar "width",CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1, CL.intTy)))  
                    else  
                           CL.mkAssign(CL.mkVar "width",CL.mkInt(0,CL.intTy))  
   
   
                 val strands_init = CL.mkCall(RN.strandInitSetup, [  
                         CL.mkVar "size", CL.mkVar "width", CL.mkVar inStateVar  
                       ])  
   
             val clGlobalBuffers = getGlobalDataBuffers(!imgGlobals,3,contextVar,errVar)  
   
   
                 (* Load the Kernel and Header Files *)  
                 val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(1,CL.intTy)),  
                                                                           CL.mkApply(RN.clLoaderFN, [CL.mkVar clFNVar])),  
            CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(0,CL.intTy)),  
                                                                           CL.mkApply(RN.clLoaderFN, [CL.mkVar headerFNVar]))]  
   
                 (* val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(1,CL.intTy)),  
                                                                           CL.mkApply(RN.clLoaderFN, [CL.mkVar clFNVar]))] *)  
   
   
                 (* Created Enqueue Statements *)  
 (* FIXME: simplify this code by function abstraction *)  
         val enqueueStm = if nDims = 1  
                         then [CL.mkAssign(CL.mkVar errVar,  
                                                           CL.mkApply("clEnqueueNDRangeKernel",  
                                                                                                 [CL.mkVar cmdVar,  
                                                                                                  CL.mkVar kernelVar,  
                                                                                                  CL.mkInt(1,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar globalVar,  
                                                                                                  CL.mkVar localVar,  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]  
                         else if nDims = 2  then  
                          [CL.mkAssign(CL.mkVar errVar,  
                                                         CL.mkApply("clEnqueueNDRangeKernel",  
                                                                                                 [CL.mkVar cmdVar,  
                                                                                                  CL.mkVar kernelVar,  
                                                                                                  CL.mkInt(2,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar globalVar,  
                                                                                                  CL.mkVar localVar,  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]  
                         else  
                           [CL.mkAssign(CL.mkVar errVar,  
                                                         CL.mkApply("clEnqueueNDRangeKernel",  
                                                                                                 [CL.mkVar cmdVar,  
                                                                                                  CL.mkVar kernelVar,  
                                                                                                  CL.mkInt(3,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar globalVar,  
                                                                                                  CL.mkVar localVar,  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]  
   
   
   
                 (* Setup Global and Local variables *)  
   
                 val globalAndlocalStms = if nDims = 1 then  
                         [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),  
                          CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),  
                                                                   CL.mkVar "16")]  
   
   
                 else if nDims = 2 then  
                         [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(1,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),  
                                                                   CL.mkVar "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(1,CL.intTy)),  
                                                                   CL.mkVar "16")]  
   
                 else  
                         [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(1,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(2,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(2,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),  
                                                                   CL.mkVar "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(1,CL.intTy)),  
                                                                   CL.mkVar "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(2,CL.intTy)),  
                                                                   CL.mkVar "16")]  
   
   
   
                 (* Setup Kernel arguments *)  
                 val kernelArguments = [CL.mkAssign(CL.mkVar errVar,CL.mkApply("clSetKernelArg",  
                                                                 [CL.mkVar kernelVar,  
                                                                  CL.mkInt(0,CL.intTy),  
                                                                  CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),  
                                                                  CL.mkUnOp(CL.%&,CL.mkVar clInstateVar)])),  
                                                             CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar, CL.|=,CL.mkApply("clSetKernelArg",  
                                                                 [CL.mkVar kernelVar,  
                                                                  CL.mkInt(1,CL.intTy),  
                                                                  CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),  
                                                                  CL.mkUnOp(CL.%&,CL.mkVar clOutStateVar)]))),  
                                                                   CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar, CL.|=,CL.mkApply("clSetKernelArg",  
                                                                 [CL.mkVar kernelVar,  
                                                                  CL.mkInt(2,CL.intTy),  
                                                                  CL.mkApply("sizeof",[CL.mkVar "int"]),  
                                                                  CL.mkUnOp(CL.%&,CL.mkVar "width")])))]  
   
            val clGlobalArguments = genGlobalArguments(!imgGlobals,3,kernelVar,errVar) @ [assertStm]  
   
                 (* Retrieve output *)  
                 val outputStm = CL.mkAssign(CL.mkVar errVar,  
                                                         CL.mkApply("clEnqueueReadBuffer",  
                                                                                                 [CL.mkVar cmdVar,  
                                                                                                  CL.mkVar clOutStateVar,  
                                                                                                  CL.mkVar "CL_TRUE",  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar stateSizeVar,  
                                                                                                  CL.mkVar outStateVar,  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar "NULL"]))  
   
                 (* Free all the objects *)  
                 val freeStms = [CL.mkCall("clReleaseKernel",[CL.mkVar kernelVar]),  
                                                 CL.mkCall("clReleaseProgram",[CL.mkVar programVar ]),  
                                                 CL.mkCall("clReleaseCommandQueue",[CL.mkVar cmdVar]),  
                                                 CL.mkCall("clReleaseContext",[CL.mkVar contextVar]),  
                                                 CL.mkCall("clReleaseMemObject",[CL.mkVar clInstateVar]),  
                                                 CL.mkCall("clReleaseMemObject",[CL.mkVar clOutStateVar])]  
   
   
                 (*Setup Strand Print Function *)  
                 val outputData = [CL.mkDecl(CL.T_Ptr(CL.T_Named("FILE")), "outS", SOME(CL.I_Exp(CL.mkApply("fopen",  
                                                 [CL.mkStr "mip.txt",  
                                                 CL.mkStr "w"])))),  
                                                 CL.mkCall(concat[name, "_print"],  
                                                                         [CL.mkVar "outS",  
                                                                          CL.mkVar "size",  
                                                                          CL.mkVar "width",  
                                                                          CL.mkVar outStateVar])]  
456    
457                val clGlobalBuffers = getGlobalDataBuffers(!imgGlobals, "context","err")
458    
459                val clGlobalArguments = genGlobalArguments(!imgGlobals,"argStart","kernel","err")
460    
461                  (* Body put all the statments together *)                  (* Body put all the statments together *)
462                  val body =  declarations @ [globalsDecl,initGlobalsCall] (*@ platformStm @ devicesStm *) @ contextStm @ commandStm @ !initially @ [strandSize] @                  val body =  clGlobalBuffers @ clGlobalArguments
                                    strandsArrays @ globalAndlocalStms @ [widthDel,strands_init]  @ clStrandObjects @ clGlobalBuffers @ sourceStms  @ create_build_stms  (*@  
                                    kernelArguments @ clGlobalArguments @ enqueueStm @  [outputStm] @ freeStms @ outputData *)  
   
463                  in                  in
464                            CL.D_Func([],CL.voidTy,RN.globalsSetupName,params,CL.mkBlock(body))
                 CL.D_Func([],CL.voidTy,RN.setupFName,params,CL.mkBlock(body))  
   
465                  end                  end
466    
467  (* generate the data and global parameters *)  (* generate the data and global parameters *)
468          fun genKeneralGlobalParams ((name,tyname)::rest) =          fun genKeneralGlobalParams ((name,tyname)::rest) =
469                  CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), concat[RN.globalsVarName]) ::                  CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), concat[RN.globalsVarName]) ::
# Line 832  Line 475 
475    
476          (*generate code for intilizing kernel global data *)          (*generate code for intilizing kernel global data *)
477          fun initKernelGlobals (globals,imgGlobals) = let          fun initKernelGlobals (globals,imgGlobals) = let
478                  fun initGlobalStruct (CL.D_Var(_, _ , name, _)::rest) =                  fun initGlobalStruct ({name,target as TargetUtil.TARGET_CL,globalTy}::rest) =
479                                  CL.mkAssign(CL.mkVar name, CL.mkIndirect(CL.mkVar RN.globalsVarName, name)) ::                                  CL.mkAssign(CL.mkVar name, CL.mkIndirect(CL.mkVar RN.globalsVarName, name)) ::
480                                  initGlobalStruct(rest)                                  initGlobalStruct(rest)
481                    | initGlobalStruct ( _::rest) = initGlobalStruct(rest)                    | initGlobalStruct ( _::rest) = initGlobalStruct(rest)
# Line 869  Line 512 
512                    val strandDecl = [CL.mkDecl(CL.T_Named tyName, inState, NONE),                    val strandDecl = [CL.mkDecl(CL.T_Named tyName, inState, NONE),
513                                                          CL.mkDecl(CL.T_Named tyName, outState,NONE)]                                                          CL.mkDecl(CL.T_Named tyName, outState,NONE)]
514                    val strandObjects  = if nDims = 1                    val strandObjects  = if nDims = 1
515                          then [CL.mkAssign(CL.mkSubscript(CL.mkVar "selfIn",CL.mkStr "x"),                          then [CL.mkAssign( CL.mkVar inState, CL.mkSubscript(CL.mkVar "selfIn",CL.mkStr "x")),
516                                                                           CL.mkVar inState),                                    CL.mkAssign(CL.mkVar outState,CL.mkSubscript(CL.mkVar "selfOut",CL.mkStr "x"))]
                                   CL.mkAssign(CL.mkSubscript(CL.mkVar "selfOut",CL.mkStr "x"),  
                                                                          CL.mkVar outState)]  
517                          else let                          else let
518                                  val index = CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")                                  val index = CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")
519                                  in                                  in
520                                          [CL.mkAssign(CL.mkSubscript(CL.mkVar "selfIn",index),                                          [CL.mkAssign(CL.mkVar inState, CL.mkSubscript(CL.mkVar "selfIn",index)),
521                                                                          CL.mkVar inState),                                           CL.mkAssign(CL.mkVar outState,CL.mkSubscript(CL.mkVar "selfOut",index))]
                                          CL.mkAssign(CL.mkSubscript(CL.mkVar "selfOut",index),  
                                                                         CL.mkVar outState)]  
522                                  end                                  end
523    
524    
# Line 896  Line 535 
535                     CL.D_Func(["__kernel"], CL.voidTy, fName, params, body)                     CL.D_Func(["__kernel"], CL.voidTy, fName, params, body)
536                  end                  end
537          (* generate a global structure from the globals *)          (* generate a global structure from the globals *)
538          fun genGlobalStruct(globals) = let          fun genGlobalStruct(target,globals) = let
539                   fun getGlobals(CL.D_Var(_,ty,globalVar,_)::rest) = (ty,globalVar)::getGlobals(rest)                   fun getGlobals(_, []) = []
540                     | getGlobals([]) = []                          | getGlobals(target',{name,globalTy,target}::rest) =
541                     | getGlobals(_::rest) = getGlobals(rest)                                  if target = target' then
542                   in                                  (globalTy,name)::getGlobals(target',rest)
543                          CL.D_StructDef(getGlobals(globals),RN.globalsTy)                                  else
544                    end                                  getGlobals(target',rest)
545                     in
546        (* generate the table of strand descriptors *)                          CL.D_StructDef(getGlobals(target,globals),RN.globalsTy)
         fun genStrandTable (ppStrm, strands) = let  
               val nStrands = length strands  
               fun genInit (Strand{name, ...}) = CL.I_Exp(CL.mkUnOp(CL.%&, CL.mkVar(RN.strandDesc name)))  
               fun genInits (_, []) = []  
                 | genInits (i, s::ss) = (i, genInit s) :: genInits(i+1, ss)  
               fun ppDecl dcl = PrintAsC.output(ppStrm, dcl)  
               in  
                 ppDecl (CL.D_Var([], CL.int32, RN.numStrands,  
                   SOME(CL.I_Exp(CL.mkInt(IntInf.fromInt nStrands, CL.int32)))));  
                 ppDecl (CL.D_Var([],  
                   CL.T_Array(CL.T_Ptr(CL.T_Named RN.strandDescTy), SOME nStrands),  
                   RN.strands,  
                   SOME(CL.I_Array(genInits (0, strands)))))  
547                end                end
548    
549       fun genGlobals(_,_, []) =
550                            ()
551              | genGlobals(declFun, target',{name,globalTy,target}::rest) =
552                            if target = target' then
553                               (declFun (CL.D_Var([], globalTy, name, NONE));
554                               genGlobals (declFun,target',rest))
555                            else
556                               genGlobals (declFun,target',rest)
557    
558          fun genSrc (baseName, Prog{double,globals, topDecls, strands, initially,imgGlobals,numDims,...}) = let          fun genSrc (baseName, Prog{double,globals, topDecls, strands, initially,imgGlobals,numDims,...}) = let
559                val clFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "cl"}                val clFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "cl"}
# Line 927  Line 561 
561                val clOutS = TextIO.openOut clFileName                val clOutS = TextIO.openOut clFileName
562                val cOutS = TextIO.openOut cFileName                val cOutS = TextIO.openOut cFileName
563  (* FIXME: need to use PrintAsC and PrintAsCL *)  (* FIXME: need to use PrintAsC and PrintAsCL *)
564                val clppStrm = PrintAsC.new clOutS                val clppStrm = PrintAsCL.new clOutS
565                val cppStrm = PrintAsC.new cOutS                val cppStrm = PrintAsC.new cOutS
566                fun cppDecl dcl = PrintAsC.output(cppStrm, dcl)                fun cppDecl dcl = PrintAsC.output(cppStrm, dcl)
567                fun clppDecl dcl = PrintAsCL.output(clppStrm, dcl)                fun clppDecl dcl = PrintAsCL.output(clppStrm, dcl)
# Line 940  Line 574 
574                        then "#define DIDEROT_DOUBLE_PRECISION"                        then "#define DIDEROT_DOUBLE_PRECISION"
575                        else "#define DIDEROT_SINGLE_PRECISION",                        else "#define DIDEROT_SINGLE_PRECISION",
576                      "#define DIDEROT_TARGET_CL",                      "#define DIDEROT_TARGET_CL",
577                      "#include \"Diderot/cl-types.h\""                      "#include \"Diderot/cl-diderot.h\""
578                    ]));                    ]));
579                  List.app clppDecl (List.rev (!globals));                  genGlobals(clppDecl,TargetUtil.TARGET_CL,!globals);
580                  clppDecl (genGlobalStruct (!globals));                  clppDecl (genGlobalStruct (TargetUtil.TARGET_CL,!globals));
581                  clppDecl (genStrandTyDef strand);                  clppDecl (genStrandTyDef (strand));
582                  List.app clppDecl (!code);                  List.app clppDecl (!code);
583                  clppDecl (genKernelFun (strand,!numDims,globals,imgGlobals));                  clppDecl (genKernelFun (strand,!numDims,globals,imgGlobals));
584    
585                (* Generate the Host file .c *)                (* Generate the Host file .c *)
586                  cppDecl (CL.D_Verbatim([                  cppDecl (CL.D_Verbatim([
587                      if double                      if double
# Line 955  Line 590 
590                      "#define DIDEROT_TARGET_CL",                      "#define DIDEROT_TARGET_CL",
591                      "#include \"Diderot/diderot.h\""                      "#include \"Diderot/diderot.h\""
592                    ]));                    ]));
593                  List.app cppDecl (List.rev (!globals));                  genGlobals(cppDecl,TargetUtil.TARGET_C,!globals);
594                  cppDecl (genGlobalStruct (!globals));                  cppDecl (genGlobalStruct (TargetUtil.TARGET_C,!globals));
595                  cppDecl (genStrandTyDef strand);                  cppDecl (genStrandTyDef (strand));
596                  cppDecl  (!init_code);                  cppDecl  (!init_code);
597                  cppDecl (genStrandInit(strand,!numDims));                  cppDecl (genStrandPrint(strand));
                 cppDecl (genStrandPrint(strand,!numDims));  
                 (* cppDecl (genKernelLoader());*)  
598                  List.app cppDecl (List.rev (!topDecls));                  List.app cppDecl (List.rev (!topDecls));
599                  cppDecl (genHostSetupFunc (strand, clFileName, !numDims, initially, imgGlobals));                  cppDecl (genGlobal_Buffers_Args (imgGlobals));
600                    cppDecl (!initially);
601                  PrintAsC.close cppStrm;                  PrintAsC.close cppStrm;
602                  PrintAsCL.close clppStrm;                  PrintAsCL.close clppStrm;
603                  TextIO.closeOut cOutS;                  TextIO.closeOut cOutS;

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

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