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

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

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