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 1459, Sun Aug 7 17:12:19 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       (* translate TreeIL types to shadow types *)
32        fun convertToShadow (ty, name) = (case ty
33               of Ty.IVecTy 1 => CL.mkAssign(
34                    CL.mkSelect(CL.mkVar(RN.shadowGlaobalsName),name),
35                    CL.mkIndirect(CL.mkVar(RN.globalsVarName), name))
36                | Ty.TensorTy[n]=> CL.mkCall(RN.convertToShadowVec n, [
37                      CL.mkSelect(CL.mkVar(RN.shadowGlaobalsName),name),
38                      CL.mkIndirect(CL.mkVar(RN.globalsVarName), name)
39                    ])
40                | Ty.ImageTy(ImageInfo.ImgInfo{dim, ...}) => CL.mkCall(RN.shadowImageFunc dim, [
41                      CL.mkVar "context",
42                      CL.mkUnOp(CL.%&,CL.mkSelect(CL.mkVar(RN.shadowGlaobalsName),name)),
43                      CL.mkIndirect(CL.mkVar(RN.globalsVarName),name)
44                    ])
45                | Ty.TensorTy[n, m] => CL.mkCall(RN.convertToShadowMat(m,n), [
46                      CL.mkSelect(CL.mkVar(RN.shadowGlaobalsName),name),
47                      CL.mkIndirect(CL.mkVar(RN.globalsVarName), name)
48                    ])
49                | _ => CL.mkAssign(
50                    CL.mkSelect(CL.mkVar(RN.shadowGlaobalsName),name),
51                    CL.mkIndirect(CL.mkVar(RN.globalsVarName), name))
52              (* end case *))
53    
54      (* helper functions for specifying parameters in various address spaces *)
55        fun clParam (spc, ty, x) = CL.PARAM([spc], ty, x)
56        fun globalParam (ty, x) = CL.PARAM(["__global"], ty, x)
57        fun constantParam (ty, x) = CL.PARAM(["__constant"], ty, x)
58        fun localParam (ty, x) = CL.PARAM(["__local"], ty, x)
59        fun privateParam (ty, x) = CL.PARAM(["__private"], ty, x)
60    
61      (* OpenCL global pointer type *)
62        fun globalPtr ty = CL.T_Qual("__global", CL.T_Ptr ty)
63    
64    (* C variable translation *)    (* C variable translation *)
65      structure TrCVar =      structure TrCVar =
# Line 25  Line 72 
72        (* 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) *)
73          fun lvalueVar (env, x) = (case V.kind x          fun lvalueVar (env, x) = (case V.kind x
74                 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))
75                  | IL.VK_State strand => raise Fail "unexpected strand context"                  | IL.VK_State strand => CL.mkIndirect(CL.mkVar "selfOut", lookup(env, x))
76                  | IL.VK_Local => CL.mkVar(lookup(env, x))                  | IL.VK_Local => CL.mkVar(lookup(env, x))
77                (* end case *))                (* end case *))
78        (* translate a variable that occurs in an r-value context *)        (* translate a variable that occurs in an r-value context *)
79          val rvalueVar = lvalueVar          fun rvalueVar (env, x) = (case V.kind x
80                   of IL.VK_Global => CL.mkIndirect(CL.mkVar RN.globalsVarName, lookup(env, x))
81                    | IL.VK_State strand => CL.mkIndirect(CL.mkVar "selfIn", lookup(env, x))
82                    | IL.VK_Local => CL.mkVar(lookup(env, x))
83                  (* end case *))
84        end        end
85    
86      structure ToC = TreeToCFn (TrCVar)      structure ToC = TreeToCFn (TrCVar)
# Line 39  Line 90 
90      type stm = CL.stm      type stm = CL.stm
91    
92    (* OpenCL specific types *)    (* OpenCL specific types *)
93        val clIntTy = CL.T_Named "cl_int"
94      val clProgramTy = CL.T_Named "cl_program"      val clProgramTy = CL.T_Named "cl_program"
95      val clKernelTy  = CL.T_Named "cl_kernel"      val clKernelTy  = CL.T_Named "cl_kernel"
96      val clCmdQueueTy = CL.T_Named "cl_command_queue"      val clCmdQueueTy = CL.T_Named "cl_command_queue"
# Line 46  Line 98 
98      val clDeviceIdTy = CL.T_Named "cl_device_id"      val clDeviceIdTy = CL.T_Named "cl_device_id"
99      val clPlatformIdTy = CL.T_Named "cl_platform_id"      val clPlatformIdTy = CL.T_Named "cl_platform_id"
100      val clMemoryTy = CL.T_Named "cl_mem"      val clMemoryTy = CL.T_Named "cl_mem"
101        val globPtrTy = CL.T_Ptr(CL.T_Named RN.globalsTy)
102    
103      (* variable or field that is mirrored between host and GPU *)
104        type mirror_var = {
105    (* FIXME: perhaps it would be cleaner to just track the TreeIL type of the variable? *)
106                hostTy : CL.ty,             (* variable type on Host (i.e., C type) *)
107                shadowTy : CL.ty,           (* host-side shadow type of GPU type *)
108                gpuTy : CL.ty,              (* variable's type on GPU (i.e., OpenCL type) *)
109                hToS: stm,                  (* the statement that converts the variable to its *)
110                                            (* shadow representation *)
111                var : CL.var                (* variable name *)
112              }
113    
114      datatype strand = Strand of {      datatype strand = Strand of {
115          name : string,          name : string,
116          tyName : string,          tyName : string,
117          state : var list ref,          state : mirror_var list ref,
118          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) *)
119          code : CL.decl list ref,          code : CL.decl list ref,
120          init_code: CL.decl ref          init_code: CL.decl ref
# Line 61  Line 125 
125          double : bool,                  (* true for double-precision support *)          double : bool,                  (* true for double-precision support *)
126          parallel : bool,                (* true for multithreaded (or multi-GPU) target *)          parallel : bool,                (* true for multithreaded (or multi-GPU) target *)
127          debug : bool,                   (* true for debug support in executable *)          debug : bool,                   (* true for debug support in executable *)
128          globals : CL.decl list ref,          globals : mirror_var list ref,
129          topDecls : CL.decl list ref,          topDecls : CL.decl list ref,
130          strands : strand AtomTable.hash_table,          strands : strand AtomTable.hash_table,
131          initially : CL.stm list ref,          initially :  CL.decl ref,
132          numDims: int ref,          numDims: int ref,               (* number of dimensions in initially iteration *)
133          imgGlobals: (string * int) list ref,          imgGlobals: (string * int) list ref,
134          prFn: CL.decl ref          prFn: CL.decl ref
135        }        }
# Line 85  Line 149 
149        | GlobalScope        | GlobalScope
150        | InitiallyScope        | InitiallyScope
151        | StrandScope of TreeIL.var list  (* strand initialization *)        | StrandScope of TreeIL.var list  (* strand initialization *)
152        | MethodScope of TreeIL.var list  (* method body; vars are state variables *)        | MethodScope of MethodName.name * TreeIL.var list  (* method body; vars are state variables *)
153    
154    (* the supprted widths of vectors of reals on the target. *)    (* the supprted widths of vectors of reals on the target. *)
155  (* FIXME: for OpenCL 1.1, 3 is also valid *)  (* FIXME: for OpenCL 1.1, 3 is also valid *)
# Line 98  Line 162 
162    (* TreeIL to target translations *)    (* TreeIL to target translations *)
163      structure Tr =      structure Tr =
164        struct        struct
       (* this function is used for the initially clause, so it generates OpenCL *)  
165          fun fragment (ENV{info, vMap, scope}, blk) = let          fun fragment (ENV{info, vMap, scope}, blk) = let
166                val (vMap, stms) = ToCL.trFragment (vMap, blk)                val (vMap, stms) = (case scope
167                         of GlobalScope => ToC.trFragment (vMap, blk)
168    (* NOTE: if we move strand initialization to the GPU, then we'll have to change the following code! *)
169                          | InitiallyScope => ToC.trFragment (vMap, blk)
170                          | _ => ToCL.trFragment (vMap, blk)
171                        (* end case *))
172                in                in
173                  (ENV{info=info, vMap=vMap, scope=scope}, stms)                  (ENV{info=info, vMap=vMap, scope=scope}, stms)
174                end                end
175          fun saveState cxt stateVars (env, args, stm) = (          fun block (ENV{vMap, scope, ...}, blk) = let
176                  fun saveState cxt stateVars trAssign (env, args, stm) = (
177                ListPair.foldrEq                ListPair.foldrEq
178                  (fn (x, e, stms) => ToCL.trAssign(env, x, e)@stms)                        (fn (x, e, stms) => trAssign(env, x, e)@stms)
179                    [stm]                    [stm]
180                      (stateVars, args)                      (stateVars, args)
181                ) handle ListPair.UnequalLengths => (                ) handle ListPair.UnequalLengths => (
182                  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"]);
183                  raise Fail(concat["saveState ", cxt, ": length mismatch"]))                  raise Fail(concat["saveState ", cxt, ": length mismatch"]))
184          fun block (ENV{vMap, scope, ...}, blk) = (case scope                in
185                 of StrandScope stateVars => ToCL.trBlock (vMap, saveState "StrandScope" stateVars, blk)                  case scope
186                  | 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! *)
187                     of StrandScope stateVars =>
188                          ToCL.trBlock (vMap, saveState "StrandScope" stateVars ToCL.trAssign, blk)
189                      | MethodScope(name, stateVars) =>
190                          ToCL.trBlock (vMap, saveState "MethodScope" stateVars ToCL.trAssign, blk)
191                  | InitiallyScope => ToCL.trBlock (vMap, fn (_, _, stm) => [stm], blk)                  | InitiallyScope => ToCL.trBlock (vMap, fn (_, _, stm) => [stm], blk)
192                  | _ => ToC.trBlock (vMap, fn (_, _, stm) => [stm], blk)                  | _ => ToC.trBlock (vMap, fn (_, _, stm) => [stm], blk)
193                (* end case *))                  (* end case *)
194                  end
195          fun exp (ENV{vMap, ...}, e) = ToCL.trExp(vMap, e)          fun exp (ENV{vMap, ...}, e) = ToCL.trExp(vMap, e)
196        end        end
197    
198    (* variables *)    (* variables *)
199      structure Var =      structure Var =
200        struct        struct
201            fun mirror (ty, name) = {
202                    hostTy = ToC.trType ty,
203                    shadowTy = shadowTy ty,
204                    gpuTy = ToCL.trType ty,
205                    hToS = convertToShadow(ty,name),
206                    var = name
207                  }
208          fun name (ToCL.V(_, name)) = name          fun name (ToCL.V(_, name)) = name
209          fun global (Prog{globals, imgGlobals, ...}, name, ty) = let          fun global (Prog{globals, imgGlobals, ...}, name, ty) = let
210                val ty' = ToCL.trType ty                val x = mirror (ty, name)
211                fun isImgGlobal (imgGlobals, Ty.ImageTy(ImageInfo.ImgInfo{dim, ...}), name) =  imgGlobals  := (name,dim):: !imgGlobals                fun isImgGlobal (Ty.ImageTy(ImageInfo.ImgInfo{dim, ...}), name) =
212                  | isImgGlobal (imgGlobals, _, _) =  ()                      imgGlobals  := (name,dim) :: !imgGlobals
213                in                  | isImgGlobal _ =  ()
214                  globals := CL.D_Var([], ty', name, NONE) :: !globals;                in
215                  isImgGlobal(imgGlobals,ty,name);                  globals := x :: !globals;
216                  ToCL.V(ty', name)                  isImgGlobal (ty, name);
217                    ToCL.V(#gpuTy x, name)
218                end                end
219          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)
220          fun state (Strand{state, ...}, x) = let          fun state (Strand{state, ...}, x) = let
221                val ty' = ToCL.trType(V.ty x)                val ty = V.ty x
222                val x' = ToCL.V(ty', V.name x)                val x' = mirror (ty, V.name x)
223                in                in
224                  state := x' :: !state;                  state := x' :: !state;
225                  x'                  ToCL.V(#gpuTy x', #var x')
226                end                end
227        end        end
228    
# Line 158  Line 240 
240          val scopeGlobal = setScope GlobalScope          val scopeGlobal = setScope GlobalScope
241          val scopeInitially = setScope InitiallyScope          val scopeInitially = setScope InitiallyScope
242          fun scopeStrand (env, svars) = setScope (StrandScope svars) env          fun scopeStrand (env, svars) = setScope (StrandScope svars) env
243          fun scopeMethod (env, svars) = setScope (MethodScope svars) env          fun scopeMethod (env, name, svars) = setScope (MethodScope(name, svars)) env
244        (* bind a TreeIL varaiable to a target variable *)        (* bind a TreeIL varaiable to a target variable *)
245          fun bind (ENV{info, vMap, scope}, x, x') = ENV{          fun bind (ENV{info, vMap, scope}, x, x') = ENV{
246                  info = info,                  info = info,
# Line 179  Line 261 
261                    globals = ref [],                    globals = ref [],
262                    topDecls = ref [],                    topDecls = ref [],
263                    strands = AtomTable.mkTable (16, Fail "strand table"),                    strands = AtomTable.mkTable (16, Fail "strand table"),
264                    initially = ref([CL.S_Comment["missing initially"]]),                    initially = ref(CL.D_Comment["missing initially"]),
265                                    numDims = ref(0),                    numDims = ref 0,
266                                    imgGlobals = ref[],                                    imgGlobals = ref[],
267                                    prFn = ref(CL.D_Comment(["No Print Function"]))                                    prFn = ref(CL.D_Comment(["No Print Function"]))
268                  })                  })
       (* 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  
269    
270        (* 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 *)
271          fun inputs (Prog{topDecls, ...}, stm) = let          fun inputs (Prog{topDecls, ...}, stm) = let
# Line 207  Line 279 
279    
280        (* register the global initialization part of a program *)        (* register the global initialization part of a program *)
281          fun init (Prog{topDecls, ...}, init) = let          fun init (Prog{topDecls, ...}, init) = let
282                val globPtrTy = CL.T_Ptr(CL.T_Named RN.globalsTy)                val globalsDecl = CL.mkAssign(CL.E_Var RN.globalsVarName,
283                        CL.mkApply("malloc", [CL.mkSizeof(CL.T_Named RN.globalsTy)]))
284                val initFn = CL.D_Func(                val initFn = CL.D_Func(
285                      [], CL.voidTy, RN.initGlobals, [CL.PARAM([], globPtrTy, RN.globalsVarName)],                      [], CL.voidTy, RN.initGlobals, [],
286                        CL.mkBlock[
287                            globalsDecl,
288                            CL.mkCall(RN.initGlobalsHelper, [CL.mkVar RN.globalsVarName])
289                          ])
290                  val initHelperFn = CL.D_Func(
291                        [], CL.voidTy, RN.initGlobalsHelper,
292                        [CL.PARAM([], globPtrTy, RN.globalsVarName)],
293                      init)                      init)
294                val shutdownFn = CL.D_Func(                val shutdownFn = CL.D_Func(
295                      [], CL.voidTy, RN.shutdown,                      [], CL.voidTy, RN.shutdown,
296                      [CL.PARAM([], CL.T_Ptr(CL.T_Named RN.worldTy), "wrld")],                      [CL.PARAM([], CL.T_Ptr(CL.T_Named RN.worldTy), "wrld")],
297                      CL.S_Block[])                      CL.S_Block[])
298                in                in
299                  topDecls := shutdownFn :: initFn :: !topDecls                  topDecls := shutdownFn :: initFn :: initHelperFn :: !topDecls
300                end                end
301    
302        (* create and register the initially function for a program *)        (* create and register the initially function for a program *)
303          fun initially {          fun initially {
304                prog = Prog{strands, initially, numDims,...},                prog = Prog{name=progName, strands, initially,numDims, ...},
305                isArray : bool,                isArray : bool,
306                iterPrefix : stm list,                iterPrefix : stm list,
307                iters : (var * exp * exp) list,                iters : (var * exp * exp) list,
# Line 231  Line 311 
311              } = let              } = let
312                val name = Atom.toString strand                val name = Atom.toString strand
313                val nDims = List.length iters                val nDims = List.length iters
314                  val worldTy = CL.T_Ptr(CL.T_Named N.worldTy)
315                fun mapi f xs = let                fun mapi f xs = let
316                      fun mapf (_, []) = []                      fun mapf (_, []) = []
317                        | 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 320 
320                      end                      end
321                val baseInit = mapi (fn (i, (_, e, _)) => (i, CL.I_Exp e)) iters                val baseInit = mapi (fn (i, (_, e, _)) => (i, CL.I_Exp e)) iters
322                val sizeInit = mapi                val sizeInit = mapi
323                      (fn (i, (ToCL.V(ty, _), lo, hi)) =>                      (fn (i, (CL.V(ty, _), lo, hi)) =>
324                          (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))))
325                      ) iters                      ) iters
326                    val numStrandsVar = "numStrandsVar"              (* code to allocate the world and initial strands *)
327                val allocCode = iterPrefix @ [                val wrld = "wrld"
328                  val allocCode = [
329                        CL.mkComment["allocate initial block of strands"],                        CL.mkComment["allocate initial block of strands"],
330                        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)),
331                        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)),
332                        CL.mkDecl(CL.int32,"numDims",SOME(CL.I_Exp(CL.mkInt(IntInf.fromInt nDims, CL.int32))))                        CL.mkDecl(worldTy, wrld,
333                      ]                          SOME(CL.I_Exp(CL.E_Apply(N.allocInitially, [
334                val numStrandsLoopBody =                              CL.mkVar "ProgramName",
335                      CL.mkExpStm(CL.mkAssignOp(CL.mkVar numStrandsVar, CL.*=,CL.mkSubscript(CL.mkVar "size",CL.mkVar "i")))                              CL.mkUnOp(CL.%&, CL.E_Var(N.strandDesc name)),
336                val numStrandsLoop =  CL.mkFor([(CL.intTy, "i", CL.mkInt(0,CL.intTy))],                              CL.E_Bool isArray,
337                      CL.mkBinOp(CL.mkVar "i", CL.#<, CL.mkVar "numDims"),                              CL.E_Int(IntInf.fromInt nDims, CL.int32),
338                      [CL.mkPostOp(CL.mkVar "i", CL.^++)], numStrandsLoopBody)                              CL.E_Var "base",
339                in                              CL.E_Var "size"
340                  numDims := nDims;                            ]))))
                 initially := allocCode @ [numStrandsLoop]  
               end  
   
   
       (***** OUTPUT *****)  
         fun genStrandInit (Strand{name,tyName,state,output,code,...}, nDims) = let  
               val params = [  
                       CL.PARAM([], CL.T_Ptr(CL.uint32), "sizes" ),  
                       CL.PARAM([], CL.intTy, "width"),  
                       CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "strands")  
                     ]  
               val body = let  
                     fun loopParams 3 = ["x", "y", "k"]  
                       | loopParams 2 = ["x", "y"]  
                       | loopParams 1 = ["x"]  
                       | loopParams _ = raise Fail "genStrandInit: missing size dim"  
                     fun mkLoopNest ([], _, nDims) = if nDims = 1  
                           then CL.mkBlock [  
                               CL.mkCall(RN.strandInit name, [  
                                 CL.mkUnOp(CL.%&,CL.mkSubscript(CL.mkVar "strands",CL.mkStr "x")),  
                                                 CL.mkVar "x"])  
341                              ]                              ]
342                            else let              (* create the loop nest for the initially iterations *)
343                              val index = CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")                val indexVar = "ix"
344                              in                val strandTy = CL.T_Ptr(CL.T_Named(N.strandTy name))
345                                CL.mkBlock([CL.mkCall(RN.strandInit name, [CL.mkUnOp(CL.%&,CL.mkSubscript(CL.mkVar "strands",index)),                fun mkLoopNest [] = CL.mkBlock(createPrefix @ [
346                                CL.mkVar "x", CL.mkVar"y"])])                        CL.mkDecl(strandTy, "sp",
347                              end                          SOME(CL.I_Exp(
348                        | mkLoopNest (param::rest,count,nDims) = let                            CL.E_Cast(strandTy,
349                            val body = mkLoopNest (rest, count + 1,nDims)                            CL.E_Apply(N.inState, [CL.E_Var "wrld", CL.mkBinOp(CL.mkVar indexVar, CL.#*, CL.mkSizeof(CL.T_Named (N.strandDesc name)))]))))),
350                          CL.mkCall(N.strandInit name, CL.E_Var "sp" :: args),
351                          CL.mkAssign(CL.E_Var indexVar, CL.mkBinOp(CL.E_Var indexVar, CL.#+, CL.E_Int(1, CL.uint32)))
352                        ])
353                    | mkLoopNest ((CL.V(ty, param), lo, hi)::iters) = let
354                        val body = mkLoopNest iters
355                            in                            in
356                              CL.mkFor(                              CL.mkFor(
357                                  [(CL.intTy, param, CL.mkInt(0,CL.intTy))],                          [(ty, param, lo)],
358                                  CL.mkBinOp(CL.mkVar param, CL.#<, CL.mkSubscript(CL.mkVar "sizes",CL.mkInt(count,CL.intTy))),                          CL.mkBinOp(CL.E_Var param, CL.#<=, hi),
359                                  [CL.mkPostOp(CL.mkVar param, CL.^++)],                          [CL.mkPostOp(CL.E_Var param, CL.^++)],
360                                  body)                                  body)
361                            end                            end
362                  val iterCode = [
363                          CL.mkComment["initially"],
364                          CL.mkDecl(CL.uint32, indexVar, SOME(CL.I_Exp(CL.E_Int(0, CL.uint32)))),
365                          mkLoopNest iters
366                        ]
367                  val body = CL.mkBlock(
368                        iterPrefix @
369                        allocCode @
370                        iterCode @
371                        [CL.mkReturn(SOME(CL.E_Var "wrld"))])
372                  val initFn = CL.D_Func([], worldTy, N.initially, [], body)
373                      in                      in
374                        [mkLoopNest ((loopParams nDims),0,nDims)]                  numDims := nDims;
375                      end                  initially := initFn
                 in  
                   CL.D_Func(["static"], CL.voidTy, RN.strandInitSetup, params,CL.mkBlock(body))  
376                  end                  end
377    
378          fun genStrandPrint (Strand{name, tyName, state, output, code,...},nDims) = let  
379          (***** OUTPUT *****)
380    (* FIXME: I think that the iteration and test for stable strands can be moved into the runtime, which
381     * will make the print function compatible with the C target version.
382     *)
383            fun genStrandPrint (Strand{name, tyName, state, output, code, ...}) = let
384              (* the print function *)              (* the print function *)
385                val prFnName = concat[name, "_print"]                val prFnName = concat[name, "Print"]
386                val prFn = let                val prFn = let
387                      val params = [                      val params = [
388                            CL.PARAM([], CL.T_Ptr(CL.T_Named "FILE"), "outS"),                            CL.PARAM([], CL.T_Ptr(CL.T_Named "FILE"), "outS"),
                           CL.PARAM([], CL.T_Ptr(CL.uint32), "sizes" ),  
                           CL.PARAM([], CL.intTy, "width"),  
389                            CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "self")                            CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "self")
390                          ]                          ]
   
391                     val SOME(ty, x) = !output                     val SOME(ty, x) = !output
392                     val outState = if nDims = 1 then                      val outState = CL.mkIndirect(CL.mkVar "self", x)
                           CL.mkSelect(CL.mkSubscript(CL.mkVar "self",CL.mkVar "x"), x)  
                         else if nDims = 2 then  
                                 CL.mkSelect(CL.mkSubscript(CL.mkVar "self",  
                                    CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")), x)  
   
                         else CL.mkSelect(CL.mkVar "self",x)  
   
393                      val prArgs = (case ty                      val prArgs = (case ty
394                             of Ty.IVecTy 1 => [CL.mkStr(!RN.gIntFormat ^ "\n"), outState]                             of Ty.IVecTy 1 => [CL.E_Str(!N.gIntFormat ^ "\n"), outState]
395                              | Ty.IVecTy d => let                              | Ty.IVecTy d => let
396                                    fun sel i = CL.mkSubscript(CL.mkSelect(outState, "s"),
397                                          CL.mkInt(IntInf.fromInt i))
398                                  val fmt = CL.mkStr(                                  val fmt = CL.mkStr(
399                                        String.concatWith " " (List.tabulate(d, fn _ => !RN.gIntFormat))                                        String.concatWith " " (List.tabulate(d, fn _ => !N.gIntFormat))
400                                        ^ "\n")                                        ^ "\n")
401                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))                                  val args = List.tabulate (d, sel)
402                                  in                                  in
403                                    fmt :: args                                    fmt :: args
404                                  end                                  end
405                              | Ty.TensorTy[] => [CL.mkStr "%f\n", outState]                              | Ty.TensorTy[] => [CL.mkStr "%f\n", outState]
406                              | Ty.TensorTy[d] => let                              | Ty.TensorTy[d] => let
407                                    fun sel i = CL.mkSubscript(CL.mkSelect(outState, "s"),
408                                          CL.mkInt(IntInf.fromInt i))
409                                  val fmt = CL.mkStr(                                  val fmt = CL.mkStr(
410                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))
411                                        ^ "\n")                                        ^ "\n")
412                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))                                  val args = List.tabulate (d, sel)
413                                  in                                  in
414                                    fmt :: args                                    fmt :: args
415                                  end                                  end
416                              | _ => raise Fail("genStrand: unsupported output type " ^ Ty.toString ty)                              | _ => raise Fail("genStrand: unsupported output type " ^ Ty.toString ty)
417                            (* end case *))                            (* end case *))
   
                           val body = let  
   
                             fun loopParams (3) =  
                                  "x"::"y"::"k"::[]  
                               | loopParams (2) =  
                                  "x"::"y"::[]  
                               | loopParams (1) =  
                                  "x"::[]  
                               | loopParams (_) =  
                                 raise Fail("genStrandPrint: unsupported output type " ^ Ty.toString ty)  
   
                            fun mkLoopNest ([],_) =  
                                                 CL.mkCall("fprintf", CL.mkVar "outS" :: prArgs)  
                                 | mkLoopNest (param::rest,count) = let  
                                         val body = mkLoopNest (rest, count + 1)  
418                                     in                                     in
419                                                  CL.mkFor(                        CL.D_Func(["static"], CL.voidTy, prFnName, params,
420                                                          [(CL.intTy, param, CL.mkInt(0,CL.intTy))],                          CL.mkCall("fprintf", CL.mkVar "outS" :: prArgs))
                                                 CL.mkBinOp(CL.mkVar param, CL.#<, CL.mkSubscript(CL.mkVar "sizes",CL.mkInt(count,CL.intTy))),  
                                                 [CL.mkPostOp(CL.mkVar param, CL.^++)],  
                                                 body)  
                                    end  
                         in  
                                 [mkLoopNest ((loopParams nDims),0)]  
                         end  
   
                     in  
                       CL.D_Func(["static"], CL.voidTy, prFnName, params,CL.mkBlock(body))  
421                      end                      end
422                in                in
423                                   prFn                                   prFn
424                end                end
425          fun genStrandTyDef (Strand{tyName, state,...}) =  
426            fun genStrandTyDef (targetTy, Strand{state,...},tyName) =
427              (* the type declaration for the strand's state struct *)              (* the type declaration for the strand's state struct *)
428                CL.D_StructDef(                CL.D_StructDef(
429                        List.rev (List.map (fn ToCL.V(ty, x) => (ty, x)) (!state)),                  List.rev (List.map (fn x => (targetTy x, #var x)) (!state)),
430                        tyName)                        tyName)
431    
432    
433          (* generates the load kernel function *)           fun genStrandCopy(Strand{tyName,name,state,...}) = let
434  (* FIXME: this code might be part of the runtime system *)                val params = [
435          fun genKernelLoader() =                        CL.PARAM(["__global"], CL.T_Ptr(CL.T_Named tyName), "selfIn"),
436                  CL.D_Verbatim ( ["/* Loads the Kernel from a file */",                        CL.PARAM(["__global"], CL.T_Ptr(CL.T_Named tyName), "selfOut")
437                                                  "char * loadKernel (const char * filename) {",                    ]
438                                                  "struct stat statbuf;",                  val assignStms = List.rev(List.map(fn x => CL.mkAssign(CL.mkIndirect(CL.E_Var "selfOut", #var x),
439                                                  "FILE *fh;",                                                                                                         CL.mkIndirect(CL.E_Var "selfIn", #var x))) (!state))
440                                                  "char *source;",                   in
441                                                  "fh = fopen(filename, \"r\");",                          CL.D_Func([""], CL.voidTy, RN.strandCopy name, params,CL.mkBlock(assignStms))
442                                                  "if (fh == 0)",                   end
443                                                  "   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;",  
                                                 "}"])  
444  (* generates the opencl buffers for the image data *)  (* generates the opencl buffers for the image data *)
445          fun getGlobalDataBuffers(globals,count,contextVar,errVar) = let          fun getGlobalDataBuffers (globals, imgGlobals, contextVar, errVar) = let
446                  val globalBuffErr = "error creating OpenCL global buffer\n"
447                  fun errorFn msg = CL.mkIfThen(CL.mkBinOp(CL.E_Var errVar, CL.#!=, CL.E_Var "CL_SUCCESS"),
448                        CL.mkBlock([CL.mkCall("fprintf",[CL.E_Var "stderr", CL.E_Str msg]),
449                        CL.mkCall("exit",[CL.mkInt 1])]))
450                  val shadowTypeDecl =
451                        CL.mkDecl(CL.T_Named(RN.shadowGlobalsTy), RN.shadowGlaobalsName, NONE)
452                  val globalToShadowStms = List.map (fn (x:mirror_var) => #hToS x ) globals
453                  val globalBufferDecl =  CL.mkDecl(clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)                  val globalBufferDecl =  CL.mkDecl(clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)
454                  val globalBuffer = CL.mkAssign(CL.mkVar(concat[RN.globalsVarName,"_cl"]), CL.mkApply("clCreateBuffer",                val globalBuffer = CL.mkAssign(CL.mkVar(concat[RN.globalsVarName,"_cl"]),
455                                                                  [CL.mkVar contextVar,                      CL.mkApply("clCreateBuffer", [
456                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",                          CL.mkVar contextVar,
457                                                                  CL.mkApply("sizeof",[CL.mkVar RN.globalsTy]),                          CL.mkBinOp(CL.mkVar "CL_MEM_READ_ONLY", CL.#|, CL.mkVar "CL_MEM_COPY_HOST_PTR"),
458                                                                  CL.mkVar RN.globalsVarName,                          CL.mkSizeof(CL.T_Named RN.shadowGlobalsTy),
459                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))                          CL.mkUnOp(CL.%&,CL.mkVar RN.shadowGlaobalsName),
460                            CL.mkUnOp(CL.%&,CL.mkVar errVar)
461                          ]))
462          fun genDataBuffers([],_,_,_) = []          fun genDataBuffers([],_,_,_) = []
463            | genDataBuffers((var,nDims)::globals,count,contextVar,errVar) = let                  | genDataBuffers ((var,nDims)::globals, contextVar, errVar,errFn) = let
464          (* FIXME: use CL constructors to  build expressions (not strings) *)                      val hostVar = CL.mkIndirect(CL.mkVar RN.globalsVarName, var)
465                    val size = if nDims = 1 then                      val size = CL.mkIndirect(hostVar, "dataSzb")
                                         CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,  
                                          CL.mkIndirect(CL.mkVar var, "size[0]"))  
                                         else if nDims = 2 then  
                                         CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,  
                                           CL.mkIndirect(CL.mkVar var, concat["size[0]", " * ", var, "->size[1]"]))  
                                         else  
                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,  
                                           CL.mkIndirect(CL.mkVar var,concat["size[0]", " * ", var, "->size[1] * ", var, "->size[2]"]))  
   
466                   in                   in
                    CL.mkDecl(clMemoryTy, RN.addBufferSuffix var ,NONE)::  
467                     CL.mkDecl(clMemoryTy, RN.addBufferSuffixData var ,NONE)::                     CL.mkDecl(clMemoryTy, RN.addBufferSuffixData var ,NONE)::
468                     CL.mkAssign(CL.mkVar(RN.addBufferSuffix var), CL.mkApply("clCreateBuffer",                        CL.mkAssign(CL.mkVar(RN.addBufferSuffixData var),
469                                                                  [CL.mkVar contextVar,                          CL.mkApply("clCreateBuffer", [
470                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",                              CL.mkVar contextVar,
471                                                                  CL.mkApply("sizeof",[CL.mkVar (RN.imageTy nDims)]),                              CL.mkVar "CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR",
                                                                 CL.mkVar var,  
                                                                 CL.mkUnOp(CL.%&,CL.mkVar errVar)])) ::  
                         CL.mkAssign(CL.mkVar(RN.addBufferSuffixData var), CL.mkApply("clCreateBuffer",  
                                                                 [CL.mkVar contextVar,  
                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",  
472                                                                  size,                                                                  size,
473                                                                  CL.mkIndirect(CL.mkVar var,"data"),                              CL.mkIndirect(hostVar, "data"),
474                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)])):: genDataBuffers(globals,count + 2,contextVar,errVar)                              CL.mkUnOp(CL.%&,CL.mkVar errVar)
475                              ])) ::
476                            errFn(concat["error in creating ",RN.addBufferSuffixData var, " global buffer\n"]) ::
477                            genDataBuffers(globals,contextVar,errVar,errFn)
478                  end                  end
479          in          in
480                  [globalBufferDecl] @ [globalBuffer] @ genDataBuffers(globals,count + 2,contextVar,errVar)                  [shadowTypeDecl] @ globalToShadowStms
481                    @ [globalBufferDecl, globalBuffer,errorFn(globalBuffErr)]
482                    @ genDataBuffers(imgGlobals,contextVar,errVar,errorFn)
483          end          end
484    
   
485  (* generates the kernel arguments for the image data *)  (* generates the kernel arguments for the image data *)
486          fun genGlobalArguments(globals,count,kernelVar,errVar) = let          fun genGlobalArguments(globals,count,kernelVar,errVar) = let
487          val globalArgument = CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,CL.mkApply("clSetKernelArg",                val globalArgErr = "error creating OpenCL global argument\n"
488                  fun errorFn msg = CL.mkIfThen(CL.mkBinOp(CL.E_Var errVar, CL.#!=, CL.E_Var "CL_SUCCESS"),
489                        CL.mkBlock([CL.mkCall("fprintf",[CL.E_Var "stderr", CL.E_Str msg]),
490                        CL.mkCall("exit",[CL.mkInt 1])]))
491                  val globalArgument = CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.&=,
492                        CL.mkApply("clSetKernelArg",
493                                                                  [CL.mkVar kernelVar,                                                                  [CL.mkVar kernelVar,
494                                                                   CL.mkInt(count,CL.intTy),                         CL.mkPostOp(CL.E_Var count, CL.^++),
495                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
496                                                                   CL.mkUnOp(CL.%&,CL.mkVar(concat[RN.globalsVarName,"_cl"]))])))                                                                   CL.mkUnOp(CL.%&,CL.mkVar(concat[RN.globalsVarName,"_cl"]))])))
497                  fun genDataArguments ([],_,_,_,_) = []
498          fun genDataArguments([],_,_,_) = []                  | genDataArguments ((var,nDims)::globals,count,kernelVar,errVar,errFn) =
499            | genDataArguments((var,nDims)::globals,count,kernelVar,errVar) =                      CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.$=,
500                          CL.mkApply("clSetKernelArg",
                 CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=, CL.mkApply("clSetKernelArg",  
501                                                                  [CL.mkVar kernelVar,                                                                  [CL.mkVar kernelVar,
502                                                                   CL.mkInt(count,CL.intTy),                           CL.mkPostOp(CL.E_Var count, CL.^++),
503                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
504                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffix var))])))::                           CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffixData var))]))) ::
505                             errFn(concat["error in creating ",RN.addBufferSuffixData var, " argument\n"]) ::
506                          CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,CL.mkApply("clSetKernelArg",                      genDataArguments (globals,count,kernelVar,errVar,errFn)
                                                                 [CL.mkVar kernelVar,  
                                                                  CL.mkInt((count + 1),CL.intTy),  
                                                                  CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),  
                                                                  CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffixData var))]))):: genDataArguments (globals, count + 2,kernelVar,errVar)  
   
         in  
   
                 [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)  
507                in                in
508                  CL.D_Func([],CL.intTy,"main",params,body)                  globalArgument :: errorFn globalArgErr ::
509                      genDataArguments(globals, count, kernelVar, errVar,errorFn)
510                end                end
511    
512        (* generates the host-side setup function *)        (* generates the globals buffers and arguments function *)
513          fun genHostSetupFunc (strand as Strand{name,tyName,...}, filename, nDims, initially, imgGlobals) = let          fun genGlobalBuffersArgs (globals,imgGlobals) = let
514              (* 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"  
515                val errVar = "err"                val errVar = "err"
516                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")])  
517                val params = [                val params = [
518                        CL.PARAM([],CL.T_Named("cl_device_id"), deviceVar)                        CL.PARAM([],CL.T_Named("cl_context"), "context"),
519                          CL.PARAM([],CL.T_Named("cl_kernel"), "kernel"),
520                          CL.PARAM([],CL.T_Named("cl_command_queue"), "cmdQ"),
521                          CL.PARAM([],CL.T_Named("int"), "argStart")
522                      ]                      ]
523                val declarations = [                val clGlobalBuffers = getGlobalDataBuffers(globals,!imgGlobals, "context", errVar)
524                      CL.mkDecl(clProgramTy, programVar, NONE),                val clGlobalArguments = genGlobalArguments(!imgGlobals, "argStart", "kernel", errVar)
525                      CL.mkDecl(clKernelTy, kernelVar, NONE),              (* Body put all the statments together *)
526                      CL.mkDecl(clCmdQueueTy, cmdVar, NONE),                val body = CL.mkDecl(clIntTy, errVar, SOME(CL.I_Exp(CL.mkInt 0)))
527                      CL.mkDecl(clContextTy, contextVar, NONE),                      :: clGlobalBuffers @ clGlobalArguments
528                      CL.mkDecl(CL.intTy, errVar, NONE),                in
529                      CL.mkDecl(CL.intTy, numStrandsVar, SOME(CL.I_Exp(CL.mkInt(1,CL.intTy)))),                  CL.D_Func([],CL.voidTy,RN.globalsSetupName,params,CL.mkBlock(body))
530                      CL.mkDecl(CL.intTy, stateSizeVar, NONE),                end
                     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])]  
   
   
531    
532                  (* Setup Global and Local variables *)        (* generate the data and global parameters *)
533            fun genKeneralGlobalParams ((name,tyname)::rest) =
534                  val globalAndlocalStms = if nDims = 1 then                globalParam (CL.T_Ptr(CL.voidTy), RN.addBufferSuffixData name) ::
535                          [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),                genKeneralGlobalParams rest
536                                                                     CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),            | genKeneralGlobalParams [] = []
                          CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),  
                                                                   CL.mkVar "16")]  
537    
538          (* generate the main kernel function for the .cl file *)
539            fun genKernelFun (strand, nDims, globals, imgGlobals) = let
540                  val Strand{name, tyName, state, output, code,...} = strand
541                  val fName = RN.kernelFuncName;
542                  val inState = "strand_in"
543                  val outState = "strand_out"
544                  val tempVar = "tmp"
545    
546                  val (workerOffset,localOffset) = if nDims = 1 then
547                            ( CL.mkApply(RN.getGroupId,[CL.mkInt 0]), CL.mkApply(RN.getLocalThreadId,[CL.mkInt 0]) )
548                  else if nDims = 2 then                  else if nDims = 2 then
549                          [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),                       (CL.mkBinOp(CL.mkBinOp(CL.mkApply(RN.getGroupId,[CL.mkInt 0]),CL.#*, CL.mkApply(RN.getNumGroups,[CL.mkInt 1])),CL.#+,CL.mkApply(RN.getGroupId,[CL.mkInt 1])),
550                                                                     CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),                              CL.mkBinOp(CL.mkBinOp(CL.mkApply(RN.getLocalThreadId,[CL.mkInt 0]),CL.#*, CL.mkApply(RN.getLocalSize,[CL.mkInt 1])),CL.#+,CL.mkApply(RN.getLocalThreadId,[CL.mkInt 1])))
                         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")]  
   
551                  else                  else
552                          [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),                            ( CL.mkBinOp(CL.mkBinOp(CL.mkBinOp(
553                                                                     CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),                              CL.mkBinOp(CL.mkApply(RN.getGroupId,[CL.mkInt 0]), CL.#*, CL.mkApply(RN.getNumGroups,[CL.mkInt 1])),CL.#*, CL.mkApply(RN.getLocalSize,[CL.mkInt 2])), CL.#+,
554                          CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1,CL.intTy)),                              CL.mkBinOp(CL.mkApply(RN.getGroupId,[CL.mkInt 1]),CL.#*,CL.mkApply(RN.getLocalSize,[CL.mkInt 1]))),CL.#+,CL.mkApply(RN.getGroupId,[CL.mkInt 2])),
555                                                                     CL.mkSubscript(CL.mkVar "size", CL.mkInt(1,CL.intTy))),                             CL.mkBinOp(CL.mkBinOp(CL.mkBinOp(
556                          CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(2,CL.intTy)),                              CL.mkBinOp(CL.mkApply(RN.getLocalThreadId,[CL.mkInt 0]), CL.#*, CL.mkApply(RN.getLocalSize,[CL.mkInt 1])),CL.#*, CL.mkApply(RN.getLocalSize,[CL.mkInt 2])), CL.#+,
557                                                                     CL.mkSubscript(CL.mkVar "size", CL.mkInt(2,CL.intTy))),                              CL.mkBinOp(CL.mkApply(RN.getLocalThreadId,[CL.mkInt 1]),CL.#*,CL.mkApply(RN.getLocalSize,[CL.mkInt 1]))),CL.#+,CL.mkApply(RN.getLocalThreadId,[CL.mkInt 2])) )
                         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")]  
   
558    
559                val params = [
560                          globalParam(CL.T_Ptr(CL.T_Named tyName), "strands"),
561                          globalParam(CL.T_Ptr(CL.T_Num(RawTypes.RT_UInt8)), "strandStatus"),
562                          globalParam(CL.T_Ptr(CL.intTy), "workerQueue"),
563                          globalParam(CL.T_Ptr(CL.intTy),"numAvail"),
564                          clParam("",CL.intTy,"numStrands"),
565                          clParam("",CL.intTy,"limit")] @
566                          [globalParam(globPtrTy, RN.globalsVarName)] @
567                          genKeneralGlobalParams(!imgGlobals)
568    
569                val index_ids = [
570                              CL.mkDecl(CL.intTy, "workerIndex",
571                                SOME(CL.I_Exp(workerOffset))),
572                              CL.mkDecl(CL.intTy, "strandIndex",
573                                SOME(CL.I_Exp(CL.mkBinOp(CL.mkBinOp(CL.mkSubscript(CL.mkVar "workQueue",CL.mkVar "workerIndex"),CL.#+,localOffset),CL.#*,CL.mkVar "limit"))))
574                            ]
575    
576                  (* Setup Kernel arguments *)              val strandDecl = [
577                  val kernelArguments = [CL.mkAssign(CL.mkVar errVar,CL.mkApply("clSetKernelArg",                        CL.mkDecl(CL.T_Named tyName, "selfIn", NONE),
578                                                                  [CL.mkVar kernelVar,                        CL.mkDecl(CL.T_Named tyName, "selfOut", NONE)
579                                                                   CL.mkInt(0,CL.intTy),                      ]
580                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),                val imageDataDecl = CL.mkDecl(CL.T_Named(RN.imageDataType),RN.globalImageDataName,NONE)
581                                                                   CL.mkUnOp(CL.%&,CL.mkVar clInstateVar)])),                val imageDataStms = List.map (fn (x,_) =>
582                                                              CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar, CL.|=,CL.mkApply("clSetKernelArg",                    CL.mkAssign(CL.mkSelect(CL.mkVar(RN.globalImageDataName),RN.imageDataName x),
583                                                                  [CL.mkVar kernelVar,                                CL.mkVar(RN.addBufferSuffixData x))) (!imgGlobals)
584                                                                   CL.mkInt(1,CL.intTy),  
585                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),              val status = [CL.mkDecl(CL.intTy, "status", NONE)]
                                                                  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])]  
586    
587    
588           fun strandCopy(inStrand, outStrand) = CL.mkCall(RN.strandCopy name,[inStrand,outStrand])
589           val updateStm =  CL.mkAssign(CL.mkVar "status",
590                            CL.mkApply(RN.strandUpdate name,
591                              [CL.mkVar inState,
592                               CL.mkVar outState,
593                               CL.mkVar RN.globalsVarName,
594                               CL.mkVar RN.globalImageDataName]))
595    
596                  (* Body put all the statments together *)                (*      CL.mkBlock([CL.mkAssign(CL.mkSubscript(CL.mkVar "strandStatus",CL.mkVar "strandIndex"),CL.mkVar RN.kStable),
597                  val body =  declarations @ [globalsDecl,initGlobalsCall] (*@ platformStm @ devicesStm *) @ contextStm @ commandStm @ !initially @ [strandSize] @                                  strandCopy(CL.mkUnOp(CL.%&,CL.mkVar "selfOut"),CL.mkBinOp(CL.mkVar "strands", CL.#+, CL.mkVar "strandIndex")),
598                                     strandsArrays @ globalAndlocalStms @ [widthDel,strands_init]  @ clStrandObjects @ clGlobalBuffers @ sourceStms  @ create_build_stms  (*@                                  CL.mkCall(RN.atom_dec,[CL.mkUnOp(CL.%&,CL.mkSubcript(CL.mkVar "numAvail",CL.mkInt 0))])]),
599                                     kernelArguments @ clGlobalArguments @ enqueueStm @  [outputStm] @ freeStms @ outputData *)                      CL.mkBlock([CL.mkIfThen(CL.mkBinOp(CL.E_Var "status", CL.#==, CL.E_Var RN.kDie),
600                                    CL.mkBlock([CL.mkAssign(CL.mkSubscript(CL.mkVar "strandStatus",CL.mkVar "strandIndex"),CL.mkVar RN.kDie),
601                                    strandCopy(CL.mkUnOp(CL.%&,CL.mkVar "selfOut"),CL.mkBinOp(CL.mkVar "strands", CL.#+, CL.mkVar "strandIndex")),
602                                    CL.mkCall(RN.atom_dec,[CL.mkUnOp(CL.%&,CL.mkSubscript(CL.mkVar "numAvail",CL.mkInt 0))])])) *)
603    
604           val statusIf = CL.mkIfThenElse(CL.mkBinOp(CL.E_Var "status", CL.#==, CL.E_Var RN.kStabilize),
605                        CL.mkBlock([CL.mkAssign(CL.mkSubscript(CL.mkVar "strandStatus",CL.mkVar "strandIndex"),CL.mkVar RN.kStable),
606                                    strandCopy(CL.mkUnOp(CL.%&,CL.mkVar "selfOut"),CL.mkBinOp(CL.mkVar "strands", CL.#+, CL.mkVar "strandIndex")),
607                                    CL.mkCall(RN.atom_dec,[CL.mkUnOp(CL.%&,CL.mkSubscript(CL.mkVar "numAvail",CL.mkInt 0))])]),
608                        CL.mkBlock([CL.mkIfThen(CL.mkBinOp(CL.E_Var "status", CL.#==, CL.E_Var RN.kDie),
609                                    CL.mkBlock([CL.mkAssign(CL.mkSubscript(CL.mkVar "strandStatus",CL.mkVar "strandIndex"),CL.mkVar RN.kDie),
610                                    strandCopy(CL.mkUnOp(CL.%&,CL.mkVar "selfOut"),CL.mkBinOp(CL.mkVar "strands", CL.#+, CL.mkVar "strandIndex")),
611                                    CL.mkCall(RN.atom_dec,[CL.mkUnOp(CL.%&,CL.mkSubscript(CL.mkVar "numAvail",CL.mkInt 0))])]))]))
612    
613           val incStrand = CL.mkExpStm(CL.mkPostOp(CL.mkVar "strandIndex",CL.^++))
614    
615            val forStablize = CL.mkFor( [(CL.intTy,"idx",CL.mkInt 0)], CL.mkBinOp(CL.mkBinOp(CL.mkVar "idx", CL.#<, CL.mkVar "limit"),CL.#&&,
616                                                                         CL.mkBinOp(CL.mkVar "strandIndex", CL.#<, CL.mkVar "numStrands")),
617                                           [CL.mkPostOp(CL.mkVar "idx", CL.^++)], CL.mkBlock(
618                                           [
619                                             strandCopy(CL.mkUnOp(CL.%&,CL.mkVar "selfIn"),CL.mkBinOp(CL.mkVar "strands", CL.#+, CL.mkVar "strandIndex")),
620                                             updateStm,
621                                             statusIf,
622                                             incStrand
623                                           ]))
624    
625                  val local_vars = index_ids
626                        @ [imageDataDecl]
627                        @ imageDataStms
628                        @ strandDecl
629                        @ status
630    
631                  val body = CL.mkBlock(local_vars @ [forStablize])
632                  in                  in
633                    CL.D_Func(["__kernel"], CL.voidTy, fName, params, body)
                 CL.D_Func([],CL.voidTy,RN.setupFName,params,CL.mkBlock(body))  
   
634                  end                  end
635  (* generate the data and global parameters *)  
636          fun genKeneralGlobalParams ((name,tyname)::rest) =        (* generate a global structure type definition from the list of globals *)
637                  CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), concat[RN.globalsVarName]) ::          fun genGlobalStruct (targetTy, globals, tyName) = let
638                  CL.PARAM([], CL.T_Ptr(CL.T_Named (RN.imageTy tyname)),RN.addBufferSuffix name) ::                val globs = List.map (fn (x : mirror_var) => (targetTy x, #var x)) globals
                 CL.PARAM([], CL.T_Ptr(CL.voidTy),RN.addBufferSuffixData name) ::  
                 genKeneralGlobalParams(rest)  
   
           | genKeneralGlobalParams ([]) = []  
   
         (*generate code for intilizing kernel global data *)  
         fun initKernelGlobals (globals,imgGlobals) = let  
                 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([]) = []  
   
                 fun initGlobalImages((name,tyname)::rest) =  
                                 CL.mkAssign(CL.mkVar name, CL.mkVar (RN.addBufferSuffix name)) ::  
                                 CL.mkAssign(CL.mkIndirect(CL.mkVar name,"data"),CL.mkVar (RN.addBufferSuffixData name)) ::  
                                 initGlobalImages(rest)  
                   | initGlobalImages([]) = []  
639                  in                  in
640                    initGlobalStruct(globals) @ initGlobalImages(imgGlobals)                  CL.D_StructDef(globs, tyName)
641                  end                  end
642    
643          (* generate the main kernel function for the .cl file *)        (* generate a global structure type definition from the image data of the image globals *)
644          fun genKernelFun(Strand{name, tyName, state, output, code,...},nDims,globals,imgGlobals) = let          fun genImageDataStruct (imgGlobals, tyName) = let
645                   val fName = RN.kernelFuncName;                val globs = List.map
646                   val inState = "strand_in"                      (fn (x, _) => (globalPtr CL.voidTy, RN.imageDataName x))
647                   val outState = "strand_out"                        imgGlobals
648               val params = [                in
649                        CL.PARAM(["__global"], CL.T_Ptr(CL.T_Named tyName), "selfIn"),                  CL.D_StructDef(globs, tyName)
                       CL.PARAM(["__global"], CL.T_Ptr(CL.T_Named tyName), "selfOut"),  
                       CL.PARAM(["__global"], CL.intTy, "width")  
                     ] @ genKeneralGlobalParams(!imgGlobals)  
                   val thread_ids = if nDims = 1  
                         then [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),  
                                   CL.mkAssign(CL.mkVar "x",CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(0,CL.intTy)]))]  
                         else  
                                 [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),  
                                  CL.mkDecl(CL.intTy, "y", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),  
                                   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)]  
650                                  end                                  end
651    
652            fun genGlobals (declFn, targetTy, globals) = let
653                  fun doVar (x : mirror_var) = declFn (CL.D_Var([], targetTy x, #var x, NONE))
654                  in
655                    List.app doVar globals
656                  end
657    
658                    val status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy))))          fun genStrandDesc (Strand{name, output, ...}) = let
659                    val local_vars =  thread_ids @ initKernelGlobals(!globals,!imgGlobals)  @ strandDecl @ strandObjects @ [status]              (* the strand's descriptor object *)
660                    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 descI = let
661                    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)])),                      fun fnPtr (ty, f) = CL.I_Exp(CL.mkCast(CL.T_Named ty, CL.mkVar f))
662                                                          CL.mkCall(RN.strandStabilize name,[ CL.mkUnOp(CL.%&,CL.mkVar inState),  CL.mkUnOp(CL.%&,CL.mkVar outState)])]                      val SOME(outTy, _) = !output
   
                   val whileBlock = [CL.mkWhile(while_exp,CL.mkBlock while_body)]  
   
                   val body = CL.mkBlock(local_vars  @ whileBlock)  
663                  in                  in
664                     CL.D_Func(["__kernel"], CL.voidTy, fName, params, body)                        CL.I_Struct[
665                              ("name", CL.I_Exp(CL.mkStr name)),
666                              ("stateSzb", CL.I_Exp(CL.mkSizeof(CL.T_Named(RN.strandTy name)))),
667    (*
668                              ("outputSzb", CL.I_Exp(CL.mkSizeof(ToC.trTy outTy))),
669    *)
670                              ("update", fnPtr("update_method_t", "0")),
671                              ("print", fnPtr("print_method_t", name ^ "Print"))
672                            ]
673                  end                  end
674          (* generate a global structure from the globals *)                val desc = CL.D_Var([], CL.T_Named N.strandDescTy, N.strandDesc name, SOME descI)
         fun genGlobalStruct(globals) = let  
                  fun getGlobals(CL.D_Var(_,ty,globalVar,_)::rest) = (ty,globalVar)::getGlobals(rest)  
                    | getGlobals([]) = []  
                    | getGlobals(_::rest) = getGlobals(rest)  
675                   in                   in
676                          CL.D_StructDef(getGlobals(globals),RN.globalsTy)                  desc
677                    end                    end
678    
679        (* generate the table of strand descriptors *)        (* generate the table of strand descriptors *)
680          fun genStrandTable (ppStrm, strands) = let          fun genStrandTable (declFn, strands) = let
681                val nStrands = length strands                val nStrands = length strands
682                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)))
683                fun genInits (_, []) = []                fun genInits (_, []) = []
684                  | 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)  
685                in                in
686                  ppDecl (CL.D_Var([], CL.int32, RN.numStrands,                  declFn (CL.D_Var([], CL.int32, N.numStrands,
687                    SOME(CL.I_Exp(CL.mkInt(IntInf.fromInt nStrands, CL.int32)))));                    SOME(CL.I_Exp(CL.E_Int(IntInf.fromInt nStrands, CL.int32)))));
688                  ppDecl (CL.D_Var([],                  declFn (CL.D_Var([],
689                    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),
690                    RN.strands,                    N.strands,
691                    SOME(CL.I_Array(genInits (0, strands)))))                    SOME(CL.I_Array(genInits (0, strands)))))
692                end                end
693    
694            fun genSrc (baseName, prog) = let
695          fun genSrc (baseName, Prog{double,globals, topDecls, strands, initially,imgGlobals,numDims,...}) = let                val Prog{name,double, globals, topDecls, strands, initially, imgGlobals, numDims, ...} = prog
696                val clFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "cl"}                val clFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "cl"}
697                val cFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "c"}                val cFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "c"}
698                val clOutS = TextIO.openOut clFileName                val clOutS = TextIO.openOut clFileName
699                val cOutS = TextIO.openOut cFileName                val cOutS = TextIO.openOut cFileName
700  (* FIXME: need to use PrintAsC and PrintAsCL *)                val clppStrm = PrintAsCL.new clOutS
               val clppStrm = PrintAsC.new clOutS  
701                val cppStrm = PrintAsC.new cOutS                val cppStrm = PrintAsC.new cOutS
702                  val progName = name
703                fun cppDecl dcl = PrintAsC.output(cppStrm, dcl)                fun cppDecl dcl = PrintAsC.output(cppStrm, dcl)
704                fun clppDecl dcl = PrintAsCL.output(clppStrm, dcl)                fun clppDecl dcl = PrintAsCL.output(clppStrm, dcl)
705                val strands = AtomTable.listItems strands                val strands = AtomTable.listItems strands
# Line 940  Line 711 
711                        then "#define DIDEROT_DOUBLE_PRECISION"                        then "#define DIDEROT_DOUBLE_PRECISION"
712                        else "#define DIDEROT_SINGLE_PRECISION",                        else "#define DIDEROT_SINGLE_PRECISION",
713                      "#define DIDEROT_TARGET_CL",                      "#define DIDEROT_TARGET_CL",
714                      "#include \"Diderot/cl-types.h\""                      "#include \"Diderot/cl-diderot.h\""
715                    ]));                    ]));
716                  List.app clppDecl (List.rev (!globals));                  clppDecl (genGlobalStruct (#gpuTy, !globals, RN.globalsTy));
717                  clppDecl (genGlobalStruct (!globals));                  clppDecl (genImageDataStruct(!imgGlobals,RN.imageDataType));
718                  clppDecl (genStrandTyDef strand);                  clppDecl (genStrandTyDef(#gpuTy, strand,tyName));
719                    clppDecl (genStrandCopy(strand));
720                  List.app clppDecl (!code);                  List.app clppDecl (!code);
721                  clppDecl (genKernelFun (strand,!numDims,globals,imgGlobals));                  clppDecl (genKernelFun (strand,!numDims,globals,imgGlobals));
722                (* Generate the Host file .c *)                (* Generate the Host C file *)
723                  cppDecl (CL.D_Verbatim([                  cppDecl (CL.D_Verbatim([
724                      if double                      if double
725                        then "#define DIDEROT_DOUBLE_PRECISION"                        then "#define DIDEROT_DOUBLE_PRECISION"
# Line 955  Line 727 
727                      "#define DIDEROT_TARGET_CL",                      "#define DIDEROT_TARGET_CL",
728                      "#include \"Diderot/diderot.h\""                      "#include \"Diderot/diderot.h\""
729                    ]));                    ]));
730                  List.app cppDecl (List.rev (!globals));                  cppDecl (CL.D_Var(["static"], CL.charPtr, "ProgramName",
731                  cppDecl (genGlobalStruct (!globals));                    SOME(CL.I_Exp(CL.mkStr progName))));
732                  cppDecl (genStrandTyDef strand);                  cppDecl (genGlobalStruct (#hostTy, !globals, RN.globalsTy));
733                    cppDecl (genGlobalStruct (#shadowTy, !globals, RN.shadowGlobalsTy));
734    (* FIXME: does this really need to be a global? *)
735                    cppDecl (CL.D_Var(["static"], globPtrTy, RN.globalsVarName, NONE));
736                    cppDecl (genStrandTyDef (#shadowTy, strand, tyName));
737                  cppDecl  (!init_code);                  cppDecl  (!init_code);
738                  cppDecl (genStrandInit(strand,!numDims));                  cppDecl (genStrandPrint strand);
                 cppDecl (genStrandPrint(strand,!numDims));  
                 (* cppDecl (genKernelLoader());*)  
739                  List.app cppDecl (List.rev (!topDecls));                  List.app cppDecl (List.rev (!topDecls));
740                  cppDecl (genHostSetupFunc (strand, clFileName, !numDims, initially, imgGlobals));                  cppDecl (genGlobalBuffersArgs (!globals,imgGlobals));
741                    List.app (fn strand => cppDecl (genStrandDesc strand)) strands;
742                    genStrandTable (cppDecl, strands);
743                    cppDecl (!initially);
744                  PrintAsC.close cppStrm;                  PrintAsC.close cppStrm;
745                  PrintAsCL.close clppStrm;                  PrintAsCL.close clppStrm;
746                  TextIO.closeOut cOutS;                  TextIO.closeOut cOutS;
747                  TextIO.closeOut clOutS                  TextIO.closeOut clOutS
748                end                end
749    
750        (* 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.  
        *)  
751          fun generate (basename, prog as Prog{double, parallel, debug, ...}) = let          fun generate (basename, prog as Prog{double, parallel, debug, ...}) = let
752                fun condCons (true, x, xs) = x::xs                fun condCons (true, x, xs) = x::xs
753                  | condCons (false, _, xs) = xs                  | condCons (false, _, xs) = xs
# Line 998  Line 773 
773                  RunCC.link (basename, ldOpts)                  RunCC.link (basename, ldOpts)
774                end                end
775    
776        end        end (* Program *)
777    
778    (* strands *)    (* strands *)
779      structure Strand =      structure Strand =
# Line 1027  Line 802 
802          fun init (Strand{name, tyName, code,init_code, ...}, params, init) = let          fun init (Strand{name, tyName, code,init_code, ...}, params, init) = let
803                val fName = RN.strandInit name                val fName = RN.strandInit name
804                val params =                val params =
805                      CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfOut") ::                      globalParam (globPtrTy, RN.globalsVarName) ::
806                        globalParam (CL.T_Ptr(CL.T_Named tyName), "selfOut") ::
807                        List.map (fn (ToCL.V(ty, x)) => CL.PARAM([], ty, x)) params                        List.map (fn (ToCL.V(ty, x)) => CL.PARAM([], ty, x)) params
808                val initFn = CL.D_Func([], CL.voidTy, fName, params, init)                val initFn = CL.D_Func([], CL.voidTy, fName, params, init)
809                in                in
# Line 1036  Line 812 
812    
813        (* register a strand method *)        (* register a strand method *)
814          fun method (Strand{name, tyName, code,...}, methName, body) = let          fun method (Strand{name, tyName, code,...}, methName, body) = let
815                val fName = concat[name, "_", methName]                val fName = concat[name, "_", MethodName.toString methName]
816                val params = [                val params = [
817                        CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfIn"),                        globalParam (CL.T_Ptr(CL.T_Named tyName), "selfIn"),
818                        CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfOut")                        globalParam (CL.T_Ptr(CL.T_Named tyName), "selfOut"),
819                          globalParam (CL.T_Ptr(CL.T_Named (RN.globalsTy)), RN.globalsVarName),
820                          CL.PARAM([],CL.T_Named(RN.imageDataType),RN.globalImageDataName)
821                      ]                      ]
822                val methFn = CL.D_Func([], CL.int32, fName, params, body)                val resTy = (case methName
823                         of MethodName.Update => CL.T_Named "StrandStatus_t"
824                          | MethodName.Stabilize => CL.voidTy
825                        (* end case *))
826                  val methFn = CL.D_Func([], resTy, fName, params, body)
827                in                in
828                  code := methFn :: !code                  code := methFn :: !code
829                end                end

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

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