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

SCM Repository

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

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

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

revision 1278, Mon Jun 6 16:27:28 2011 UTC revision 1307, Sat Jun 11 13:58:02 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      type var = ToCL.var    (* 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 *)
43        structure TrCVar =
44          struct
45            type env = CL.typed_var TreeIL.Var.Map.map
46            fun lookup (env, x) = (case V.Map.find (env, x)
47                   of SOME(CL.V(_, x')) => x'
48                    | NONE => raise Fail(concat["TrCVar.lookup(_, ", V.name x, ")"])
49                  (* end case *))
50          (* translate a variable that occurs in an l-value context (i.e., as the target of an assignment) *)
51            fun lvalueVar (env, x) = (case V.kind x
52                   of IL.VK_Global => CL.mkIndirect(CL.mkVar RN.globalsVarName, lookup(env, x))
53                    | IL.VK_State strand => raise Fail "unexpected strand context"
54                    | IL.VK_Local => CL.mkVar(lookup(env, x))
55                  (* end case *))
56          (* translate a variable that occurs in an r-value context *)
57            val rvalueVar = lvalueVar
58          end
59    
60        structure ToC = TreeToCFn (TrCVar)
61    
62        type var = CL.typed_var
63      type exp = CL.exp      type exp = CL.exp
64      type stm = CL.stm      type stm = CL.stm
65    
66      (* OpenCL specific types *)
67        val clProgramTy = CL.T_Named "cl_program"
68        val clKernelTy  = CL.T_Named "cl_kernel"
69        val clCmdQueueTy = CL.T_Named "cl_command_queue"
70        val clContextTy = CL.T_Named "cl_context"
71        val clDeviceIdTy = CL.T_Named "cl_device_id"
72        val clPlatformIdTy = CL.T_Named "cl_platform_id"
73        val clMemoryTy = CL.T_Named "cl_mem"
74    
75      (* variable or field that is mirrored between host and GPU *)
76        type mirror_var = {
77                hostTy : CL.ty,             (* variable type on Host (i.e., C type) *)
78                gpuTy : CL.ty,              (* variable's type on GPU (i.e., OpenCL type) *)
79                var : CL.var                (* variable name *)
80              }
81    
82      datatype strand = Strand of {      datatype strand = Strand of {
83          name : string,          name : string,
84          tyName : string,          tyName : string,
85          state : var list ref,          state : mirror_var list ref,
86          output : (Ty.ty * CL.var) option ref,   (* the strand's output variable (only one for now) *)          output : (Ty.ty * CL.var) option ref,   (* the strand's output variable (only one for now) *)
87          code : CL.decl list ref,          code : CL.decl list ref,
88          init_code: CL.decl ref          init_code: CL.decl ref
# Line 32  Line 93 
93          double : bool,                  (* true for double-precision support *)          double : bool,                  (* true for double-precision support *)
94          parallel : bool,                (* true for multithreaded (or multi-GPU) target *)          parallel : bool,                (* true for multithreaded (or multi-GPU) target *)
95          debug : bool,                   (* true for debug support in executable *)          debug : bool,                   (* true for debug support in executable *)
96          globals : CL.decl list ref,          globals : mirror_var list ref,
97          topDecls : CL.decl list ref,          topDecls : CL.decl list ref,
98          strands : strand AtomTable.hash_table,          strands : strand AtomTable.hash_table,
99          initially : CL.stm list ref,          initially :  CL.decl ref,
100          numDims: int ref,          numDims: int ref,
101          imgGlobals: (string * int) list ref,          imgGlobals: (string * int) list ref,
102          prFn: CL.decl ref          prFn: CL.decl ref
# Line 69  Line 130 
130    (* TreeIL to target translations *)    (* TreeIL to target translations *)
131      structure Tr =      structure Tr =
132        struct        struct
133          (* this function is used for the initially clause, so it generates OpenCL *)
134          fun fragment (ENV{info, vMap, scope}, blk) = let          fun fragment (ENV{info, vMap, scope}, blk) = let
135                val (vMap, stms) = ToCL.trFragment (vMap, blk)                val (vMap, stms) = ToCL.trFragment (vMap, blk)
136                in                in
# Line 85  Line 147 
147          fun block (ENV{vMap, scope, ...}, blk) = (case scope          fun block (ENV{vMap, scope, ...}, blk) = (case scope
148                 of StrandScope stateVars => ToCL.trBlock (vMap, saveState "StrandScope" stateVars, blk)                 of StrandScope stateVars => ToCL.trBlock (vMap, saveState "StrandScope" stateVars, blk)
149                  | MethodScope stateVars => ToCL.trBlock (vMap, saveState "MethodScope" stateVars, blk)                  | MethodScope stateVars => ToCL.trBlock (vMap, saveState "MethodScope" stateVars, blk)
150                  | _ => ToCL.trBlock (vMap, fn (_, _, stm) => [stm], blk)                  | InitiallyScope => ToCL.trBlock (vMap, fn (_, _, stm) => [stm], blk)
151                    | _ => ToC.trBlock (vMap, fn (_, _, stm) => [stm], blk)
152                (* end case *))                (* end case *))
153          fun exp (ENV{vMap, ...}, e) = ToCL.trExp(vMap, e)          fun exp (ENV{vMap, ...}, e) = ToCL.trExp(vMap, e)
154        end        end
# Line 95  Line 158 
158        struct        struct
159          fun name (ToCL.V(_, name)) = name          fun name (ToCL.V(_, name)) = name
160          fun global (Prog{globals, imgGlobals, ...}, name, ty) = let          fun global (Prog{globals, imgGlobals, ...}, name, ty) = let
161                val ty' = ToCL.trType ty                val x = {hostTy = ToC.trType ty, gpuTy = ToCL.trType ty, var = name}
162                fun isImgGlobal (imgGlobals, Ty.ImageTy(ImageInfo.ImgInfo{dim, ...}), name) =  imgGlobals  := (name,dim):: !imgGlobals                fun isImgGlobal (Ty.ImageTy(ImageInfo.ImgInfo{dim, ...}), name) =
163                  | isImgGlobal (imgGlobals, _, _) =  ()                      imgGlobals  := (name,dim) :: !imgGlobals
164                in                  | isImgGlobal _ =  ()
165                  globals := CL.D_Var([], ty', name, NONE) :: !globals;                in
166                  isImgGlobal(imgGlobals,ty,name);                  globals := x :: !globals;
167               ToCL.V(ty', name)                  isImgGlobal (ty, name);
168                    ToCL.V(#gpuTy x, name)
169                end                end
170          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)
171          fun state (Strand{state, ...}, x) = let          fun state (Strand{state, ...}, x) = let
172                val ty' = ToCL.trType(V.ty x)                val ty = V.ty x
173                val x' = ToCL.V(ty', V.name x)                val x' = {hostTy = ToC.trType ty, gpuTy = ToCL.trType ty, var = V.name x}
174                in                in
175                  state := x' :: !state;                  state := x' :: !state;
176                  x'                  ToCL.V(#gpuTy x', #var x')
177                end                end
178        end        end
179    
# Line 141  Line 205 
205        struct        struct
206          fun new {name, double, parallel, debug} = (          fun new {name, double, parallel, debug} = (
207                RN.initTargetSpec double;                RN.initTargetSpec double;
208                  CNames.initTargetSpec double;
209                Prog{                Prog{
210                    name = name,                    name = name,
211                    double = double, parallel = parallel, debug = debug,                    double = double, parallel = parallel, debug = debug,
212                    globals = ref [],                    globals = ref [],
213                    topDecls = ref [],                    topDecls = ref [],
214                    strands = AtomTable.mkTable (16, Fail "strand table"),                    strands = AtomTable.mkTable (16, Fail "strand table"),
215                    initially = ref([CL.S_Comment["missing initially"]]),                    initially = ref(CL.D_Comment["missing initially"]),
216                                    numDims = ref(0),                                    numDims = ref(0),
217                                    imgGlobals = ref[],                                    imgGlobals = ref[],
218                                    prFn = ref(CL.D_Comment(["No Print Function"]))                                    prFn = ref(CL.D_Comment(["No Print Function"]))
219                  })                  })
220        (* register the global initialization part of a program *)        (* register the global initialization part of a program *)
221            fun globalIndirects (globals,stms) = let            fun globalIndirects (globals,stms) = let
222                  fun getGlobals(CL.D_Var(_,_,globalVar,_)::rest) = CL.mkAssign(CL.mkIndirect(CL.mkVar RN.globalsVarName,globalVar),CL.mkVar globalVar)::getGlobals(rest)                  fun getGlobals ({name,target as TargetUtil.TARGET_CL}::rest) =
223                    | getGlobals([]) = []                        CL.mkAssign(CL.mkIndirect(CL.mkVar RN.globalsVarName,name),CL.mkVar name)
224                    | getGlobals(_::rest) = getGlobals(rest)                          ::getGlobals rest
225                      | getGlobals [] = []
226                      | getGlobals (_::rest) = getGlobals rest
227                  in                  in
228                    stms @ getGlobals(globals)                    stms @ getGlobals globals
229                  end                  end
230    
231        (* 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 171  Line 238 
238                  topDecls := inputsFn :: !topDecls                  topDecls := inputsFn :: !topDecls
239                end                end
240    
241          fun init (Prog{globals, topDecls,...}, CL.S_Block(init)) = let        (* register the global initialization part of a program *)
242                val params = [          fun init (Prog{topDecls, ...}, init) = let
243                        CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName)                val globPtrTy = CL.T_Ptr(CL.T_Named RN.globalsTy)
244                      ]                val initFn = CL.D_Func(
245                val body = CL.S_Block(globalIndirects(!globals,init))                      [], CL.voidTy, RN.initGlobals, [CL.PARAM([], globPtrTy, RN.globalsVarName)],
246                val initFn = CL.D_Func([], CL.voidTy, RN.initGlobals, params, body)                      init)
247                in                val shutdownFn = CL.D_Func(
248                  topDecls := initFn :: !topDecls                      [], CL.voidTy, RN.shutdown,
249                end                      [CL.PARAM([], CL.T_Ptr(CL.T_Named RN.worldTy), "wrld")],
250            | 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)  
251                in                in
252                  topDecls := initFn :: !topDecls                  topDecls := shutdownFn :: initFn :: !topDecls
253                end                end
   
254        (* create and register the initially function for a program *)        (* create and register the initially function for a program *)
255          fun initially {          fun initially {
256                prog = Prog{strands, initially, numDims,...},                prog = Prog{name=progName, strands, initially, ...},
257                isArray : bool,                isArray : bool,
258                iterPrefix : stm list,                iterPrefix : stm list,
259                iters : (var * exp * exp) list,                iters : (var * exp * exp) list,
# Line 201  Line 263 
263              } = let              } = let
264                val name = Atom.toString strand                val name = Atom.toString strand
265                val nDims = List.length iters                val nDims = List.length iters
266                  val worldTy = CL.T_Ptr(CL.T_Named N.worldTy)
267                fun mapi f xs = let                fun mapi f xs = let
268                      fun mapf (_, []) = []                      fun mapf (_, []) = []
269                        | mapf (i, x::xs) = f(i, x) :: mapf(i+1, xs)                        | mapf (i, x::xs) = f(i, x) :: mapf(i+1, xs)
# Line 209  Line 272 
272                      end                      end
273                val baseInit = mapi (fn (i, (_, e, _)) => (i, CL.I_Exp e)) iters                val baseInit = mapi (fn (i, (_, e, _)) => (i, CL.I_Exp e)) iters
274                val sizeInit = mapi                val sizeInit = mapi
275                      (fn (i, (ToCL.V(ty, _), lo, hi)) =>                      (fn (i, (CL.V(ty, _), lo, hi)) =>
276                          (i, CL.I_Exp(CL.mkBinOp(CL.mkBinOp(hi, CL.#-, lo), CL.#+, CL.mkInt(1, ty))))                          (i, CL.I_Exp(CL.mkBinOp(CL.mkBinOp(hi, CL.#-, lo), CL.#+, CL.E_Int(1, ty))))
277                      ) iters                      ) iters
278                    val numStrandsVar = "numStrandsVar"              (* code to allocate the world and initial strands *)
279                val allocCode = iterPrefix @ [                val wrld = "wrld"
280                  val allocCode = [
281                        CL.mkComment["allocate initial block of strands"],                        CL.mkComment["allocate initial block of strands"],
282                        CL.mkDecl(CL.T_Array(CL.int32, SOME nDims), "base", SOME(CL.I_Array baseInit)),                        CL.mkDecl(CL.T_Array(CL.int32, SOME nDims), "base", SOME(CL.I_Array baseInit)),
283                        CL.mkDecl(CL.T_Array(CL.uint32, SOME nDims), "size", SOME(CL.I_Array sizeInit)),                        CL.mkDecl(CL.T_Array(CL.uint32, SOME nDims), "size", SOME(CL.I_Array sizeInit)),
284                        CL.mkDecl(CL.int32,"numDims",SOME(CL.I_Exp(CL.mkInt(IntInf.fromInt nDims, CL.int32))))                        CL.mkDecl(worldTy, wrld,
285                      ]                          SOME(CL.I_Exp(CL.E_Apply(N.allocInitially, [
286                val numStrandsLoopBody = CL.mkExpStm(CL.mkAssignOp(CL.mkVar numStrandsVar, CL.*=,CL.mkSubscript(CL.mkVar "size",CL.mkVar "i")))                              CL.mkVar "ProgramName",
287                val numStrandsLoop =  CL.mkFor([(CL.intTy, "i", CL.mkInt(0,CL.intTy))],                              CL.mkUnOp(CL.%&, CL.E_Var(N.strandDesc name)),
288                      CL.mkBinOp(CL.mkVar "i", CL.#<, CL.mkVar "numDims"),                              CL.E_Bool isArray,
289                      [CL.mkPostOp(CL.mkVar "i", CL.^++)], numStrandsLoopBody)                              CL.E_Int(IntInf.fromInt nDims, CL.int32),
290                in                              CL.E_Var "base",
291                  numDims := nDims;                              CL.E_Var "size"
292                  initially := allocCode @ [numStrandsLoop]                            ]))))
               end  
   
   
       (***** OUTPUT *****)  
         fun genStrandInit(Strand{name,tyName,state,output,code,...},nDims) = let  
               val params = [  
                       CL.PARAM([], CL.T_Ptr(CL.uint32), "sizes" ),  
                       CL.PARAM([], CL.intTy, "width"),  
                       CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "strands")  
                     ]  
               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"])  
293                              ]                              ]
294                                          else let              (* create the loop nest for the initially iterations *)
295                                                  val index = CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")                val indexVar = "ix"
296                                          in                val strandTy = CL.T_Ptr(CL.T_Named(N.strandTy name))
297                                                  CL.mkBlock([CL.mkCall(RN.strandInit name, [CL.mkUnOp(CL.%&,CL.mkSubscript(CL.mkVar "strands",index)),                fun mkLoopNest [] = CL.mkBlock(createPrefix @ [
298                                                  CL.mkVar "x", CL.mkVar"y"])])                        CL.mkDecl(strandTy, "sp",
299                                          end                          SOME(CL.I_Exp(
300                              CL.E_Cast(strandTy,
301                                  | mkLoopNest (param::rest,count,nDims) = let                            CL.E_Apply(N.inState, [CL.E_Var "wrld", CL.E_Var indexVar]))))),
302                                          val body = mkLoopNest (rest, count + 1,nDims)                        CL.mkCall(N.strandInit name, CL.E_Var "sp" :: args),
303                          CL.mkAssign(CL.E_Var indexVar, CL.mkBinOp(CL.E_Var indexVar, CL.#+, CL.E_Int(1, CL.uint32)))
304                        ])
305                    | mkLoopNest ((CL.V(ty, param), lo, hi)::iters) = let
306                        val body = mkLoopNest iters
307                                     in                                     in
308                                          CL.mkFor(                                          CL.mkFor(
309                                                          [(CL.intTy, param, CL.mkInt(0,CL.intTy))],                          [(ty, param, lo)],
310                                                  CL.mkBinOp(CL.mkVar param, CL.#<, CL.mkSubscript(CL.mkVar "sizes",CL.mkInt(count,CL.intTy))),                          CL.mkBinOp(CL.E_Var param, CL.#<=, hi),
311                                                  [CL.mkPostOp(CL.mkVar param, CL.^++)],                          [CL.mkPostOp(CL.E_Var param, CL.^++)],
312                                                  body)                                                  body)
313                                     end                                     end
314                  val iterCode = [
315                          CL.mkComment["initially"],
316                          CL.mkDecl(CL.uint32, indexVar, SOME(CL.I_Exp(CL.E_Int(0, CL.uint32)))),
317                          mkLoopNest iters
318                        ]
319                  val body = CL.mkBlock(
320                        iterPrefix @
321                        allocCode @
322                        iterCode @
323                        [CL.mkReturn(SOME(CL.E_Var "wrld"))])
324                  val initFn = CL.D_Func([], worldTy, N.initially, [], body)
325                          in                          in
326                                  [mkLoopNest ((loopParams nDims),0,nDims)]                  initially := initFn
                         end  
                 in  
                         CL.D_Func(["static"], CL.voidTy, RN.strandInitSetup, params,CL.mkBlock(body))  
327                  end                  end
328          fun genStrandPrint (Strand{name, tyName, state, output, code,...},nDims) = let  
329          (***** OUTPUT *****)
330            fun genStrandPrint (Strand{name, tyName, state, output, code,...}) = let
331              (* the print function *)              (* the print function *)
332                val prFnName = concat[name, "_print"]                val prFnName = concat[name, "_print"]
   
333                val prFn = let                val prFn = let
334                      val params = [                      val params = [
335                            CL.PARAM([], CL.T_Ptr(CL.T_Named "FILE"), "outS"),                            CL.PARAM([], CL.T_Ptr(CL.T_Named "FILE"), "outS"),
                           CL.PARAM([], CL.T_Ptr(CL.uint32), "sizes" ),  
                           CL.PARAM([], CL.intTy, "width"),  
336                            CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "self")                            CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "self")
337                          ]                          ]
   
338                     val SOME(ty, x) = !output                     val SOME(ty, x) = !output
339                     val outState = if nDims = 1 then                      val outState = CL.mkIndirect(CL.mkVar "self", x)
                           CL.mkSelect(CL.mkSubscript(CL.mkVar "self",CL.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)  
   
340                      val prArgs = (case ty                      val prArgs = (case ty
341                             of Ty.IVecTy 1 => [CL.mkStr(!RN.gIntFormat ^ "\n"), outState]                             of Ty.IVecTy 1 => [CL.E_Str(!N.gIntFormat ^ "\n"), outState]
342                              | Ty.IVecTy d => let                              | Ty.IVecTy d => let
343                                  val fmt = CL.mkStr(                                  val fmt = CL.E_Str(
344                                        String.concatWith " " (List.tabulate(d, fn _ => !RN.gIntFormat))                                        String.concatWith " " (List.tabulate(d, fn _ => !N.gIntFormat))
345                                        ^ "\n")                                        ^ "\n")
346                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))                                  val args = List.tabulate (d, fn i => ToC.ivecIndex(outState, d, i))
347                                  in                                  in
348                                    fmt :: args                                    fmt :: args
349                                  end                                  end
350                              | Ty.TensorTy[] => [CL.mkStr "%f\n", outState]                              | Ty.TensorTy[] => [CL.E_Str "%f\n", outState]
351                              | Ty.TensorTy[d] => let                              | Ty.TensorTy[d] => let
352                                  val fmt = CL.mkStr(                                  val fmt = CL.E_Str(
353                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))
354                                        ^ "\n")                                        ^ "\n")
355                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))                                  val args = List.tabulate (d, fn i => ToC.vecIndex(outState, d, i))
356                                  in                                  in
357                                    fmt :: args                                    fmt :: args
358                                  end                                  end
359                              | _ => raise Fail("genStrand: unsupported output type " ^ Ty.toString ty)                              | _ => raise Fail("genStrand: unsupported output type " ^ Ty.toString ty)
360                            (* end case *))                            (* end case *))
   
                           val body = let  
   
                             fun loopParams (3) =  
                                  "x"::"y"::"k"::[]  
                               | loopParams (2) =  
                                  "x"::"y"::[]  
                               | loopParams (1) =  
                                  "x"::[]  
                               | loopParams (_) =  
                                 raise Fail("genStrandPrint: unsupported output type " ^ Ty.toString ty)  
   
                            fun mkLoopNest ([],_) =  
                                                 CL.mkCall("fprintf", CL.mkVar "outS" :: prArgs)  
                                 | mkLoopNest (param::rest,count) = let  
                                         val body = mkLoopNest (rest, count + 1)  
361                                     in                                     in
362                                                  CL.mkFor(                        CL.D_Func(["static"], CL.voidTy, prFnName, params,
363                                                          [(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))  
364                      end                      end
365                in                in
366                                   prFn                                   prFn
367                end                end
368          fun genStrandTyDef (Strand{tyName, state,...}) =  
369            fun genStrandTyDef (targetTy, Strand{tyName, state,...}) =
370              (* the type declaration for the strand's state struct *)              (* the type declaration for the strand's state struct *)
371                CL.D_StructDef(                CL.D_StructDef(
372                        List.rev (List.map (fn ToCL.V(ty, x) => (ty, x)) (!state)),                  List.rev (List.map (fn x => (targetTy x, #var x)) (!state)),
373                        tyName)                        tyName)
374    
375    
376          (* generates the load kernel function *)          (* generates the load kernel function *)
377  (* 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;",  
                                                 "}"])  
378  (* generates the opencl buffers for the image data *)  (* generates the opencl buffers for the image data *)
379          fun getGlobalDataBuffers(globals,count,contextVar,errVar) = let          fun getGlobalDataBuffers(globals,contextVar,errVar) = let
380                  val globalBufferDecl =  CL.mkDecl(CL.clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)                val globalBufferDecl =  CL.mkDecl(clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)
381                  val globalBuffer = CL.mkAssign(CL.mkVar(concat[RN.globalsVarName,"_cl"]), CL.mkApply("clCreateBuffer",                val globalBuffer = CL.mkAssign(CL.mkVar(concat[RN.globalsVarName,"_cl"]),
382                                                                  [CL.mkVar contextVar,                      CL.mkApply("clCreateBuffer", [
383                            CL.mkVar contextVar,
384                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",
385                                                                  CL.mkApply("sizeof",[CL.mkVar RN.globalsTy]),                                                                  CL.mkApply("sizeof",[CL.mkVar RN.globalsTy]),
386                                                                  CL.mkVar RN.globalsVarName,                                                                  CL.mkVar RN.globalsVarName,
387                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))                          CL.mkUnOp(CL.%&,CL.mkVar errVar)
388                          ]))
389    
390          fun genDataBuffers([],_,_,_) = []          fun genDataBuffers([],_,_) = []
391            | genDataBuffers((var,nDims)::globals,count,contextVar,errVar) = let            | genDataBuffers((var,nDims)::globals,contextVar,errVar) = let
392          (* FIXME: use CL constructors to  build expressions (not strings) *)          (* FIXME: use CL constructors to  build expressions (not strings) *)
393                    val size = if nDims = 1 then                val size = if nDims = 1
394                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,                      then CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,
395                                           CL.mkIndirect(CL.mkVar var, "size[0]"))                                           CL.mkIndirect(CL.mkVar var, "size[0]"))
396                                          else if nDims = 2 then                                          else if nDims = 2 then
397                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,
# Line 392  Line 401 
401                                            CL.mkIndirect(CL.mkVar var,concat["size[0]", " * ", var, "->size[1] * ", var, "->size[2]"]))                                            CL.mkIndirect(CL.mkVar var,concat["size[0]", " * ", var, "->size[1] * ", var, "->size[2]"]))
402    
403                   in                   in
404                     CL.mkDecl(CL.clMemoryTy,RN.addBufferSuffix var ,NONE)::                     CL.mkDecl(clMemoryTy, RN.addBufferSuffix var ,NONE)::
405                     CL.mkDecl(CL.clMemoryTy,RN.addBufferSuffixData var ,NONE)::                     CL.mkDecl(clMemoryTy, RN.addBufferSuffixData var ,NONE)::
406                     CL.mkAssign(CL.mkVar(RN.addBufferSuffix var), CL.mkApply("clCreateBuffer",                     CL.mkAssign(CL.mkVar(RN.addBufferSuffix var), CL.mkApply("clCreateBuffer",
407                                                                  [CL.mkVar contextVar,                                                                  [CL.mkVar contextVar,
408                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",
# Line 405  Line 414 
414                                                                   CL.mkVar "CL_MEM_COPY_HOST_PTR",                                                                   CL.mkVar "CL_MEM_COPY_HOST_PTR",
415                                                                  size,                                                                  size,
416                                                                  CL.mkIndirect(CL.mkVar var,"data"),                                                                  CL.mkIndirect(CL.mkVar var,"data"),
417                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)])):: genDataBuffers(globals,count + 2,contextVar,errVar)                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)])):: genDataBuffers(globals,contextVar,errVar)
418                  end                  end
419          in          in
420                  [globalBufferDecl] @ [globalBuffer] @ genDataBuffers(globals,count + 2,contextVar,errVar)                  [globalBufferDecl] @ [globalBuffer] @ genDataBuffers(globals,contextVar,errVar)
421          end          end
422    
423    
# Line 416  Line 425 
425          fun genGlobalArguments(globals,count,kernelVar,errVar) = let          fun genGlobalArguments(globals,count,kernelVar,errVar) = let
426          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",
427                                                                  [CL.mkVar kernelVar,                                                                  [CL.mkVar kernelVar,
428                                                                   CL.mkInt(count,CL.intTy),                                                                   CL.mkPostOp(CL.E_Var count, CL.^++),
429                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
430                                                                   CL.mkUnOp(CL.%&,CL.mkVar(concat[RN.globalsVarName,"_cl"]))])))                                                                   CL.mkUnOp(CL.%&,CL.mkVar(concat[RN.globalsVarName,"_cl"]))])))
431    
# Line 425  Line 434 
434    
435                  CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=, CL.mkApply("clSetKernelArg",                  CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=, CL.mkApply("clSetKernelArg",
436                                                                  [CL.mkVar kernelVar,                                                                  [CL.mkVar kernelVar,
437                                                                   CL.mkInt(count,CL.intTy),                                   CL.mkPostOp(CL.E_Var count, CL.^++),
438                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
439                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffix var))])))::                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffix var))])))::
440    
441                          CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,CL.mkApply("clSetKernelArg",                          CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,CL.mkApply("clSetKernelArg",
442                                                                  [CL.mkVar kernelVar,                                                                  [CL.mkVar kernelVar,
443                                                                   CL.mkInt((count + 1),CL.intTy),                                   CL.mkPostOp(CL.E_Var count, CL.^++),
444                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
445                                                                   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)
446    
447          in          in
448    
449                  [globalArgument] @ genDataArguments(globals,count + 1,kernelVar,errVar)                  [globalArgument] @ genDataArguments(globals,count,kernelVar,errVar)
   
         end  
450    
         (* generates the main function of host code *)  
         fun genHostMain() = 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)  
451                end                end
452    
453        (* generates the host-side setup function *)        (* generates the globals buffers and arguments function *)
454          fun genHostSetupFunc (strand as Strand{name,tyName,...}, filename, nDims, initially, imgGlobals) = let          fun genGlobalBuffersArgs (imgGlobals) = let
455              (* 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"  
456                val errVar = "err"                val errVar = "err"
457                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")])  
458                val params = [                val params = [
459                        CL.PARAM([],CL.T_Named("cl_device_id"), deviceVar)                        CL.PARAM([],CL.T_Named("cl_context"), "context"),
460                      ]                        CL.PARAM([],CL.T_Named("cl_kernel"), "kernel"),
461                val declarations = [                        CL.PARAM([],CL.T_Named("int"), "argStart")
                     CL.mkDecl(CL.clProgramTy, programVar, NONE),  
                     CL.mkDecl(CL.clKernelTy, kernelVar, NONE),  
                     CL.mkDecl(CL.clCmdQueueTy, cmdVar, NONE),  
                     CL.mkDecl(CL.clContextTy, contextVar, NONE),  
                     CL.mkDecl(CL.intTy, errVar, NONE),  
                     CL.mkDecl(CL.intTy, numStrandsVar, SOME(CL.I_Exp(CL.mkInt(1,CL.intTy)))),  
                     CL.mkDecl(CL.intTy, stateSizeVar, NONE),  
                     CL.mkDecl(CL.intTy, "width", NONE),  
                     CL.mkDecl(CL.intTy, imgDataSizeVar, NONE),  
                     (*CL.mkDecl(CL.clDeviceIdTy, deviceVar, NONE), *)  
                     CL.mkDecl(CL.T_Ptr(CL.T_Named tyName), inStateVar,NONE),  
                     CL.mkDecl(CL.clMemoryTy,clInstateVar,NONE),  
                     CL.mkDecl(CL.clMemoryTy,clOutStateVar,NONE),  
                     CL.mkDecl(CL.T_Ptr(CL.T_Named tyName), outStateVar,NONE),  
                     CL.mkDecl(CL.charPtr, clFNVar,SOME(CL.I_Exp(CL.mkStr filename))),  
 (* FIXME:  use Paths.diderotInclude *)  
                     CL.mkDecl(CL.charPtr, headerFNVar,SOME(CL.I_Exp(CL.mkStr "../src/include/Diderot/cl-types.h"))),  
                     CL.mkDecl(CL.T_Array(CL.charPtr,SOME(2)),sourcesVar,NONE),  
                     CL.mkDecl(CL.T_Array(CL.T_Named "size_t",SOME(nDims)),globalVar,NONE),  
                     CL.mkDecl(CL.T_Array(CL.T_Named "size_t",SOME(nDims)),localVar,NONE),  
                     CL.mkDecl(CL.intTy,numDevicesVar,SOME(CL.I_Exp(CL.mkInt(~1,CL.intTy)))),  
                     CL.mkDecl(CL.T_Array(CL.T_Named "cl_platform_id", SOME(1)), platformsVar, NONE),  
                     CL.mkDecl(CL.intTy,"num_platforms",SOME(CL.I_Exp(CL.mkInt(~1,CL.intTy))))  
462                  ]                  ]
463              (* Setup Global Variables *)                val clGlobalBuffers = getGlobalDataBuffers(!imgGlobals, "context", "err")
464                val globalsDecl = CL.mkDecl(                val clGlobalArguments = genGlobalArguments(!imgGlobals, "argStart", "kernel", "err")
                     CL.T_Ptr(CL.T_Named RN.globalsTy),  
                     RN.globalsVarName,  
                     SOME(CL.I_Exp(CL.mkApply("malloc", [CL.mkApply("sizeof",[CL.mkVar RN.globalsTy])]))))  
               val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.mkVar RN.globalsVarName])  
   
                 (* Retrieve the platforms  
                 val platformStm = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetPlatformIDs",  
                                                   [CL.mkInt(10,CL.intTy),  
                                                    CL.mkVar platformsVar,  
                                                    CL.mkUnOp(CL.%&,CL.mkVar numPlatformsVar)])),  
                                                    assertStm]  
   
                 val devicesStm = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetDeviceIDs",  
                                                   [CL.mkSubscript(CL.mkVar platformsVar,CL.mkInt(0,CL.intTy)),  
                                                    CL.mkVar "CL_DEVICE_TYPE_GPU",  
                                                    CL.mkInt(1,CL.intTy),  
                                                    CL.mkUnOp(CL.%&,CL.mkVar deviceVar),  
                                                    CL.mkUnOp(CL.%&,CL.mkVar numDevicesVar)])),  
                                                    assertStm] *)  
   
                 (* Create Context *)  
                 val contextStm = [CL.mkAssign(CL.mkVar contextVar, CL.mkApply("clCreateContext",  
                                                   [CL.mkInt(0,CL.intTy),  
                                                   CL.mkInt(1,CL.intTy),  
                                                   CL.mkUnOp(CL.%&,CL.mkVar deviceVar),  
                                                   CL.mkVar "NULL",  
                                                   CL.mkVar "NULL",  
                                                   CL.mkUnOp(CL.%&,CL.mkVar errVar)])),  
                                                   assertStm]  
   
                 (* Create Command Queue *)  
                 val commandStm = [CL.mkAssign(CL.mkVar cmdVar, CL.mkApply("clCreateCommandQueue",  
                                                   [CL.mkVar contextVar,  
                                                   CL.mkVar deviceVar,  
                                                   CL.mkInt(0,CL.intTy),  
                                                   CL.mkUnOp(CL.%&,CL.mkVar errVar)])),  
                                                   assertStm]  
   
   
                 (*Create Program/Build/Kernel with Source statement *)  
                 val createProgStm = CL.mkAssign(CL.mkVar programVar, CL.mkApply("clCreateProgramWithSource",  
                                                                                                                 [CL.mkVar contextVar,  
                                                                                                                  CL.mkInt(2,CL.intTy),  
                                                                                                                  CL.mkCast(CL.T_Ptr(CL.T_Named("const char *")),CL.mkUnOp(CL.%&,CL.mkVar sourcesVar)),  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))  
   
                 (* FIXME: Remove after testing purposes, Build Log for OpenCL*)  
                 val buildLog = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clBuildProgram",  
                                                                                                                 [CL.mkVar programVar,  
                                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkVar "NULL"])),  
                                           CL.mkDecl(CL.charPtr, "build", NONE),  
                                           CL.mkDecl(CL.T_Named("size_t"),"ret_val_size",NONE),  
                                            CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetProgramBuildInfo",  
                                                                                                                 [CL.mkVar programVar,  
                                                                                                                 CL.mkVar deviceVar,  
                                                                                                                  CL.mkVar "CL_PROGRAM_BUILD_LOG",  
                                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkUnOp(CL.%&,CL.mkVar "ret_val_size")])),  
                                           CL.mkAssign(CL.mkVar "build", CL.mkApply("malloc", [CL.mkVar "ret_val_size"])),  
                                                 CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetProgramBuildInfo",  
                                                                                                                 [CL.mkVar programVar,  
                                                                                                                 CL.mkVar deviceVar,  
                                                                                                                  CL.mkVar "CL_PROGRAM_BUILD_LOG",  
                                                                                                                  CL.mkVar "ret_val_size",  
                                                                                                                  CL.mkVar "build",  
                                                                                                                  CL.mkVar "NULL"])),  
                                                 CL.mkAssign(CL.mkSubscript(CL.mkVar "build",CL.mkVar "ret_val_size"),CL.mkVar ("'\\" ^ "0'")),  
                                                 CL.mkCall("printf",[CL.mkStr ( "Build Log:" ^ "\n" ^ "%s" ^ "\n"), CL.mkVar "build"])]  
   
   
   
   
                 val createKernel = CL.mkAssign(CL.mkVar kernelVar, CL.mkApply("clCreateKernel",  
                                                                                                                 [CL.mkVar programVar,  
                                                                                                                  CL.mkStr RN.kernelFuncName,  
                                                                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))  
   
   
                 val create_build_stms = [createProgStm,assertStm] @ buildLog @ [assertStm,createKernel,assertStm]  
   
   
   
                 (* Create Memory Buffers for Strand States and Globals *)  
                 val strandSize = CL.mkAssign(CL.mkVar stateSizeVar,CL.mkBinOp(CL.mkApply("sizeof",  
                                                                         [CL.mkVar tyName]), CL.#*,CL.mkVar numStrandsVar))  
   
                 val clStrandObjects = [CL.mkAssign(CL.mkVar clInstateVar, CL.mkApply("clCreateBuffer",  
                                                                 [CL.mkVar contextVar,  
                                                                 CL.mkVar "CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR",  
                                                                 CL.mkVar stateSizeVar,  
                                                                 CL.mkVar "NULL",  
                                                                 CL.mkUnOp(CL.%&,CL.mkVar errVar)])),  
                                                          CL.mkAssign(CL.mkVar clOutStateVar, CL.mkApply("clCreateBuffer",  
                                                                 [CL.mkVar contextVar,  
                                                                 CL.mkVar "CL_MEM_READ_WRITE",  
                                                                 CL.mkVar stateSizeVar,  
                                                                 CL.mkVar "NULL",  
                                                                 CL.mkUnOp(CL.%&,CL.mkVar errVar)]))]  
   
   
                 (* Setup up selfOut variable *)  
                 val strandsArrays = [CL.mkAssign(CL.mkVar outStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.mkVar numStrandsVar,  
                                                                         CL.#*, CL.mkApply("sizeof",[CL.mkVar tyName]))])),  
                                                                 CL.mkAssign(CL.mkVar inStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.mkVar numStrandsVar,  
                                                                         CL.#*, CL.mkApply("sizeof",[CL.mkVar tyName]))]))]  
   
   
                 (* Initialize Width Parameter *)  
                 val widthDel = if nDims = 2 then  
                           CL.mkAssign(CL.mkVar "width",CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1, CL.intTy)))  
                    else  
                           CL.mkAssign(CL.mkVar "width",CL.mkInt(0,CL.intTy))  
   
   
                 val strands_init = CL.mkCall(RN.strandInitSetup, [  
                         CL.mkVar "size", CL.mkVar "width", CL.mkVar inStateVar  
                       ])  
   
             val clGlobalBuffers = getGlobalDataBuffers(!imgGlobals,3,contextVar,errVar)  
   
   
                 (* Load the Kernel and Header Files *)  
                 val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(1,CL.intTy)),  
                                                                           CL.mkApply(RN.clLoaderFN, [CL.mkVar clFNVar])),  
            CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(0,CL.intTy)),  
                                                                           CL.mkApply(RN.clLoaderFN, [CL.mkVar headerFNVar]))]  
   
                 (* val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(1,CL.intTy)),  
                                                                           CL.mkApply(RN.clLoaderFN, [CL.mkVar clFNVar]))] *)  
   
   
                 (* Created Enqueue Statements *)  
 (* FIXME: simplify this code by function abstraction *)  
         val enqueueStm = if nDims = 1  
                         then [CL.mkAssign(CL.mkVar errVar,  
                                                           CL.mkApply("clEnqueueNDRangeKernel",  
                                                                                                 [CL.mkVar cmdVar,  
                                                                                                  CL.mkVar kernelVar,  
                                                                                                  CL.mkInt(1,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar globalVar,  
                                                                                                  CL.mkVar localVar,  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]  
                         else if nDims = 2  then  
                          [CL.mkAssign(CL.mkVar errVar,  
                                                         CL.mkApply("clEnqueueNDRangeKernel",  
                                                                                                 [CL.mkVar cmdVar,  
                                                                                                  CL.mkVar kernelVar,  
                                                                                                  CL.mkInt(2,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar globalVar,  
                                                                                                  CL.mkVar localVar,  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]  
                         else  
                           [CL.mkAssign(CL.mkVar errVar,  
                                                         CL.mkApply("clEnqueueNDRangeKernel",  
                                                                                                 [CL.mkVar cmdVar,  
                                                                                                  CL.mkVar kernelVar,  
                                                                                                  CL.mkInt(3,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar globalVar,  
                                                                                                  CL.mkVar localVar,  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]  
   
   
   
                 (* Setup Global and Local variables *)  
   
                 val globalAndlocalStms = if nDims = 1 then  
                         [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),  
                          CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),  
                                                                   CL.mkVar "16")]  
   
   
                 else if nDims = 2 then  
                         [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(1,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),  
                                                                   CL.mkVar "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(1,CL.intTy)),  
                                                                   CL.mkVar "16")]  
   
                 else  
                         [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(1,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(2,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(2,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),  
                                                                   CL.mkVar "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(1,CL.intTy)),  
                                                                   CL.mkVar "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(2,CL.intTy)),  
                                                                   CL.mkVar "16")]  
   
   
   
                 (* Setup Kernel arguments *)  
                 val kernelArguments = [CL.mkAssign(CL.mkVar errVar,CL.mkApply("clSetKernelArg",  
                                                                 [CL.mkVar kernelVar,  
                                                                  CL.mkInt(0,CL.intTy),  
                                                                  CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),  
                                                                  CL.mkUnOp(CL.%&,CL.mkVar clInstateVar)])),  
                                                             CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar, CL.|=,CL.mkApply("clSetKernelArg",  
                                                                 [CL.mkVar kernelVar,  
                                                                  CL.mkInt(1,CL.intTy),  
                                                                  CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),  
                                                                  CL.mkUnOp(CL.%&,CL.mkVar clOutStateVar)]))),  
                                                                   CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar, CL.|=,CL.mkApply("clSetKernelArg",  
                                                                 [CL.mkVar kernelVar,  
                                                                  CL.mkInt(2,CL.intTy),  
                                                                  CL.mkApply("sizeof",[CL.mkVar "int"]),  
                                                                  CL.mkUnOp(CL.%&,CL.mkVar "width")])))]  
   
            val clGlobalArguments = genGlobalArguments(!imgGlobals,3,kernelVar,errVar) @ [assertStm]  
   
                 (* Retrieve output *)  
                 val outputStm = CL.mkAssign(CL.mkVar errVar,  
                                                         CL.mkApply("clEnqueueReadBuffer",  
                                                                                                 [CL.mkVar cmdVar,  
                                                                                                  CL.mkVar clOutStateVar,  
                                                                                                  CL.mkVar "CL_TRUE",  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar stateSizeVar,  
                                                                                                  CL.mkVar outStateVar,  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar "NULL"]))  
   
                 (* Free all the objects *)  
                 val freeStms = [CL.mkCall("clReleaseKernel",[CL.mkVar kernelVar]),  
                                                 CL.mkCall("clReleaseProgram",[CL.mkVar programVar ]),  
                                                 CL.mkCall("clReleaseCommandQueue",[CL.mkVar cmdVar]),  
                                                 CL.mkCall("clReleaseContext",[CL.mkVar contextVar]),  
                                                 CL.mkCall("clReleaseMemObject",[CL.mkVar clInstateVar]),  
                                                 CL.mkCall("clReleaseMemObject",[CL.mkVar clOutStateVar])]  
   
   
                 (*Setup Strand Print Function *)  
                 val outputData = [CL.mkDecl(CL.T_Ptr(CL.T_Named("FILE")), "outS", SOME(CL.I_Exp(CL.mkApply("fopen",  
                                                 [CL.mkStr "mip.txt",  
                                                 CL.mkStr "w"])))),  
                                                 CL.mkCall(concat[name, "_print"],  
                                                                         [CL.mkVar "outS",  
                                                                          CL.mkVar "size",  
                                                                          CL.mkVar "width",  
                                                                          CL.mkVar outStateVar])]  
   
   
   
465                  (* Body put all the statments together *)                  (* Body put all the statments together *)
466                  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 *)  
   
467                  in                  in
468                    CL.D_Func([],CL.voidTy,RN.globalsSetupName,params,CL.mkBlock(body))
                 CL.D_Func([],CL.voidTy,RN.setupFName,params,CL.mkBlock(body))  
   
469                  end                  end
470    
471  (* generate the data and global parameters *)  (* generate the data and global parameters *)
472          fun genKeneralGlobalParams ((name,tyname)::rest) =          fun genKeneralGlobalParams ((name,tyname)::rest) =
473                  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]) ::
474                  CL.PARAM([], CL.T_Ptr(CL.T_Named (RN.imageTy tyname)),RN.addBufferSuffix name) ::                  CL.PARAM([], CL.T_Ptr(CL.T_Named (RN.imageTy tyname)),RN.addBufferSuffix name) ::
475                  CL.PARAM([], CL.T_Ptr(CL.voidTy),RN.addBufferSuffixData name) ::                  CL.PARAM([], CL.T_Ptr(CL.voidTy),RN.addBufferSuffixData name) ::
476                  genKeneralGlobalParams(rest)                genKeneralGlobalParams rest
477              | genKeneralGlobalParams [] = []
           | genKeneralGlobalParams ([]) = []  
478    
479          (*generate code for intilizing kernel global data *)          (*generate code for intilizing kernel global data *)
480          fun initKernelGlobals (globals,imgGlobals) = let          fun initKernelGlobals (globals,imgGlobals) = let
481                  fun initGlobalStruct (CL.D_Var(_, _ , name, _)::rest) =  (* FIXME: should use List.map here *)
482                                  CL.mkAssign(CL.mkVar name, CL.mkIndirect(CL.mkVar RN.globalsVarName, name)) ::                fun initGlobalStruct ({hostTy, gpuTy, var}::rest) =
483                                  initGlobalStruct(rest)                      CL.mkAssign(CL.mkVar var, CL.mkIndirect(CL.mkVar RN.globalsVarName, var)) ::
484                    | initGlobalStruct ( _::rest) = initGlobalStruct(rest)                      initGlobalStruct rest
485                    | initGlobalStruct([]) = []                  | initGlobalStruct [] = []
   
486                  fun initGlobalImages((name,tyname)::rest) =                  fun initGlobalImages((name,tyname)::rest) =
487                                  CL.mkAssign(CL.mkVar name, CL.mkVar (RN.addBufferSuffix name)) ::                                  CL.mkAssign(CL.mkVar name, CL.mkVar (RN.addBufferSuffix name)) ::
488                                  CL.mkAssign(CL.mkIndirect(CL.mkVar name,"data"),CL.mkVar (RN.addBufferSuffixData name)) ::                                  CL.mkAssign(CL.mkIndirect(CL.mkVar name,"data"),CL.mkVar (RN.addBufferSuffixData name)) ::
489                                  initGlobalImages(rest)                      initGlobalImages rest
490                    | initGlobalImages([]) = []                    | initGlobalImages [] = []
491                  in                  in
492                    initGlobalStruct(globals) @ initGlobalImages(imgGlobals)                  initGlobalStruct globals @ initGlobalImages(imgGlobals)
493                  end                  end
494    
495          (* generate the main kernel function for the .cl file *)          (* generate the main kernel function for the .cl file *)
496          fun genKernelFun(Strand{name, tyName, state, output, code,...},nDims,globals,imgGlobals) = let          fun genKernelFun (strand, nDims, globals, imgGlobals) = let
497                  val Strand{name, tyName, state, output, code,...} = strand
498                   val fName = RN.kernelFuncName;                   val fName = RN.kernelFuncName;
499                   val inState = "strand_in"                   val inState = "strand_in"
500                   val outState = "strand_out"                   val outState = "strand_out"
# Line 828  Line 504 
504                        CL.PARAM(["__global"], CL.intTy, "width")                        CL.PARAM(["__global"], CL.intTy, "width")
505                      ] @ genKeneralGlobalParams(!imgGlobals)                      ] @ genKeneralGlobalParams(!imgGlobals)
506                    val thread_ids = if nDims = 1                    val thread_ids = if nDims = 1
507                          then [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),                      then [
508                                    CL.mkAssign(CL.mkVar "x",CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(0,CL.intTy)]))]                          CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),
509                          else                          CL.mkAssign(CL.mkVar "x",CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(0,CL.intTy)]))
510                                  [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),                        ]
511                        else [
512                            CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),
513                                   CL.mkDecl(CL.intTy, "y", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),                                   CL.mkDecl(CL.intTy, "y", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),
514                                    CL.mkAssign(CL.mkVar "x",  CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(0,CL.intTy)])),                                    CL.mkAssign(CL.mkVar "x",  CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(0,CL.intTy)])),
515                                    CL.mkAssign(CL.mkVar "y",CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(1,CL.intTy)]))]                          CL.mkAssign(CL.mkVar "y",CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(1,CL.intTy)]))
516                          ]
517                    val strandDecl = [CL.mkDecl(CL.T_Named tyName, inState, NONE),                val strandDecl = [
518                        CL.mkDecl(CL.T_Named tyName, inState, NONE),
519                                                          CL.mkDecl(CL.T_Named tyName, outState,NONE)]                                                          CL.mkDecl(CL.T_Named tyName, outState,NONE)]
520                    val strandObjects  = if nDims = 1                    val strandObjects  = if nDims = 1
521                          then [CL.mkAssign(CL.mkSubscript(CL.mkVar "selfIn",CL.mkStr "x"),                          then [
522                                                                           CL.mkVar inState),                              CL.mkAssign( CL.mkVar inState, CL.mkSubscript(CL.mkVar "selfIn", CL.mkStr "x")),
523                                    CL.mkAssign(CL.mkSubscript(CL.mkVar "selfOut",CL.mkStr "x"),                              CL.mkAssign(CL.mkVar outState,CL.mkSubscript(CL.mkVar "selfOut", CL.mkStr "x"))
524                                                                           CL.mkVar outState)]                            ]
525                          else let                          else let
526                                  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")
527                                  in                                  in
528                                          [CL.mkAssign(CL.mkSubscript(CL.mkVar "selfIn",index),                                          [CL.mkAssign(CL.mkVar inState, CL.mkSubscript(CL.mkVar "selfIn",index)),
529                                                                          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)]  
530                                  end                                  end
   
   
531                    val status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy))))                    val status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy))))
532                    val local_vars =  thread_ids @ initKernelGlobals(!globals,!imgGlobals)  @ strandDecl @ strandObjects @ [status]                    val local_vars =  thread_ids @ initKernelGlobals(!globals,!imgGlobals)  @ strandDecl @ strandObjects @ [status]
533                    val while_exp = CL.mkBinOp(CL.mkBinOp(CL.mkVar "status",CL.#!=, CL.mkVar RN.kStabilize),CL.#||,CL.mkBinOp(CL.mkVar "status", CL.#!=, CL.mkVar RN.kDie))                    val while_exp = CL.mkBinOp(CL.mkBinOp(CL.mkVar "status",CL.#!=, CL.mkVar RN.kStabilize),CL.#||,CL.mkBinOp(CL.mkVar "status", CL.#!=, CL.mkVar RN.kDie))
534                    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)])),                val whileBody = CL.mkBlock [
535                                                          CL.mkCall(RN.strandStabilize name,[ CL.mkUnOp(CL.%&,CL.mkVar inState),  CL.mkUnOp(CL.%&,CL.mkVar outState)])]                        CL.mkAssign(CL.mkVar "status",
536                            CL.mkApply(RN.strandUpdate name,
537                    val whileBlock = [CL.mkWhile(while_exp,CL.mkBlock while_body)]                            [CL.mkUnOp(CL.%&,CL.mkVar inState), CL.mkUnOp(CL.%&,CL.mkVar outState)])),
538                          CL.mkCall(RN.strandStabilize name,
539                            [CL.mkUnOp(CL.%&,CL.mkVar inState), CL.mkUnOp(CL.%&,CL.mkVar outState)])
540                        ]
541                  val whileBlock = [CL.mkWhile(while_exp, whileBody)]
542                    val body = CL.mkBlock(local_vars  @ whileBlock)                    val body = CL.mkBlock(local_vars  @ whileBlock)
543                  in                  in
544                     CL.D_Func(["__kernel"], CL.voidTy, fName, params, body)                     CL.D_Func(["__kernel"], CL.voidTy, fName, params, body)
545                  end                  end
546          (* generate a global structure from the globals *)          (* generate a global structure from the globals *)
547          fun genGlobalStruct(globals) = let          fun genGlobalStruct (targetTy, globals) = let
548                   fun getGlobals(CL.D_Var(_,ty,globalVar,_)::rest) = (ty,globalVar)::getGlobals(rest)                val globs = List.map (fn (x : mirror_var) => (targetTy x, #var x)) globals
549                     | getGlobals([]) = []                in
550                     | getGlobals(_::rest) = getGlobals(rest)                  CL.D_StructDef(globs, RN.globalsTy)
551                   in                end
552                          CL.D_StructDef(getGlobals(globals),RN.globalsTy)          fun genGlobals (declFn, targetTy, globals) = let
553                    end                fun doVar (x : mirror_var) = declFn (CL.D_Var([], targetTy x, #var x, NONE))
554                  in
555        (* generate the table of strand descriptors *)                  List.app doVar globals
         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)))))  
556                end                end
   
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 897  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 = PrintAsC.output(clppStrm, dcl)                fun clppDecl dcl = PrintAsCL.output(clppStrm, dcl)
568                val strands = AtomTable.listItems strands                val strands = AtomTable.listItems strands
569                val [strand as Strand{name, tyName, code,init_code, ...}] = strands                val [strand as Strand{name, tyName, code,init_code, ...}] = strands
570                in                in
# Line 910  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, #gpuTy, !globals);
580                  clppDecl (genGlobalStruct (!globals));                  clppDecl (genGlobalStruct (#gpuTy, !globals));
581                  clppDecl (genStrandTyDef strand);                  clppDecl (genStrandTyDef(#gpuTy, strand));
582                  List.app clppDecl (!code);                  List.app clppDecl (!code);
583                  clppDecl (genKernelFun (strand,!numDims,globals,imgGlobals));                  clppDecl (genKernelFun (strand,!numDims,globals,imgGlobals));
584                (* Generate the Host file .c *)  
585                  (* Generate the Host C file *)
586                  cppDecl (CL.D_Verbatim([                  cppDecl (CL.D_Verbatim([
587                      if double                      if double
588                        then "#define DIDEROT_DOUBLE_PRECISION"                        then "#define DIDEROT_DOUBLE_PRECISION"
# Line 925  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, #hostTy, !globals);
594                  cppDecl (genGlobalStruct (!globals));                  cppDecl (genGlobalStruct (#hostTy, !globals));
595                  cppDecl (genStrandTyDef strand);                  cppDecl (genStrandTyDef (#gpuTy, 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 (genGlobalBuffersArgs (imgGlobals));
600                    cppDecl (!initially);
601                  PrintAsC.close cppStrm;                  PrintAsC.close cppStrm;
602                  PrintAsC.close clppStrm;                  PrintAsCL.close clppStrm;
603                  TextIO.closeOut cOutS;                  TextIO.closeOut cOutS;
604                  TextIO.closeOut clOutS                  TextIO.closeOut clOutS
605                end                end

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

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