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 1351, Mon Jun 20 14:48:00 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      (* helper functions for specifying parameters in various address spaces *)
19        fun clParam (spc, ty, x) = CL.PARAM([spc], ty, x)
20        fun globalParam (ty, x) = CL.PARAM(["__global"], ty, x)
21        fun constantParam (ty, x) = CL.PARAM(["__constant"], ty, x)
22        fun localParam (ty, x) = CL.PARAM(["__local"], ty, x)
23        fun privateParam (ty, x) = CL.PARAM(["__private"], ty, x)
24    
25    (* C variable translation *)    (* C variable translation *)
26      structure TrCVar =      structure TrCVar =
# Line 25  Line 33 
33        (* 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) *)
34          fun lvalueVar (env, x) = (case V.kind x          fun lvalueVar (env, x) = (case V.kind x
35                 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))
36                  | IL.VK_State strand => raise Fail "unexpected strand context"                  | IL.VK_State strand => CL.mkIndirect(CL.mkVar "selfOut", lookup(env, x))
37                  | IL.VK_Local => CL.mkVar(lookup(env, x))                  | IL.VK_Local => CL.mkVar(lookup(env, x))
38                (* end case *))                (* end case *))
39        (* translate a variable that occurs in an r-value context *)        (* translate a variable that occurs in an r-value context *)
40          val rvalueVar = lvalueVar          fun rvalueVar (env, x) = (case V.kind x
41                   of IL.VK_Global => CL.mkIndirect(CL.mkVar RN.globalsVarName, lookup(env, x))
42                    | IL.VK_State strand => CL.mkIndirect(CL.mkVar "selfIn", lookup(env, x))
43                    | IL.VK_Local => CL.mkVar(lookup(env, x))
44                  (* end case *))
45        end        end
46    
47      structure ToC = TreeToCFn (TrCVar)      structure ToC = TreeToCFn (TrCVar)
# Line 39  Line 51 
51      type stm = CL.stm      type stm = CL.stm
52    
53    (* OpenCL specific types *)    (* OpenCL specific types *)
54        val clIntTy = CL.T_Named "cl_int"
55      val clProgramTy = CL.T_Named "cl_program"      val clProgramTy = CL.T_Named "cl_program"
56      val clKernelTy  = CL.T_Named "cl_kernel"      val clKernelTy  = CL.T_Named "cl_kernel"
57      val clCmdQueueTy = CL.T_Named "cl_command_queue"      val clCmdQueueTy = CL.T_Named "cl_command_queue"
# Line 46  Line 59 
59      val clDeviceIdTy = CL.T_Named "cl_device_id"      val clDeviceIdTy = CL.T_Named "cl_device_id"
60      val clPlatformIdTy = CL.T_Named "cl_platform_id"      val clPlatformIdTy = CL.T_Named "cl_platform_id"
61      val clMemoryTy = CL.T_Named "cl_mem"      val clMemoryTy = CL.T_Named "cl_mem"
62        val globPtrTy = CL.T_Ptr(CL.T_Named RN.globalsTy)
63    
64      (* variable or field that is mirrored between host and GPU *)
65        type mirror_var = {
66                hostTy : CL.ty,             (* variable type on Host (i.e., C type) *)
67                gpuTy : CL.ty,              (* variable's type on GPU (i.e., OpenCL type) *)
68                var : CL.var                (* variable name *)
69              }
70    
71      datatype strand = Strand of {      datatype strand = Strand of {
72          name : string,          name : string,
73          tyName : string,          tyName : string,
74          state : var list ref,          state : mirror_var list ref,
75          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) *)
76          code : CL.decl list ref,          code : CL.decl list ref,
77          init_code: CL.decl ref          init_code: CL.decl ref
# Line 61  Line 82 
82          double : bool,                  (* true for double-precision support *)          double : bool,                  (* true for double-precision support *)
83          parallel : bool,                (* true for multithreaded (or multi-GPU) target *)          parallel : bool,                (* true for multithreaded (or multi-GPU) target *)
84          debug : bool,                   (* true for debug support in executable *)          debug : bool,                   (* true for debug support in executable *)
85          globals : CL.decl list ref,          globals : mirror_var list ref,
86          topDecls : CL.decl list ref,          topDecls : CL.decl list ref,
87          strands : strand AtomTable.hash_table,          strands : strand AtomTable.hash_table,
88          initially : CL.stm list ref,          initially :  CL.decl ref,
89          numDims: int ref,          numDims: int ref,               (* number of dimensions in initially iteration *)
90          imgGlobals: (string * int) list ref,          imgGlobals: (string * int) list ref,
91          prFn: CL.decl ref          prFn: CL.decl ref
92        }        }
# Line 98  Line 119 
119    (* TreeIL to target translations *)    (* TreeIL to target translations *)
120      structure Tr =      structure Tr =
121        struct        struct
       (* this function is used for the initially clause, so it generates OpenCL *)  
122          fun fragment (ENV{info, vMap, scope}, blk) = let          fun fragment (ENV{info, vMap, scope}, blk) = let
123                val (vMap, stms) = ToCL.trFragment (vMap, blk)                val (vMap, stms) = (case scope
124                         of GlobalScope => ToC.trFragment (vMap, blk)
125    (* NOTE: if we move strand initialization to the GPU, then we'll have to change the following code! *)
126                          | InitiallyScope => ToC.trFragment (vMap, blk)
127                          | _ => ToCL.trFragment (vMap, blk)
128                        (* end case *))
129                in                in
130                  (ENV{info=info, vMap=vMap, scope=scope}, stms)                  (ENV{info=info, vMap=vMap, scope=scope}, stms)
131                end                end
132          fun saveState cxt stateVars (env, args, stm) = (          fun block (ENV{vMap, scope, ...}, blk) = let
133                  fun saveState cxt stateVars trAssign (env, args, stm) = (
134                ListPair.foldrEq                ListPair.foldrEq
135                  (fn (x, e, stms) => ToCL.trAssign(env, x, e)@stms)                        (fn (x, e, stms) => trAssign(env, x, e)@stms)
136                    [stm]                    [stm]
137                      (stateVars, args)                      (stateVars, args)
138                ) handle ListPair.UnequalLengths => (                ) handle ListPair.UnequalLengths => (
139                  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"]);
140                  raise Fail(concat["saveState ", cxt, ": length mismatch"]))                  raise Fail(concat["saveState ", cxt, ": length mismatch"]))
141          fun block (ENV{vMap, scope, ...}, blk) = (case scope                in
142                 of StrandScope stateVars => ToCL.trBlock (vMap, saveState "StrandScope" stateVars, blk)                  case scope
143                  | 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! *)
144                     of StrandScope stateVars =>
145                          ToCL.trBlock (vMap, saveState "StrandScope" stateVars ToC.trAssign, blk)
146                      | MethodScope stateVars =>
147                          ToCL.trBlock (vMap, saveState "MethodScope" stateVars ToCL.trAssign, blk)
148                  | InitiallyScope => ToCL.trBlock (vMap, fn (_, _, stm) => [stm], blk)                  | InitiallyScope => ToCL.trBlock (vMap, fn (_, _, stm) => [stm], blk)
149                  | _ => ToC.trBlock (vMap, fn (_, _, stm) => [stm], blk)                  | _ => ToC.trBlock (vMap, fn (_, _, stm) => [stm], blk)
150                (* end case *))                  (* end case *)
151                  end
152          fun exp (ENV{vMap, ...}, e) = ToCL.trExp(vMap, e)          fun exp (ENV{vMap, ...}, e) = ToCL.trExp(vMap, e)
153        end        end
154    
# Line 126  Line 157 
157        struct        struct
158          fun name (ToCL.V(_, name)) = name          fun name (ToCL.V(_, name)) = name
159          fun global (Prog{globals, imgGlobals, ...}, name, ty) = let          fun global (Prog{globals, imgGlobals, ...}, name, ty) = let
160                val ty' = ToCL.trType ty                val x = {hostTy = ToC.trType ty, gpuTy = ToCL.trType ty, var = name}
161                fun isImgGlobal (imgGlobals, Ty.ImageTy(ImageInfo.ImgInfo{dim, ...}), name) =  imgGlobals  := (name,dim):: !imgGlobals                fun isImgGlobal (Ty.ImageTy(ImageInfo.ImgInfo{dim, ...}), name) =
162                  | isImgGlobal (imgGlobals, _, _) =  ()                      imgGlobals  := (name,dim) :: !imgGlobals
163                in                  | isImgGlobal _ =  ()
164                  globals := CL.D_Var([], ty', name, NONE) :: !globals;                in
165                  isImgGlobal(imgGlobals,ty,name);                  globals := x :: !globals;
166                  ToCL.V(ty', name)                  isImgGlobal (ty, name);
167                    ToCL.V(#gpuTy x, name)
168                end                end
169          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)
170          fun state (Strand{state, ...}, x) = let          fun state (Strand{state, ...}, x) = let
171                val ty' = ToCL.trType(V.ty x)                val ty = V.ty x
172                val x' = ToCL.V(ty', V.name x)                val x' = {hostTy = ToC.trType ty, gpuTy = ToCL.trType ty, var = V.name x}
173                in                in
174                  state := x' :: !state;                  state := x' :: !state;
175                  x'                  ToCL.V(#gpuTy x', #var x')
176                end                end
177        end        end
178    
# Line 179  Line 211 
211                    globals = ref [],                    globals = ref [],
212                    topDecls = ref [],                    topDecls = ref [],
213                    strands = AtomTable.mkTable (16, Fail "strand table"),                    strands = AtomTable.mkTable (16, Fail "strand table"),
214                    initially = ref([CL.S_Comment["missing initially"]]),                    initially = ref(CL.D_Comment["missing initially"]),
215                                    numDims = ref(0),                    numDims = ref 0,
216                                    imgGlobals = ref[],                                    imgGlobals = ref[],
217                                    prFn = ref(CL.D_Comment(["No Print Function"]))                                    prFn = ref(CL.D_Comment(["No Print Function"]))
218                  })                  })
219        (* register the global initialization part of a program *)        (* register the global initialization part of a program *)
220    (* FIXME: unused code; can this be removed??
221            fun globalIndirects (globals,stms) = let            fun globalIndirects (globals,stms) = let
222                  fun getGlobals (CL.D_Var(_,_,globalVar,_)::rest) =                  fun getGlobals ({name,target as TargetUtil.TARGET_CL}::rest) =
223                        CL.mkAssign(CL.mkIndirect(CL.mkVar RN.globalsVarName,globalVar),CL.mkVar globalVar)                        CL.mkAssign(CL.mkIndirect(CL.mkVar RN.globalsVarName,name),CL.mkVar name)
224                          ::getGlobals rest                          ::getGlobals rest
225                    | getGlobals [] = []                    | getGlobals [] = []
226                    | getGlobals (_::rest) = getGlobals rest                    | 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 *)
232          fun inputs (Prog{topDecls, ...}, stm) = let          fun inputs (Prog{topDecls, ...}, stm) = let
233                val inputsFn = CL.D_Func(                val inputsFn = CL.D_Func(
# Line 207  Line 240 
240    
241        (* register the global initialization part of a program *)        (* register the global initialization part of a program *)
242          fun init (Prog{topDecls, ...}, init) = let          fun init (Prog{topDecls, ...}, init) = let
243                val globPtrTy = CL.T_Ptr(CL.T_Named RN.globalsTy)                val globalsDecl = CL.mkAssign(CL.E_Var RN.globalsVarName,
244                        CL.mkApply("malloc", [CL.mkApply("sizeof",[CL.mkVar RN.globalsTy])]))
245                  val initGlobalsCall = CL.mkCall(RN.initGlobalsHelper,[])
246                val initFn = CL.D_Func(                val initFn = CL.D_Func(
247                      [], CL.voidTy, RN.initGlobals, [CL.PARAM([], globPtrTy, RN.globalsVarName)],                      [], CL.voidTy, RN.initGlobals, [],
248                        CL.mkBlock([globalsDecl,initGlobalsCall]))
249                  val initFn_helper = CL.D_Func(
250                        [], CL.voidTy, RN.initGlobalsHelper, [],
251                      init)                      init)
252                val shutdownFn = CL.D_Func(                val shutdownFn = CL.D_Func(
253                      [], CL.voidTy, RN.shutdown,                      [], CL.voidTy, RN.shutdown,
254                      [CL.PARAM([], CL.T_Ptr(CL.T_Named RN.worldTy), "wrld")],                      [CL.PARAM([], CL.T_Ptr(CL.T_Named RN.worldTy), "wrld")],
255                      CL.S_Block[])                      CL.S_Block[])
256                in                in
257                  topDecls := shutdownFn :: initFn :: !topDecls                  topDecls := shutdownFn :: initFn :: initFn_helper :: !topDecls
258                end                end
   
259        (* create and register the initially function for a program *)        (* create and register the initially function for a program *)
260          fun initially {          fun initially {
261                prog = Prog{strands, initially, numDims,...},                prog = Prog{name=progName, strands, initially, numDims, ...},
262                isArray : bool,                isArray : bool,
263                iterPrefix : stm list,                iterPrefix : stm list,
264                iters : (var * exp * exp) list,                iters : (var * exp * exp) list,
# Line 231  Line 268 
268              } = let              } = let
269                val name = Atom.toString strand                val name = Atom.toString strand
270                val nDims = List.length iters                val nDims = List.length iters
271                  val worldTy = CL.T_Ptr(CL.T_Named N.worldTy)
272                fun mapi f xs = let                fun mapi f xs = let
273                      fun mapf (_, []) = []                      fun mapf (_, []) = []
274                        | 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 277 
277                      end                      end
278                val baseInit = mapi (fn (i, (_, e, _)) => (i, CL.I_Exp e)) iters                val baseInit = mapi (fn (i, (_, e, _)) => (i, CL.I_Exp e)) iters
279                val sizeInit = mapi                val sizeInit = mapi
280                      (fn (i, (ToCL.V(ty, _), lo, hi)) =>                      (fn (i, (CL.V(ty, _), lo, hi)) =>
281                          (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))))
282                      ) iters                      ) iters
283                    val numStrandsVar = "numStrandsVar"              (* code to allocate the world and initial strands *)
284                val allocCode = iterPrefix @ [                val wrld = "wrld"
285                  val allocCode = [
286                        CL.mkComment["allocate initial block of strands"],                        CL.mkComment["allocate initial block of strands"],
287                        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)),
288                        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)),
289                        CL.mkDecl(CL.int32,"numDims",SOME(CL.I_Exp(CL.mkInt(IntInf.fromInt nDims, CL.int32))))                        CL.mkDecl(worldTy, wrld,
290                            SOME(CL.I_Exp(CL.E_Apply(RN.allocInitially, [
291                                CL.mkVar "ProgramName",
292                                CL.mkUnOp(CL.%&, CL.E_Var(N.strandDesc name)),
293                                CL.E_Bool isArray,
294                                CL.E_Int(IntInf.fromInt nDims, CL.int32),
295                                CL.E_Var "base",
296                                CL.E_Var "size"
297                              ]))))
298                      ]                      ]
299                val numStrandsLoopBody =              (* create the loop nest for the initially iterations
300                      CL.mkExpStm(CL.mkAssignOp(CL.mkVar numStrandsVar, CL.*=,CL.mkSubscript(CL.mkVar "size",CL.mkVar "i")))                val indexVar = "ix"
301                val numStrandsLoop =  CL.mkFor([(CL.intTy, "i", CL.mkInt(0,CL.intTy))],                val strandTy = CL.T_Ptr(CL.T_Named(N.strandTy name))
302                      CL.mkBinOp(CL.mkVar "i", CL.#<, CL.mkVar "numDims"),                fun mkLoopNest [] = CL.mkBlock(createPrefix @ [
303                      [CL.mkPostOp(CL.mkVar "i", CL.^++)], numStrandsLoopBody)                        CL.mkDecl(strandTy, "sp",
304                in                          SOME(CL.I_Exp(
305                  numDims := nDims;                            CL.E_Cast(strandTy,
306                  initially := allocCode @ [numStrandsLoop]                            CL.E_Apply(N.inState, [CL.E_Var "wrld", CL.E_Var indexVar]))))),
307                end                        CL.mkCall(N.strandInit name,
308                            CL.E_Var RN.globalsVarName :: CL.E_Var "sp" :: args),
309                          CL.mkAssign(CL.E_Var indexVar, CL.mkBinOp(CL.E_Var indexVar, CL.#+, CL.E_Int(1, CL.uint32)))
310        (***** OUTPUT *****)                      ])
311          fun genStrandInit (Strand{name,tyName,state,output,code,...}, nDims) = let                  | mkLoopNest ((CL.V(ty, param), lo, hi)::iters) = let
312                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)  
313                            in                            in
314                              CL.mkFor(                              CL.mkFor(
315                                  [(CL.intTy, param, CL.mkInt(0,CL.intTy))],                          [(ty, param, lo)],
316                                  CL.mkBinOp(CL.mkVar param, CL.#<, CL.mkSubscript(CL.mkVar "sizes",CL.mkInt(count,CL.intTy))),                          CL.mkBinOp(CL.E_Var param, CL.#<=, hi),
317                                  [CL.mkPostOp(CL.mkVar param, CL.^++)],                          [CL.mkPostOp(CL.E_Var param, CL.^++)],
318                                  body)                                  body)
319                            end                            end
320                  val iterCode = [
321                          CL.mkComment["initially"],
322                          CL.mkDecl(CL.uint32, indexVar, SOME(CL.I_Exp(CL.E_Int(0, CL.uint32)))),
323                          mkLoopNest iters
324                        ] *)
325                  val body = CL.mkBlock(
326                        iterPrefix @
327                        allocCode @
328                        [CL.mkReturn(SOME(CL.E_Var "wrld"))])
329                  val initFn = CL.D_Func([], worldTy, N.initially, [], body)
330                      in                      in
331                        [mkLoopNest ((loopParams nDims),0,nDims)]                  numDims := nDims;
332                      end                  initially := initFn
                 in  
                   CL.D_Func(["static"], CL.voidTy, RN.strandInitSetup, params,CL.mkBlock(body))  
333                  end                  end
334    
335          fun genStrandPrint (Strand{name, tyName, state, output, code,...},nDims) = let        (***** OUTPUT *****)
336            fun genStrandPrint (Strand{name, tyName, state, output, code,...}) = let
337              (* the print function *)              (* the print function *)
338                val prFnName = concat[name, "_print"]                val prFnName = concat[name, "_print"]
339                val prFn = let                val prFn = let
340                      val params = [                      val params = [
341                            CL.PARAM([], CL.T_Ptr(CL.T_Named "FILE"), "outS"),                            CL.PARAM([], CL.T_Ptr(CL.T_Named "FILE"), "outS"),
342                            CL.PARAM([], CL.T_Ptr(CL.uint32), "sizes" ),                              CL.PARAM([], CL.T_Ptr(CL.T_Num(RawTypes.RT_UInt8)),"status"),
343                            CL.PARAM([], CL.intTy, "width"),                              CL.PARAM([], CL.intTy,"numStrands"),
344                            CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "self")                            CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "self")
345                          ]                          ]
   
346                     val SOME(ty, x) = !output                     val SOME(ty, x) = !output
347                     val outState = if nDims = 1 then                      val outState = CL.mkSelect(CL.mkSubscript(CL.mkVar "self",CL.E_Var "i"), 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)  
   
348                      val prArgs = (case ty                      val prArgs = (case ty
349                             of Ty.IVecTy 1 => [CL.mkStr(!RN.gIntFormat ^ "\n"), outState]                             of Ty.IVecTy 1 => [CL.E_Str(!N.gIntFormat ^ "\n"), outState]
350                              | Ty.IVecTy d => let                              | Ty.IVecTy d => let
351                                  val fmt = CL.mkStr(                                  val fmt = CL.mkStr(
352                                        String.concatWith " " (List.tabulate(d, fn _ => !RN.gIntFormat))                                        String.concatWith " " (List.tabulate(d, fn _ => !N.gIntFormat))
353                                        ^ "\n")                                        ^ "\n")
354                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))                                  val args = List.tabulate (d, fn i => ToC.ivecIndex(outState, d, i))
355                                  in                                  in
356                                    fmt :: args                                    fmt :: args
357                                  end                                  end
# Line 335  Line 360 
360                                  val fmt = CL.mkStr(                                  val fmt = CL.mkStr(
361                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))
362                                        ^ "\n")                                        ^ "\n")
363                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))                                  val args = List.tabulate (d, fn i => ToC.vecIndex(outState, d, i))
364                                  in                                  in
365                                    fmt :: args                                    fmt :: args
366                                  end                                  end
367                              | _ => raise Fail("genStrand: unsupported output type " ^ Ty.toString ty)                              | _ => raise Fail("genStrand: unsupported output type " ^ Ty.toString ty)
368                            (* end case *))                            (* end case *))
369                        val forBody = CL.mkIfThen(CL.mkBinOp(CL.mkSubscript(CL.E_Var "status",CL.E_Var "i"), CL.#==, CL.E_Var "DIDEROT_STABILIZE"),
370                            val body = let                                                CL.mkBlock([CL.mkCall("fprintf", CL.mkVar "outS" :: prArgs)]))
371                        val body =  CL.mkFor(
372                              fun loopParams (3) =                          [(CL.intTy, "i", CL.mkInt 0)],
373                                   "x"::"y"::"k"::[]                          CL.mkBinOp(CL.E_Var "i", CL.#<, CL.E_Var "numStrands"),
374                                | loopParams (2) =                          [CL.mkPostOp(CL.E_Var "i", CL.^++)],
375                                   "x"::"y"::[]                          forBody)
                               | 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)  
376                                     in                                     in
377                                                  CL.mkFor(                        CL.D_Func(["static"], CL.voidTy, prFnName, params,
                                                         [(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.^++)],  
378                                                  body)                                                  body)
379                                     end                                     end
380                          in                          in
                                 [mkLoopNest ((loopParams nDims),0)]  
                         end  
   
                     in  
                       CL.D_Func(["static"], CL.voidTy, prFnName, params,CL.mkBlock(body))  
                     end  
               in  
381                                   prFn                                   prFn
382                end                end
383          fun genStrandTyDef (Strand{tyName, state,...}) =  
384            fun genStrandTyDef (targetTy, Strand{tyName, state,...}) =
385              (* the type declaration for the strand's state struct *)              (* the type declaration for the strand's state struct *)
386                CL.D_StructDef(                CL.D_StructDef(
387                        List.rev (List.map (fn ToCL.V(ty, x) => (ty, x)) (!state)),                  List.rev (List.map (fn x => (targetTy x, #var x)) (!state)),
388                        tyName)                        tyName)
389    
390            fun genStrandCopy(Strand{tyName,name,state,...}) = let
391                  val params = [
392                          CL.PARAM(["__global"], CL.T_Ptr(CL.T_Named tyName), "selfIn"),
393                          CL.PARAM(["__global"], CL.T_Ptr(CL.T_Named tyName), "selfOut")
394                      ]
395                    val assignStms = List.rev(List.map(fn x => CL.mkAssign(CL.mkIndirect(CL.E_Var "selfOut", #var x),
396                                                                                                                                    CL.mkIndirect(CL.E_Var "selfIn", #var x))) (!state))
397                     in
398                            CL.D_Func([""], CL.voidTy, RN.strandCopy name, params,CL.mkBlock(assignStms))
399                     end
400    
401          (* generates the load kernel function *)          (* generates the load kernel function *)
402  (* 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;",  
                                                 "}"])  
403  (* generates the opencl buffers for the image data *)  (* generates the opencl buffers for the image data *)
404          fun getGlobalDataBuffers(globals,count,contextVar,errVar) = let          fun getGlobalDataBuffers (globals,contextVar,errVar) = let
405                  val globalBuffErr = "error creating OpenCL global buffer"
406                  fun errorFn msg = CL.mkIfThen(CL.mkBinOp(CL.E_Var errVar, CL.#!=, CL.E_Var "CL_SUCCESS"),
407                        CL.mkBlock([CL.mkCall("fprintf",[CL.E_Var "stderr", CL.E_Str msg]),
408                        CL.mkCall("exit",[CL.mkInt 1])]))
409                  val globalBufferDecl =  CL.mkDecl(clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)                  val globalBufferDecl =  CL.mkDecl(clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)
410                  val globalBuffer = CL.mkAssign(CL.mkVar(concat[RN.globalsVarName,"_cl"]), CL.mkApply("clCreateBuffer",                val globalBuffer = CL.mkAssign(CL.mkVar(concat[RN.globalsVarName,"_cl"]),
411                                                                  [CL.mkVar contextVar,                      CL.mkApply("clCreateBuffer", [
412                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",                          CL.mkVar contextVar,
413                            CL.mkVar "CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR",
414                                                                  CL.mkApply("sizeof",[CL.mkVar RN.globalsTy]),                                                                  CL.mkApply("sizeof",[CL.mkVar RN.globalsTy]),
415                                                                  CL.mkVar RN.globalsVarName,                                                                  CL.mkVar RN.globalsVarName,
416                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))                          CL.mkUnOp(CL.%&,CL.mkVar errVar)
417                          ]))
418          fun genDataBuffers([],_,_,_) = []          fun genDataBuffers([],_,_,_) = []
419            | genDataBuffers((var,nDims)::globals,count,contextVar,errVar) = let                  | genDataBuffers ((var,nDims)::globals, contextVar, errVar,errFn) = let
420                        val hostVar = CL.mkIndirect(CL.mkVar RN.globalsVarName, var)
421          (* FIXME: use CL constructors to  build expressions (not strings) *)          (* FIXME: use CL constructors to  build expressions (not strings) *)
422                    val size = if nDims = 1 then                      fun sizeExp i = CL.mkSubscript(CL.mkIndirect(hostVar, "size"), CL.mkInt i)
423                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,                      val size = CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*, sizeExp 0)
424                                           CL.mkIndirect(CL.mkVar var, "size[0]"))                      val size = if (nDims > 1)
425                                          else if nDims = 2 then                            then CL.mkBinOp(size, CL.#*, sizeExp 1)
426                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,                            else size
427                                            CL.mkIndirect(CL.mkVar var, concat["size[0]", " * ", var, "->size[1]"]))                      val size = if (nDims > 2)
428                                          else                            then CL.mkBinOp(size, CL.#*, sizeExp 2)
429                                           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]"]))  
   
430                   in                   in
431                     CL.mkDecl(clMemoryTy, RN.addBufferSuffix var ,NONE)::                     CL.mkDecl(clMemoryTy, RN.addBufferSuffix var ,NONE)::
432                     CL.mkDecl(clMemoryTy, RN.addBufferSuffixData var ,NONE)::                     CL.mkDecl(clMemoryTy, RN.addBufferSuffixData var ,NONE)::
433                     CL.mkAssign(CL.mkVar(RN.addBufferSuffix var), CL.mkApply("clCreateBuffer",                        CL.mkAssign(CL.mkVar(RN.addBufferSuffix var),
434                                                                  [CL.mkVar contextVar,                          CL.mkApply("clCreateBuffer", [
435                                CL.mkVar contextVar,
436                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",
437                                                                  CL.mkApply("sizeof",[CL.mkVar (RN.imageTy nDims)]),                                                                  CL.mkApply("sizeof",[CL.mkVar (RN.imageTy nDims)]),
438                                                                  CL.mkVar var,                              hostVar,
439                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)])) ::                              CL.mkUnOp(CL.%&,CL.mkVar errVar)
440                          CL.mkAssign(CL.mkVar(RN.addBufferSuffixData var), CL.mkApply("clCreateBuffer",                            ])) ::
441                                                                  [CL.mkVar contextVar,                        errFn(concat["error in creating ",RN.addBufferSuffix var, " global buffer"]) ::
442                          CL.mkAssign(CL.mkVar(RN.addBufferSuffixData var),
443                            CL.mkApply("clCreateBuffer", [
444                                CL.mkVar contextVar,
445                                                                   CL.mkVar "CL_MEM_COPY_HOST_PTR",                                                                   CL.mkVar "CL_MEM_COPY_HOST_PTR",
446                                                                  size,                                                                  size,
447                                                                  CL.mkIndirect(CL.mkVar var,"data"),                              CL.mkIndirect(hostVar, "data"),
448                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)])):: genDataBuffers(globals,count + 2,contextVar,errVar)                              CL.mkUnOp(CL.%&,CL.mkVar errVar)
449                              ])) ::
450                            errFn(concat["error in creating ",RN.addBufferSuffixData var, " global buffer"]) ::
451                            genDataBuffers(globals,contextVar,errVar,errFn)
452                  end                  end
453          in          in
454                  [globalBufferDecl] @ [globalBuffer] @ genDataBuffers(globals,count + 2,contextVar,errVar)                  globalBufferDecl
455                    :: globalBuffer
456                    :: errorFn(globalBuffErr)
457                    :: genDataBuffers(globals,contextVar,errVar,errorFn)
458          end          end
459    
   
460  (* generates the kernel arguments for the image data *)  (* generates the kernel arguments for the image data *)
461          fun genGlobalArguments(globals,count,kernelVar,errVar) = let          fun genGlobalArguments(globals,count,kernelVar,errVar) = let
462          val globalArgument = CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,CL.mkApply("clSetKernelArg",                val globalArgErr = "error creating OpenCL global argument"
463                  fun errorFn msg = CL.mkIfThen(CL.mkBinOp(CL.E_Var errVar, CL.#!=, CL.E_Var "CL_SUCCESS"),
464                        CL.mkBlock([CL.mkCall("fprintf",[CL.E_Var "stderr", CL.E_Str msg]),
465                        CL.mkCall("exit",[CL.mkInt 1])]))
466                  val globalArgument = CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,
467                        CL.mkApply("clSetKernelArg",
468                                                                  [CL.mkVar kernelVar,                                                                  [CL.mkVar kernelVar,
469                                                                   CL.mkInt(count,CL.intTy),                         CL.mkPostOp(CL.E_Var count, CL.^++),
470                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
471                                                                   CL.mkUnOp(CL.%&,CL.mkVar(concat[RN.globalsVarName,"_cl"]))])))                                                                   CL.mkUnOp(CL.%&,CL.mkVar(concat[RN.globalsVarName,"_cl"]))])))
472                  fun genDataArguments ([],_,_,_,_) = []
473          fun genDataArguments([],_,_,_) = []                  | genDataArguments ((var,nDims)::globals,count,kernelVar,errVar,errFn) =
474            | genDataArguments((var,nDims)::globals,count,kernelVar,errVar) =                      CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.$=,
475                          CL.mkApply("clSetKernelArg",
                 CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=, CL.mkApply("clSetKernelArg",  
476                                                                  [CL.mkVar kernelVar,                                                                  [CL.mkVar kernelVar,
477                                                                   CL.mkInt(count,CL.intTy),                           CL.mkPostOp(CL.E_Var count, CL.^++),
478                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
479                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffix var))])))::                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffix var))])))::
480                             errFn(concat["error in creating ",RN.addBufferSuffix var, " argument"]) ::
481                          CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,CL.mkApply("clSetKernelArg",                      CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.$=,
482                          CL.mkApply("clSetKernelArg",
483                                                                  [CL.mkVar kernelVar,                                                                  [CL.mkVar kernelVar,
484                                                                   CL.mkInt((count + 1),CL.intTy),                           CL.mkPostOp(CL.E_Var count, CL.^++),
485                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
486                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffixData var))]))):: genDataArguments (globals, count + 2,kernelVar,errVar)                           CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffixData var))]))) ::
487                             errFn(concat["error in creating ",RN.addBufferSuffixData var, " argument"]) ::
488          in                      genDataArguments (globals,count,kernelVar,errVar,errFn)
   
                 [globalArgument] @ genDataArguments(globals,count + 1,kernelVar,errVar)  
   
         end  
   
         (* 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)  
489                in                in
490                  CL.D_Func([],CL.intTy,"main",params,body)                 [globalArgument,errorFn(globalArgErr)] @ genDataArguments(globals, count, kernelVar, errVar,errorFn)
491                end                end
492    
493        (* generates the host-side setup function *)        (* generates the globals buffers and arguments function *)
494          fun genHostSetupFunc (strand as Strand{name,tyName,...}, filename, nDims, initially, imgGlobals) = let          fun genGlobalBuffersArgs (imgGlobals) = let
495              (* 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"  
496                val errVar = "err"                val errVar = "err"
497                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")])  
498                val params = [                val params = [
499                        CL.PARAM([],CL.T_Named("cl_device_id"), deviceVar)                        CL.PARAM([],CL.T_Named("cl_context"), "context"),
500                          CL.PARAM([],CL.T_Named("cl_kernel"), "kernel"),
501                          CL.PARAM([],CL.T_Named("cl_command_queue"), "cmdQ"),
502                          CL.PARAM([],CL.T_Named("int"), "argStart")
503                      ]                      ]
504                val declarations = [                val clGlobalBuffers = getGlobalDataBuffers(!imgGlobals, "context", errVar)
505                      CL.mkDecl(clProgramTy, programVar, NONE),                val clGlobalArguments = genGlobalArguments(!imgGlobals, "argStart", "kernel", errVar)
                     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))))  
                 ]  
             (* Setup Global Variables *)  
               val globalsDecl = CL.mkDecl(  
                     CL.T_Ptr(CL.T_Named RN.globalsTy),  
                     RN.globalsVarName,  
                     SOME(CL.I_Exp(CL.mkApply("malloc", [CL.mkApply("sizeof",[CL.mkVar RN.globalsTy])]))))  
               val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.mkVar RN.globalsVarName])  
   
                 (* Retrieve the platforms  
                 val platformStm = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetPlatformIDs",  
                                                   [CL.mkInt(10,CL.intTy),  
                                                    CL.mkVar platformsVar,  
                                                    CL.mkUnOp(CL.%&,CL.mkVar numPlatformsVar)])),  
                                                    assertStm]  
   
                 val devicesStm = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetDeviceIDs",  
                                                   [CL.mkSubscript(CL.mkVar platformsVar,CL.mkInt(0,CL.intTy)),  
                                                    CL.mkVar "CL_DEVICE_TYPE_GPU",  
                                                    CL.mkInt(1,CL.intTy),  
                                                    CL.mkUnOp(CL.%&,CL.mkVar deviceVar),  
                                                    CL.mkUnOp(CL.%&,CL.mkVar numDevicesVar)])),  
                                                    assertStm] *)  
   
                 (* Create Context *)  
                 val contextStm = [CL.mkAssign(CL.mkVar contextVar, CL.mkApply("clCreateContext",  
                                                   [CL.mkInt(0,CL.intTy),  
                                                   CL.mkInt(1,CL.intTy),  
                                                   CL.mkUnOp(CL.%&,CL.mkVar deviceVar),  
                                                   CL.mkVar "NULL",  
                                                   CL.mkVar "NULL",  
                                                   CL.mkUnOp(CL.%&,CL.mkVar errVar)])),  
                                                   assertStm]  
   
                 (* 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])]  
   
   
   
506                  (* Body put all the statments together *)                  (* Body put all the statments together *)
507                  val body =  declarations @ [globalsDecl,initGlobalsCall] (*@ platformStm @ devicesStm *) @ contextStm @ commandStm @ !initially @ [strandSize] @                val body = CL.mkDecl(clIntTy, errVar, SOME(CL.I_Exp(CL.mkInt 0)))
508                                     strandsArrays @ globalAndlocalStms @ [widthDel,strands_init]  @ clStrandObjects @ clGlobalBuffers @ sourceStms  @ create_build_stms  (*@                      :: clGlobalBuffers @ clGlobalArguments
                                    kernelArguments @ clGlobalArguments @ enqueueStm @  [outputStm] @ freeStms @ outputData *)  
   
509                  in                  in
510                    CL.D_Func([],CL.voidTy,RN.globalsSetupName,params,CL.mkBlock(body))
                 CL.D_Func([],CL.voidTy,RN.setupFName,params,CL.mkBlock(body))  
   
511                  end                  end
512    
513  (* generate the data and global parameters *)  (* generate the data and global parameters *)
514          fun genKeneralGlobalParams ((name,tyname)::rest) =          fun genKeneralGlobalParams ((name,tyname)::rest) =
515                  CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), concat[RN.globalsVarName]) ::                globalParam (CL.T_Ptr(CL.T_Named (RN.imageTy tyname)), RN.addBufferSuffix name) ::
516                  CL.PARAM([], CL.T_Ptr(CL.T_Named (RN.imageTy tyname)),RN.addBufferSuffix name) ::                globalParam (CL.T_Ptr(CL.voidTy), RN.addBufferSuffixData name) ::
517                  CL.PARAM([], CL.T_Ptr(CL.voidTy),RN.addBufferSuffixData name) ::                genKeneralGlobalParams rest
518                  genKeneralGlobalParams(rest)            | genKeneralGlobalParams [] = []
   
           | genKeneralGlobalParams ([]) = []  
519    
520          (*generate code for intilizing kernel global data *)          (*generate code for intilizing kernel global data *)
521          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([]) = []  
   
522                  fun initGlobalImages((name,tyname)::rest) =                  fun initGlobalImages((name,tyname)::rest) =
523                                  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)) ::
524                                  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)) ::
525                                  initGlobalImages(rest)               initGlobalImages rest
526                    | initGlobalImages([]) = []            | initGlobalImages [] = []
                 in  
                   initGlobalStruct(globals) @ initGlobalImages(imgGlobals)  
                 end  
527    
528          (* generate the main kernel function for the .cl file *)          (* generate the main kernel function for the .cl file *)
529          fun genKernelFun(Strand{name, tyName, state, output, code,...},nDims,globals,imgGlobals) = let          fun genKernelFun (strand, nDims, globals, imgGlobals) = let
530                  val Strand{name, tyName, state, output, code,...} = strand
531                   val fName = RN.kernelFuncName;                   val fName = RN.kernelFuncName;
532                   val inState = "strand_in"                   val inState = "strand_in"
533                   val outState = "strand_out"                   val outState = "strand_out"
534                  val tempVar = "tmp"
535               val params = [               val params = [
536                        CL.PARAM(["__global"], CL.T_Ptr(CL.T_Named tyName), "selfIn"),                        CL.PARAM(["__global"], CL.T_Ptr(CL.T_Named tyName), "selfIn"),
537                        CL.PARAM(["__global"], CL.T_Ptr(CL.T_Named tyName), "selfOut"),                        CL.PARAM(["__global"], CL.T_Ptr(CL.T_Named tyName), "selfOut"),
538                        CL.PARAM(["__global"], CL.intTy, "width")                        CL.PARAM(["__global"], CL.T_Ptr(CL.T_Num(RawTypes.RT_UInt8)), "strandStatus"),
539                          CL.PARAM(["__global"], CL.intTy, "width"),
540                          CL.PARAM(["__global"], CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName)
541                      ] @ genKeneralGlobalParams(!imgGlobals)                      ] @ genKeneralGlobalParams(!imgGlobals)
542                    val thread_ids = if nDims = 1                    val thread_ids = if nDims = 1
543                          then [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),                        then [
544                                    CL.mkAssign(CL.mkVar "x",CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(0,CL.intTy)]))]                            CL.mkDecl(CL.intTy, "x",
545                                SOME(CL.I_Exp(CL.mkApply(RN.getGlobalThreadId,[CL.mkInt 0]))))
546                            ]
547                        else if nDims = 2
548                          then [
549                              CL.mkDecl(CL.intTy, "x",
550                                SOME(CL.I_Exp(CL.mkApply(RN.getGlobalThreadId,[CL.mkInt 0])))),
551                              CL.mkDecl(CL.intTy, "y",
552                                SOME(CL.I_Exp(CL.mkApply(RN.getGlobalThreadId,[CL.mkInt 1]))))
553                            ]
554                        else raise Fail "nDims > 2"
555                  val strandDecl = [
556                          CL.mkDecl(CL.T_Ptr(CL.T_Named (concat["__global ",tyName])), inState, NONE),
557                          CL.mkDecl(CL.T_Ptr(CL.T_Named (concat["__global ",tyName])), outState, NONE),
558                          CL.mkDecl(CL.T_Ptr(CL.T_Named (concat["__global ",tyName])), tempVar, NONE)
559                        ]
560                  val barrierCode = CL.mkCall(RN.strandCopy name, [CL.E_Var outState, CL.E_Var inState])
561                  val barrierStm = CL.mkCall("barrier",[CL.E_Var "CLK_LOCAL_MEM_FENCE"])
562                  val index = if nDims = 1 then
563                            CL.mkStr "x"
564                          else                          else
565                                  [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),                          CL.mkBinOp(
566                                   CL.mkDecl(CL.intTy, "y", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),                              CL.mkBinOp(CL.mkVar "x", CL.#*, CL.mkVar "width"), CL.#+, CL.mkVar "y")
                                   CL.mkAssign(CL.mkVar "x",  CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(0,CL.intTy)])),  
                                   CL.mkAssign(CL.mkVar "y",CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(1,CL.intTy)]))]  
   
                   val strandDecl = [CL.mkDecl(CL.T_Named tyName, inState, NONE),  
                                                         CL.mkDecl(CL.T_Named tyName, outState,NONE)]  
                   val strandObjects  = if nDims = 1  
                         then [CL.mkAssign(CL.mkSubscript(CL.mkVar "selfIn",CL.mkStr "x"),  
                                                                          CL.mkVar inState),  
                                   CL.mkAssign(CL.mkSubscript(CL.mkVar "selfOut",CL.mkStr "x"),  
                                                                          CL.mkVar outState)]  
                         else let  
                                 val index = CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")  
                                 in  
                                         [CL.mkAssign(CL.mkSubscript(CL.mkVar "selfIn",index),  
                                                                         CL.mkVar inState),  
                                          CL.mkAssign(CL.mkSubscript(CL.mkVar "selfOut",index),  
                                                                         CL.mkVar outState)]  
                                 end  
567    
568                  val strandObjects =
569                         [ CL.mkAssign(CL.mkVar inState,  CL.mkBinOp(CL.mkVar "selfIn",CL.#+,index)),
570                           CL.mkAssign(CL.mkVar outState, CL.mkBinOp(CL.mkVar "selfOut",CL.#+,index))
571                         ]
572    
573                    val status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy))))                  val stabalizeStm = CL.mkAssign(CL.mkSubscript(CL.mkVar "strandStatus",index),
574                    val local_vars =  thread_ids @ initKernelGlobals(!globals,!imgGlobals)  @ strandDecl @ strandObjects @ [status]                                                                          CL.E_Var "status")
575                    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 status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.mkSubscript(CL.mkVar "strandStatus",index))))
576                    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 strandInitStm = CL.mkCall(RN.strandInit name, [
577                                                          CL.mkCall(RN.strandStabilize name,[ CL.mkUnOp(CL.%&,CL.mkVar inState),  CL.mkUnOp(CL.%&,CL.mkVar outState)])]                        CL.E_Var RN.globalsVarName,
578                          CL.E_Var outState,
579                    val whileBlock = [CL.mkWhile(while_exp,CL.mkBlock while_body)]                        CL.E_Var "x",
580    (* FIXME: if nDims = 1, then "y" is not defined! the arguments to this call should really come from
581                    val body = CL.mkBlock(local_vars  @ whileBlock)   * the initially code!
582     *)
583                          CL.E_Var "y"])
584                  val local_vars = thread_ids
585                        @ initGlobalImages(!imgGlobals)
586                        @ strandDecl
587                        @ strandObjects
588                        @ [strandInitStm,status]
589                  val while_exp = CL.mkBinOp(CL.mkVar "status",CL.#==, CL.mkVar RN.kActive)
590                  val whileBody = CL.mkBlock ([barrierCode,barrierStm] @ [
591                          CL.mkAssign(CL.mkVar "status",
592                            CL.mkApply(RN.strandUpdate name,
593                              [CL.mkVar inState, CL.mkVar outState,CL.E_Var RN.globalsVarName]))] )
594                  val whileBlock = [CL.mkWhile(while_exp, whileBody)]
595                  val body = CL.mkBlock(local_vars @ whileBlock @ [stabalizeStm])
596                  in                  in
597                     CL.D_Func(["__kernel"], CL.voidTy, fName, params, body)                     CL.D_Func(["__kernel"], CL.voidTy, fName, params, body)
598                  end                  end
599          (* generate a global structure from the globals *)          (* generate a global structure from the globals *)
600          fun genGlobalStruct(globals) = let          fun genGlobalStruct (targetTy, globals) = let
601                   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
                    | getGlobals([]) = []  
                    | getGlobals(_::rest) = getGlobals(rest)  
602                   in                   in
603                          CL.D_StructDef(getGlobals(globals),RN.globalsTy)                  CL.D_StructDef(globs, RN.globalsTy)
604                  end
605    
606            fun genGlobals (declFn, targetTy, globals) = let
607                  fun doVar (x : mirror_var) = declFn (CL.D_Var([], targetTy x, #var x, NONE))
608                  in
609                    List.app doVar globals
610                  end
611    
612            fun genStrandDesc (Strand{name, output, ...}) = let
613                (* the strand's descriptor object *)
614                  val descI = let
615                        fun fnPtr (ty, f) = CL.I_Exp(CL.mkCast(CL.T_Named ty, CL.mkVar f))
616                        val SOME(outTy, _) = !output
617                        in
618                          CL.I_Struct[
619                              ("name", CL.I_Exp(CL.mkStr name)),
620                              ("stateSzb", CL.I_Exp(CL.mkSizeof(CL.T_Named(N.strandTy name)))),
621    (*
622                              ("outputSzb", CL.I_Exp(CL.mkSizeof(ToC.trTy outTy))),
623    *)
624                              ("update", fnPtr("update_method_t", "0")),
625                              ("print", fnPtr("print_method_cl_t", name ^ "_print"))
626                            ]
627                        end
628                  val desc = CL.D_Var([], CL.T_Named N.strandDescTy, N.strandDesc name, SOME descI)
629                  in
630                    desc
631                    end                    end
632    
633        (* generate the table of strand descriptors *)        (* generate the table of strand descriptors *)
634          fun genStrandTable (ppStrm, strands) = let          fun genStrandTable (declFn, strands) = let
635                val nStrands = length strands                val nStrands = length strands
636                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)))
637                fun genInits (_, []) = []                fun genInits (_, []) = []
638                  | 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)  
639                in                in
640                  ppDecl (CL.D_Var([], CL.int32, RN.numStrands,                  declFn (CL.D_Var([], CL.int32, N.numStrands,
641                    SOME(CL.I_Exp(CL.mkInt(IntInf.fromInt nStrands, CL.int32)))));                    SOME(CL.I_Exp(CL.E_Int(IntInf.fromInt nStrands, CL.int32)))));
642                  ppDecl (CL.D_Var([],                  declFn (CL.D_Var([],
643                    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),
644                    RN.strands,                    N.strands,
645                    SOME(CL.I_Array(genInits (0, strands)))))                    SOME(CL.I_Array(genInits (0, strands)))))
646                end                end
647    
648            fun genSrc (baseName, prog) = let
649          fun genSrc (baseName, Prog{double,globals, topDecls, strands, initially,imgGlobals,numDims,...}) = let                val Prog{name,double, globals, topDecls, strands, initially, imgGlobals, numDims, ...} = prog
650                val clFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "cl"}                val clFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "cl"}
651                val cFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "c"}                val cFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "c"}
652                val clOutS = TextIO.openOut clFileName                val clOutS = TextIO.openOut clFileName
653                val cOutS = TextIO.openOut cFileName                val cOutS = TextIO.openOut cFileName
654  (* FIXME: need to use PrintAsC and PrintAsCL *)                val clppStrm = PrintAsCL.new clOutS
               val clppStrm = PrintAsC.new clOutS  
655                val cppStrm = PrintAsC.new cOutS                val cppStrm = PrintAsC.new cOutS
656                  val progName = name
657                fun cppDecl dcl = PrintAsC.output(cppStrm, dcl)                fun cppDecl dcl = PrintAsC.output(cppStrm, dcl)
658                fun clppDecl dcl = PrintAsCL.output(clppStrm, dcl)                fun clppDecl dcl = PrintAsCL.output(clppStrm, dcl)
659                val strands = AtomTable.listItems strands                val strands = AtomTable.listItems strands
# Line 940  Line 665 
665                        then "#define DIDEROT_DOUBLE_PRECISION"                        then "#define DIDEROT_DOUBLE_PRECISION"
666                        else "#define DIDEROT_SINGLE_PRECISION",                        else "#define DIDEROT_SINGLE_PRECISION",
667                      "#define DIDEROT_TARGET_CL",                      "#define DIDEROT_TARGET_CL",
668                      "#include \"Diderot/cl-types.h\""                      "#include \"Diderot/cl-diderot.h\""
669                    ]));                    ]));
670                  List.app clppDecl (List.rev (!globals));                  clppDecl (genGlobalStruct (#gpuTy, !globals));
671                  clppDecl (genGlobalStruct (!globals));                  clppDecl (genStrandTyDef(#gpuTy, strand));
672                  clppDecl (genStrandTyDef strand);                  clppDecl  (!init_code);
673                    clppDecl  (genStrandCopy(strand));
674                  List.app clppDecl (!code);                  List.app clppDecl (!code);
675                  clppDecl (genKernelFun (strand,!numDims,globals,imgGlobals));                  clppDecl (genKernelFun (strand,!numDims,globals,imgGlobals));
676                (* Generate the Host file .c *)                (* Generate the Host C file *)
677                  cppDecl (CL.D_Verbatim([                  cppDecl (CL.D_Verbatim([
678                      if double                      if double
679                        then "#define DIDEROT_DOUBLE_PRECISION"                        then "#define DIDEROT_DOUBLE_PRECISION"
# Line 955  Line 681 
681                      "#define DIDEROT_TARGET_CL",                      "#define DIDEROT_TARGET_CL",
682                      "#include \"Diderot/diderot.h\""                      "#include \"Diderot/diderot.h\""
683                    ]));                    ]));
684                  List.app cppDecl (List.rev (!globals));                  cppDecl (CL.D_Var(["static"], CL.charPtr, "ProgramName",
685                  cppDecl (genGlobalStruct (!globals));                    SOME(CL.I_Exp(CL.mkStr progName))));
686                  cppDecl (genStrandTyDef strand);                  cppDecl (genGlobalStruct (#hostTy, !globals));
687                  cppDecl  (!init_code);                  cppDecl (CL.D_Var(["static"], CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName, NONE));
688                  cppDecl (genStrandInit(strand,!numDims));                  cppDecl (genStrandTyDef (#hostTy, strand));
689                  cppDecl (genStrandPrint(strand,!numDims));                  cppDecl (genStrandPrint strand);
                 (* cppDecl (genKernelLoader());*)  
690                  List.app cppDecl (List.rev (!topDecls));                  List.app cppDecl (List.rev (!topDecls));
691                  cppDecl (genHostSetupFunc (strand, clFileName, !numDims, initially, imgGlobals));                  cppDecl (genGlobalBuffersArgs imgGlobals);
692                    List.app (fn strand => cppDecl (genStrandDesc strand)) strands;
693                    genStrandTable (cppDecl, strands);
694                    cppDecl (!initially);
695                  PrintAsC.close cppStrm;                  PrintAsC.close cppStrm;
696                  PrintAsCL.close clppStrm;                  PrintAsCL.close clppStrm;
697                  TextIO.closeOut cOutS;                  TextIO.closeOut cOutS;
# Line 1027  Line 755 
755          fun init (Strand{name, tyName, code,init_code, ...}, params, init) = let          fun init (Strand{name, tyName, code,init_code, ...}, params, init) = let
756                val fName = RN.strandInit name                val fName = RN.strandInit name
757                val params =                val params =
758                      CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfOut") ::                      globalParam (globPtrTy, RN.globalsVarName) ::
759                        globalParam (CL.T_Ptr(CL.T_Named tyName), "selfOut") ::
760                        List.map (fn (ToCL.V(ty, x)) => CL.PARAM([], ty, x)) params                        List.map (fn (ToCL.V(ty, x)) => CL.PARAM([], ty, x)) params
761                val initFn = CL.D_Func([], CL.voidTy, fName, params, init)                val initFn = CL.D_Func([], CL.voidTy, fName, params, init)
762                in                in
# Line 1038  Line 767 
767          fun method (Strand{name, tyName, code,...}, methName, body) = let          fun method (Strand{name, tyName, code,...}, methName, body) = let
768                val fName = concat[name, "_", methName]                val fName = concat[name, "_", methName]
769                val params = [                val params = [
770                        CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfIn"),                        globalParam (CL.T_Ptr(CL.T_Named tyName), "selfIn"),
771                        CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfOut")                        globalParam (CL.T_Ptr(CL.T_Named tyName), "selfOut"),
772                          globalParam (CL.T_Ptr(CL.T_Named (RN.globalsTy)), RN.globalsVarName)
773                      ]                      ]
774                val methFn = CL.D_Func([], CL.int32, fName, params, body)                val methFn = CL.D_Func([], CL.int32, fName, params, body)
775                in                in

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

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