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

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

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