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 1285, Tue Jun 7 10:33:17 2011 UTC revision 1460, Sun Aug 7 20:31:00 2011 UTC
# Line 1  Line 1 
1  (* c-target.sml  (* cl-target.sml
2   *   *
3   * COPYRIGHT (c) 2011 The Diderot Project (http://diderot-language.cs.uchicago.edu)   * COPYRIGHT (c) 2011 The Diderot Project (http://diderot-language.cs.uchicago.edu)
4   * All rights reserved.   * All rights reserved.
# Line 13  Line 13 
13      structure CL = CLang      structure CL = CLang
14      structure RN = RuntimeNames      structure RN = RuntimeNames
15      structure ToCL = TreeToCL      structure ToCL = TreeToCL
16        structure N = CNames
17    
18      (* 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 24  Line 71 
71                (* end case *))                (* end case *))
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 "diderotGlobals", 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 99  Line 163 
163      structure Tr =      structure Tr =
164        struct        struct
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                  | _ => ToCL.trBlock (vMap, fn (_, _, stm) => [stm], blk)                   of StrandScope stateVars =>
188                (* end case *))                        ToC.trBlock (vMap, saveState "StrandScope" stateVars ToC.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)
192                      | _ => ToC.trBlock (vMap, fn (_, _, stm) => [stm], blk)
193                    (* 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 156  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 170  Line 254 
254        struct        struct
255          fun new {name, double, parallel, debug} = (          fun new {name, double, parallel, debug} = (
256                RN.initTargetSpec double;                RN.initTargetSpec double;
257                  CNames.initTargetSpec double;
258                Prog{                Prog{
259                    name = name,                    name = name,
260                    double = double, parallel = parallel, debug = debug,                    double = double, parallel = parallel, debug = debug,
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 202  Line 277 
277                  topDecls := inputsFn :: !topDecls                  topDecls := inputsFn :: !topDecls
278                end                end
279    
280          fun init (Prog{globals, topDecls,...}, CL.S_Block(init)) = let        (* register the global initialization part of a program *)
281                val params = [          fun init (Prog{topDecls, ...}, init) = let
282                        CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName)                val globalsDecl = CL.mkAssign(CL.E_Var RN.globalsVarName,
283                      ]                      CL.mkApply("malloc", [CL.mkSizeof(CL.T_Named RN.globalsTy)]))
284                val body = CL.S_Block(globalIndirects(!globals,init))                val initFn = CL.D_Func(
285                val initFn = CL.D_Func([], CL.voidTy, RN.initGlobals, params, body)                      [], CL.voidTy, RN.initGlobals, [],
286                in                      CL.mkBlock[
287                  topDecls := initFn :: !topDecls                          globalsDecl,
288                end                          CL.mkCall(RN.initGlobalsHelper, [CL.mkVar RN.globalsVarName])
289            | init (Prog{globals,topDecls,...}, init) = let                        ])
290                val params = [                val initHelperFn = CL.D_Func(
291                        CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName)                      [], CL.voidTy, RN.initGlobalsHelper,
292                      ]                      [CL.PARAM([], globPtrTy, RN.globalsVarName)],
293                val initFn = CL.D_Func([], CL.voidTy, RN.initGlobals, params, init)                      init)
294                  val shutdownFn = CL.D_Func(
295                        [], CL.voidTy, RN.shutdown,
296                        [CL.PARAM([], CL.T_Ptr(CL.T_Named RN.worldTy), "wrld")],
297                        CL.S_Block[])
298                in                in
299                  topDecls := 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 232  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 240  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                                  val fmt = CL.mkStr(                                  val fmt = CL.E_Str(
397                                        String.concatWith " " (List.tabulate(d, fn _ => !RN.gIntFormat))                                        String.concatWith " " (List.tabulate(d, fn _ => !N.gIntFormat))
398                                        ^ "\n")                                        ^ "\n")
399                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))                                  val args = List.tabulate (d, fn i => ToC.ivecIndex(outState, d, i))
400                                  in                                  in
401                                    fmt :: args                                    fmt :: args
402                                  end                                  end
403                              | Ty.TensorTy[] => [CL.mkStr "%f\n", outState]                              | Ty.TensorTy[] => [CL.E_Str "%f\n", outState]
404                              | Ty.TensorTy[d] => let                              | Ty.TensorTy[d] => let
405                                  val fmt = CL.mkStr(                                  val fmt = CL.E_Str(
406                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))
407                                        ^ "\n")                                        ^ "\n")
408                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))                                  val args = List.tabulate (d, fn i => ToC.vecIndex(outState, d, i))
409                                  in                                  in
410                                    fmt :: args                                    fmt :: args
411                                  end                                  end
412                              | _ => raise Fail("genStrand: unsupported output type " ^ Ty.toString ty)                              | _ => raise Fail("genStrand: unsupported output type " ^ Ty.toString ty)
413                            (* 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)  
414                                     in                                     in
415                                                  CL.mkFor(                        CL.D_Func(["static"], CL.voidTy, prFnName, params,
416                                                          [(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))  
417                      end                      end
418                in                in
419                                   prFn                                   prFn
420                end                end
421          fun genStrandTyDef (Strand{tyName, state,...}) =  
422            fun genStrandTyDef (targetTy, Strand{state,...},tyName) =
423              (* the type declaration for the strand's state struct *)              (* the type declaration for the strand's state struct *)
424                CL.D_StructDef(                CL.D_StructDef(
425                        List.rev (List.map (fn ToCL.V(ty, x) => (ty, x)) (!state)),                  List.rev (List.map (fn x => (targetTy x, #var x)) (!state)),
426                        tyName)                        tyName)
427    
428    
429          (* generates the load kernel function *)           fun genStrandCopy(Strand{tyName,name,state,...}) = let
430  (* FIXME: this code might be part of the runtime system *)                val params = [
431          fun genKernelLoader() =                        CL.PARAM([""], CL.T_Ptr(CL.T_Named tyName), "selfIn"),
432                  CL.D_Verbatim ( ["/* Loads the Kernel from a file */",                        CL.PARAM([""], CL.T_Ptr(CL.T_Named tyName), "selfOut")
433                                                  "char * loadKernel (const char * filename) {",                    ]
434                                                  "struct stat statbuf;",                  val assignStms = List.rev(List.map(fn x => CL.mkAssign(CL.mkIndirect(CL.E_Var "selfOut", #var x),
435                                                  "FILE *fh;",                                                                                                         CL.mkIndirect(CL.E_Var "selfIn", #var x))) (!state))
436                                                  "char *source;",                   in
437                                                  "fh = fopen(filename, \"r\");",                          CL.D_Func([""], CL.voidTy, RN.strandCopy name, params,CL.mkBlock(assignStms))
438                                                  "if (fh == 0)",                   end
439                                                  "   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;",  
                                                 "}"])  
440  (* generates the opencl buffers for the image data *)  (* generates the opencl buffers for the image data *)
441          fun getGlobalDataBuffers(globals,count,contextVar,errVar) = let          fun getGlobalDataBuffers (globals, imgGlobals, contextVar, errVar) = let
442                  val globalBuffErr = "error creating OpenCL global buffer\n"
443                  fun errorFn msg = CL.mkIfThen(CL.mkBinOp(CL.E_Var errVar, CL.#!=, CL.E_Var "CL_SUCCESS"),
444                        CL.mkBlock([CL.mkCall("fprintf",[CL.E_Var "stderr", CL.E_Str msg]),
445                        CL.mkCall("exit",[CL.mkInt 1])]))
446                  val shadowTypeDecl =
447                        CL.mkDecl(CL.T_Named(RN.shadowGlobalsTy), RN.shadowGlaobalsName, NONE)
448                  val globalToShadowStms = List.map (fn (x:mirror_var) => #hToS x ) globals
449                  val globalBufferDecl =  CL.mkDecl(clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)                  val globalBufferDecl =  CL.mkDecl(clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)
450                  val globalBuffer = CL.mkAssign(CL.mkVar(concat[RN.globalsVarName,"_cl"]), CL.mkApply("clCreateBuffer",                val globalBuffer = CL.mkAssign(CL.mkVar(concat[RN.globalsVarName,"_cl"]),
451                                                                  [CL.mkVar contextVar,                      CL.mkApply("clCreateBuffer", [
452                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",                          CL.mkVar contextVar,
453                                                                  CL.mkApply("sizeof",[CL.mkVar RN.globalsTy]),                          CL.mkBinOp(CL.mkVar "CL_MEM_READ_ONLY", CL.#|, CL.mkVar "CL_MEM_COPY_HOST_PTR"),
454                                                                  CL.mkVar RN.globalsVarName,                          CL.mkSizeof(CL.T_Named RN.shadowGlobalsTy),
455                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))                          CL.mkUnOp(CL.%&,CL.mkVar RN.shadowGlaobalsName),
456                            CL.mkUnOp(CL.%&,CL.mkVar errVar)
457                          ]))
458          fun genDataBuffers([],_,_,_) = []          fun genDataBuffers([],_,_,_) = []
459            | genDataBuffers((var,nDims)::globals,count,contextVar,errVar) = let                  | genDataBuffers ((var,nDims)::globals, contextVar, errVar,errFn) = let
460          (* FIXME: use CL constructors to  build expressions (not strings) *)                      val hostVar = CL.mkIndirect(CL.mkVar RN.globalsVarName, var)
461                    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]"]))  
   
462                   in                   in
                    CL.mkDecl(clMemoryTy, RN.addBufferSuffix var ,NONE)::  
463                     CL.mkDecl(clMemoryTy, RN.addBufferSuffixData var ,NONE)::                     CL.mkDecl(clMemoryTy, RN.addBufferSuffixData var ,NONE)::
464                     CL.mkAssign(CL.mkVar(RN.addBufferSuffix var), CL.mkApply("clCreateBuffer",                        CL.mkAssign(CL.mkVar(RN.addBufferSuffixData var),
465                                                                  [CL.mkVar contextVar,                          CL.mkApply("clCreateBuffer", [
466                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",                              CL.mkVar contextVar,
467                                                                  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",  
468                                                                  size,                                                                  size,
469                                                                  CL.mkIndirect(CL.mkVar var,"data"),                              CL.mkIndirect(hostVar, "data"),
470                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)])):: genDataBuffers(globals,count + 2,contextVar,errVar)                              CL.mkUnOp(CL.%&,CL.mkVar errVar)
471                              ])) ::
472                            errFn(concat["error in creating ",RN.addBufferSuffixData var, " global buffer\n"]) ::
473                            genDataBuffers(globals,contextVar,errVar,errFn)
474                  end                  end
475          in          in
476                  [globalBufferDecl] @ [globalBuffer] @ genDataBuffers(globals,count + 2,contextVar,errVar)                  [shadowTypeDecl] @ globalToShadowStms
477                    @ [globalBufferDecl, globalBuffer,errorFn(globalBuffErr)]
478                    @ genDataBuffers(imgGlobals,contextVar,errVar,errorFn)
479          end          end
480    
   
481  (* generates the kernel arguments for the image data *)  (* generates the kernel arguments for the image data *)
482          fun genGlobalArguments(globals,count,kernelVar,errVar) = let          fun genGlobalArguments(globals,count,kernelVar,errVar) = let
483          val globalArgument = CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,CL.mkApply("clSetKernelArg",                val globalArgErr = "error creating OpenCL global argument\n"
484                  fun errorFn msg = CL.mkIfThen(CL.mkBinOp(CL.E_Var errVar, CL.#!=, CL.E_Var "CL_SUCCESS"),
485                        CL.mkBlock([CL.mkCall("fprintf",[CL.E_Var "stderr", CL.E_Str msg]),
486                        CL.mkCall("exit",[CL.mkInt 1])]))
487                  val globalArgument = CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.&=,
488                        CL.mkApply("clSetKernelArg",
489                                                                  [CL.mkVar kernelVar,                                                                  [CL.mkVar kernelVar,
490                                                                   CL.mkInt(count,CL.intTy),                         CL.mkPostOp(CL.E_Var count, CL.^++),
491                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
492                                                                   CL.mkUnOp(CL.%&,CL.mkVar(concat[RN.globalsVarName,"_cl"]))])))                                                                   CL.mkUnOp(CL.%&,CL.mkVar(concat[RN.globalsVarName,"_cl"]))])))
493                  fun genDataArguments ([],_,_,_,_) = []
494          fun genDataArguments([],_,_,_) = []                  | genDataArguments ((var,nDims)::globals,count,kernelVar,errVar,errFn) =
495            | genDataArguments((var,nDims)::globals,count,kernelVar,errVar) =                      CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.$=,
496                          CL.mkApply("clSetKernelArg",
                 CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=, CL.mkApply("clSetKernelArg",  
497                                                                  [CL.mkVar kernelVar,                                                                  [CL.mkVar kernelVar,
498                                                                   CL.mkInt(count,CL.intTy),                           CL.mkPostOp(CL.E_Var count, CL.^++),
499                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
500                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffix var))])))::                           CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffixData var))]))) ::
501                             errFn(concat["error in creating ",RN.addBufferSuffixData var, " argument\n"]) ::
502                          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)  
   
503          in          in
504                    globalArgument :: errorFn globalArgErr ::
505                  [globalArgument] @ genDataArguments(globals,count + 1,kernelVar,errVar)                    genDataArguments(globals, count, kernelVar, errVar,errorFn)
   
506          end          end
507    
508          (* generates the main function of host code *)        (* generates the globals buffers and arguments function *)
509          fun genHostMain() = let          fun genGlobalBuffersArgs (globals,imgGlobals) = let
               val setupCall = [CL.mkCall(RN.setupFName,[CL.mkVar RN.globalsVarName])]  
               val globalsDecl = CL.mkDecl(  
                     CL.T_Ptr(CL.T_Named RN.globalsTy),  
                     RN.globalsVarName,  
                     SOME(CL.I_Exp(CL.mkApply("malloc", [CL.mkApply("sizeof",[CL.mkVar RN.globalsTy])]))))  
               val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.mkVar RN.globalsVarName])  
               val returnStm = [CL.mkReturn(SOME(CL.mkInt(0,CL.intTy)))]  
               val params = [  
                      CL.PARAM([],CL.intTy, "argc"),  
                      CL.PARAM([],CL.charArrayPtr,"argv")  
                    ]  
               val body = CL.mkBlock([globalsDecl] @ [initGlobalsCall]  @ setupCall @ returnStm)  
               in  
                 CL.D_Func([],CL.intTy,"main",params,body)  
               end  
   
       (* generates the host-side setup function *)  
         fun genHostSetupFunc (strand as Strand{name,tyName,...}, filename, nDims, initially, imgGlobals) = let  
510              (* 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"  
511                val errVar = "err"                val errVar = "err"
512                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")])  
513                val params = [                val params = [
514                        CL.PARAM([],CL.T_Named("cl_device_id"), deviceVar)                        CL.PARAM([],CL.T_Named("cl_context"), "context"),
515                          CL.PARAM([],CL.T_Named("cl_kernel"), "kernel"),
516                          CL.PARAM([],CL.T_Named("cl_command_queue"), "cmdQ"),
517                          CL.PARAM([],CL.T_Named("int"), "argStart")
518                      ]                      ]
519                val declarations = [                val clGlobalBuffers = getGlobalDataBuffers(globals,!imgGlobals, "context", errVar)
520                      CL.mkDecl(clProgramTy, programVar, NONE),                val clGlobalArguments = genGlobalArguments(!imgGlobals, "argStart", "kernel", errVar)
521                      CL.mkDecl(clKernelTy, kernelVar, NONE),              (* Body put all the statments together *)
522                      CL.mkDecl(clCmdQueueTy, cmdVar, NONE),                val body = CL.mkDecl(clIntTy, errVar, SOME(CL.I_Exp(CL.mkInt 0)))
523                      CL.mkDecl(clContextTy, contextVar, NONE),                      :: clGlobalBuffers @ clGlobalArguments
524                      CL.mkDecl(CL.intTy, errVar, NONE),                in
525                      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))
526                      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  
                       ])  
527    
528              val clGlobalBuffers = getGlobalDataBuffers(!imgGlobals,3,contextVar,errVar)        (* generate the data and global parameters *)
529            fun genKeneralGlobalParams ((name,tyname)::rest) =
530                  globalParam (CL.T_Ptr(CL.voidTy), RN.addBufferSuffixData name) ::
531                  genKeneralGlobalParams rest
532              | genKeneralGlobalParams [] = []
533    
534          (* generate the main kernel function for the .cl file *)
535            fun genKernelFun (strand, nDims, globals, imgGlobals) = let
536                  val Strand{name, tyName, state, output, code,...} = strand
537                  val fName = RN.kernelFuncName;
538                  val inState = "selfIn"
539                  val outState = "selfOut"
540                  val tempVar = "tmp"
541    
542                  (* Load the Kernel and Header Files *)                val (workerOffset,localOffset) = if nDims = 1 then
543                  val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(1,CL.intTy)),                          ( CL.mkApply(RN.getGroupId,[CL.mkInt 0]), CL.mkApply(RN.getLocalThreadId,[CL.mkInt 0]) )
                                                                           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])]  
544                          else if nDims = 2  then                          else if nDims = 2  then
545                           [CL.mkAssign(CL.mkVar errVar,                       (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])),
546                                                          CL.mkApply("clEnqueueNDRangeKernel",                              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.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])]  
547                          else                          else
548                            [CL.mkAssign(CL.mkVar errVar,                            ( CL.mkBinOp(CL.mkBinOp(CL.mkBinOp(
549                                                          CL.mkApply("clEnqueueNDRangeKernel",                              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.#+,
550                                                                                                  [CL.mkVar cmdVar,                              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])),
551                                                                                                   CL.mkVar kernelVar,                             CL.mkBinOp(CL.mkBinOp(CL.mkBinOp(
552                                                                                                   CL.mkInt(3,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.#+,
553                                                                                                   CL.mkVar "NULL",                              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.mkVar globalVar,  
                                                                                                  CL.mkVar localVar,  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]  
   
   
   
                 (* Setup Global and Local variables *)  
   
                 val globalAndlocalStms = if nDims = 1 then  
                         [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),  
                          CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),  
                                                                   CL.mkVar "16")]  
   
   
                 else if nDims = 2 then  
                         [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(1,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),  
                                                                   CL.mkVar "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(1,CL.intTy)),  
                                                                   CL.mkVar "16")]  
   
                 else  
                         [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(1,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(2,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(2,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),  
                                                                   CL.mkVar "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(1,CL.intTy)),  
                                                                   CL.mkVar "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(2,CL.intTy)),  
                                                                   CL.mkVar "16")]  
   
554    
555                val params = [
556                          globalParam(CL.T_Ptr(CL.T_Named tyName), "strands"),
557                          globalParam(CL.T_Ptr(CL.intTy), "strandStatus"),
558                          globalParam(CL.T_Ptr(CL.intTy), "workerQueue"),
559                          globalParam(CL.T_Ptr(CL.intTy),"numAvail"),
560                          clParam("",CL.intTy,"numStrands"),
561                          clParam("",CL.intTy,"limit")] @
562                          [globalParam(globPtrTy, RN.globalsVarName)] @
563                          genKeneralGlobalParams(!imgGlobals)
564    
565                val index_ids = [
566                              CL.mkDecl(CL.intTy, "workerIndex",
567                                SOME(CL.I_Exp(workerOffset))),
568                              CL.mkDecl(CL.intTy, "strandIndex",
569                                SOME(CL.I_Exp(CL.mkBinOp(CL.mkSubscript(CL.mkVar "workQueue",CL.mkVar "workerIndex"),CL.#+,CL.mkBinOp(localOffset,CL.#*,CL.mkVar "limit")))))
570                            ]
571    
572                  (* Setup Kernel arguments *)              val strandDecl = [
573                  val kernelArguments = [CL.mkAssign(CL.mkVar errVar,CL.mkApply("clSetKernelArg",                        CL.mkDecl(CL.T_Named tyName, "selfIn", NONE),
574                                                                  [CL.mkVar kernelVar,                        CL.mkDecl(CL.T_Named tyName, "selfOut", NONE)
575                                                                   CL.mkInt(0,CL.intTy),                      ]
576                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),                val imageDataDecl = CL.mkDecl(CL.T_Named(RN.imageDataType),RN.globalImageDataName,NONE)
577                                                                   CL.mkUnOp(CL.%&,CL.mkVar clInstateVar)])),                val imageDataStms = List.map (fn (x,_) =>
578                                                              CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar, CL.|=,CL.mkApply("clSetKernelArg",                    CL.mkAssign(CL.mkSelect(CL.mkVar(RN.globalImageDataName),RN.imageDataName x),
579                                                                  [CL.mkVar kernelVar,                                CL.mkVar(RN.addBufferSuffixData x))) (!imgGlobals)
580                                                                   CL.mkInt(1,CL.intTy),  
581                                                                   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])]  
582    
583    
584           fun strandCopy(inStrand, outStrand) = CL.mkCall(RN.strandCopy name,[inStrand,outStrand])
585           val updateStm =  CL.mkAssign(CL.mkVar "status",
586                            CL.mkApply(RN.strandUpdate name,
587                              [CL.mkUnOp(CL.%&,CL.mkVar inState),
588                               CL.mkUnOp(CL.%&,CL.mkVar outState),
589                               CL.mkVar RN.globalsVarName,
590                               CL.mkVar RN.globalImageDataName]))
591    
592                  (* Body put all the statments together *)                (*      CL.mkBlock([CL.mkAssign(CL.mkSubscript(CL.mkVar "strandStatus",CL.mkVar "strandIndex"),CL.mkVar RN.kStable),
593                  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")),
594                                     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))])]),
595                                     kernelArguments @ clGlobalArguments @ enqueueStm @  [outputStm] @ freeStms @ outputData *)                      CL.mkBlock([CL.mkIfThen(CL.mkBinOp(CL.E_Var "status", CL.#==, CL.E_Var RN.kDie),
596                                    CL.mkBlock([CL.mkAssign(CL.mkSubscript(CL.mkVar "strandStatus",CL.mkVar "strandIndex"),CL.mkVar RN.kDie),
597                                    strandCopy(CL.mkUnOp(CL.%&,CL.mkVar "selfOut"),CL.mkBinOp(CL.mkVar "strands", CL.#+, CL.mkVar "strandIndex")),
598                                    CL.mkCall(RN.atom_dec,[CL.mkUnOp(CL.%&,CL.mkSubscript(CL.mkVar "numAvail",CL.mkInt 0))])])) *)
599    
600           val statusIf = CL.mkIfThenElse(CL.mkBinOp(CL.E_Var "status", CL.#==, CL.E_Var RN.kStabilize),
601                        CL.mkBlock([CL.mkAssign(CL.mkSubscript(CL.mkVar "strandStatus",CL.mkVar "strandIndex"),CL.mkVar RN.kStable),
602                                    strandCopy(CL.mkUnOp(CL.%&,CL.mkVar "selfOut"),CL.mkBinOp(CL.mkVar "strands", CL.#+, CL.mkVar "strandIndex")),
603                                    CL.mkCall(RN.atom_dec,[CL.mkUnOp(CL.%&,CL.mkSubscript(CL.mkVar "numAvail",CL.mkInt 0))])]),
604                        CL.mkBlock([CL.mkIfThen(CL.mkBinOp(CL.E_Var "status", CL.#==, CL.E_Var RN.kDie),
605                                    CL.mkBlock([CL.mkAssign(CL.mkSubscript(CL.mkVar "strandStatus",CL.mkVar "strandIndex"),CL.mkVar RN.kDie),
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    
609           val incStrand = CL.mkExpStm(CL.mkPostOp(CL.mkVar "strandIndex",CL.^++))
610    
611            val forStablize = CL.mkFor( [(CL.intTy,"idx",CL.mkInt 0)], CL.mkBinOp(CL.mkBinOp(CL.mkVar "idx", CL.#<, CL.mkVar "limit"),CL.#&&,
612                                                                         CL.mkBinOp(CL.mkVar "strandIndex", CL.#<, CL.mkVar "numStrands")),
613                                           [CL.mkPostOp(CL.mkVar "idx", CL.^++)], CL.mkBlock(
614                                           [
615                                             strandCopy(CL.mkBinOp(CL.mkVar "strands", CL.#+, CL.mkVar "strandIndex"),CL.mkUnOp(CL.%&,CL.mkVar "selfIn")),
616                                             updateStm,
617                                             statusIf,
618                                             incStrand
619                                           ]))
620    
621                  val local_vars = index_ids
622                        @ [imageDataDecl]
623                        @ imageDataStms
624                        @ strandDecl
625                        @ status
626    
627                  val body = CL.mkBlock(local_vars @ [forStablize])
628                  in                  in
629                    CL.D_Func(["__kernel"], CL.voidTy, fName, params, body)
                 CL.D_Func([],CL.voidTy,RN.setupFName,params,CL.mkBlock(body))  
   
630                  end                  end
631  (* generate the data and global parameters *)  
632          fun genKeneralGlobalParams ((name,tyname)::rest) =        (* generate a global structure type definition from the list of globals *)
633                  CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), concat[RN.globalsVarName]) ::          fun genGlobalStruct (targetTy, globals, tyName) = let
634                  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([]) = []  
635                  in                  in
636                    initGlobalStruct(globals) @ initGlobalImages(imgGlobals)                  CL.D_StructDef(globs, tyName)
637                  end                  end
638    
639          (* generate the main kernel function for the .cl file *)        (* generate a global structure type definition from the image data of the image globals *)
640          fun genKernelFun(Strand{name, tyName, state, output, code,...},nDims,globals,imgGlobals) = let          fun genImageDataStruct (imgGlobals, tyName) = let
641                   val fName = RN.kernelFuncName;                val globs = List.map
642                   val inState = "strand_in"                      (fn (x, _) => (globalPtr CL.voidTy, RN.imageDataName x))
643                   val outState = "strand_out"                        imgGlobals
644               val params = [                in
645                        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)]  
646                                  end                                  end
647    
648            fun genGlobals (declFn, targetTy, globals) = let
649                  fun doVar (x : mirror_var) = declFn (CL.D_Var([], targetTy x, #var x, NONE))
650                  in
651                    List.app doVar globals
652                  end
653    
654                    val status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy))))          fun genStrandDesc (Strand{name, output, ...}) = let
655                    val local_vars =  thread_ids @ initKernelGlobals(!globals,!imgGlobals)  @ strandDecl @ strandObjects @ [status]              (* the strand's descriptor object *)
656                    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
657                    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))
658                                                          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)  
659                  in                  in
660                     CL.D_Func(["__kernel"], CL.voidTy, fName, params, body)                        CL.I_Struct[
661                              ("name", CL.I_Exp(CL.mkStr name)),
662                              ("stateSzb", CL.I_Exp(CL.mkSizeof(CL.T_Named(RN.strandTy name)))),
663    (*
664                              ("outputSzb", CL.I_Exp(CL.mkSizeof(ToC.trTy outTy))),
665    *)
666                              ("update", fnPtr("update_method_t", "0")),
667                              ("print", fnPtr("print_method_t", name ^ "_print"))
668                            ]
669                  end                  end
670          (* 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)  
671                   in                   in
672                          CL.D_StructDef(getGlobals(globals),RN.globalsTy)                  desc
673                    end                    end
674    
675        (* generate the table of strand descriptors *)        (* generate the table of strand descriptors *)
676          fun genStrandTable (ppStrm, strands) = let          fun genStrandTable (declFn, strands) = let
677                val nStrands = length strands                val nStrands = length strands
678                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)))
679                fun genInits (_, []) = []                fun genInits (_, []) = []
680                  | 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)  
681                in                in
682                  ppDecl (CL.D_Var([], CL.int32, RN.numStrands,                  declFn (CL.D_Var([], CL.int32, N.numStrands,
683                    SOME(CL.I_Exp(CL.mkInt(IntInf.fromInt nStrands, CL.int32)))));                    SOME(CL.I_Exp(CL.E_Int(IntInf.fromInt nStrands, CL.int32)))));
684                  ppDecl (CL.D_Var([],                  declFn (CL.D_Var([],
685                    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),
686                    RN.strands,                    N.strands,
687                    SOME(CL.I_Array(genInits (0, strands)))))                    SOME(CL.I_Array(genInits (0, strands)))))
688                end                end
689    
690            fun genSrc (baseName, prog) = let
691          fun genSrc (baseName, Prog{double,globals, topDecls, strands, initially,imgGlobals,numDims,...}) = let                val Prog{name,double, globals, topDecls, strands, initially, imgGlobals, numDims, ...} = prog
692                val clFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "cl"}                val clFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "cl"}
693                val cFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "c"}                val cFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "c"}
694                val clOutS = TextIO.openOut clFileName                val clOutS = TextIO.openOut clFileName
695                val cOutS = TextIO.openOut cFileName                val cOutS = TextIO.openOut cFileName
696  (* FIXME: need to use PrintAsC and PrintAsCL *)                val clppStrm = PrintAsCL.new clOutS
               val clppStrm = PrintAsC.new clOutS  
697                val cppStrm = PrintAsC.new cOutS                val cppStrm = PrintAsC.new cOutS
698                  val progName = name
699                fun cppDecl dcl = PrintAsC.output(cppStrm, dcl)                fun cppDecl dcl = PrintAsC.output(cppStrm, dcl)
700                fun clppDecl dcl = PrintAsCL.output(clppStrm, dcl)                fun clppDecl dcl = PrintAsCL.output(clppStrm, dcl)
701                val strands = AtomTable.listItems strands                val strands = AtomTable.listItems strands
# Line 941  Line 707 
707                        then "#define DIDEROT_DOUBLE_PRECISION"                        then "#define DIDEROT_DOUBLE_PRECISION"
708                        else "#define DIDEROT_SINGLE_PRECISION",                        else "#define DIDEROT_SINGLE_PRECISION",
709                      "#define DIDEROT_TARGET_CL",                      "#define DIDEROT_TARGET_CL",
710                      "#include \"Diderot/cl-types.h\""                      "#include \"Diderot/cl-diderot.h\""
711                    ]));                    ]));
712                  List.app clppDecl (List.rev (!globals));                  clppDecl (genGlobalStruct (#gpuTy, !globals, RN.globalsTy));
713                  clppDecl (genGlobalStruct (!globals));                  clppDecl (genImageDataStruct(!imgGlobals,RN.imageDataType));
714                  clppDecl (genStrandTyDef strand);                  clppDecl (genStrandTyDef(#gpuTy, strand,tyName));
715                    clppDecl (genStrandCopy(strand));
716                  List.app clppDecl (!code);                  List.app clppDecl (!code);
717                  clppDecl (genKernelFun (strand,!numDims,globals,imgGlobals));                  clppDecl (genKernelFun (strand,!numDims,globals,imgGlobals));
718                (* Generate the Host file .c *)                (* Generate the Host C file *)
719                  cppDecl (CL.D_Verbatim([                  cppDecl (CL.D_Verbatim([
720                      if double                      if double
721                        then "#define DIDEROT_DOUBLE_PRECISION"                        then "#define DIDEROT_DOUBLE_PRECISION"
# Line 956  Line 723 
723                      "#define DIDEROT_TARGET_CL",                      "#define DIDEROT_TARGET_CL",
724                      "#include \"Diderot/diderot.h\""                      "#include \"Diderot/diderot.h\""
725                    ]));                    ]));
726                  List.app cppDecl (List.rev (!globals));                  cppDecl (CL.D_Var(["static"], CL.charPtr, "ProgramName",
727                  cppDecl (genGlobalStruct (!globals));                    SOME(CL.I_Exp(CL.mkStr progName))));
728                  cppDecl (genStrandTyDef strand);                  cppDecl (genGlobalStruct (#hostTy, !globals, RN.globalsTy));
729                    cppDecl (genGlobalStruct (#shadowTy, !globals, RN.shadowGlobalsTy));
730    (* FIXME: does this really need to be a global? *)
731                    cppDecl (CL.D_Var(["static"], globPtrTy, RN.globalsVarName, NONE));
732                    cppDecl (genStrandTyDef (#hostTy, strand, tyName));
733                  cppDecl  (!init_code);                  cppDecl  (!init_code);
734                  cppDecl (genStrandInit(strand,!numDims));                  cppDecl (genStrandPrint strand);
                 cppDecl (genStrandPrint(strand,!numDims));  
                 (* cppDecl (genKernelLoader());*)  
735                  List.app cppDecl (List.rev (!topDecls));                  List.app cppDecl (List.rev (!topDecls));
736                  cppDecl (genHostSetupFunc (strand, clFileName, !numDims, initially, imgGlobals));                  cppDecl (genGlobalBuffersArgs (!globals,imgGlobals));
737                    List.app (fn strand => cppDecl (genStrandDesc strand)) strands;
738                    genStrandTable (cppDecl, strands);
739                    cppDecl (!initially);
740                  PrintAsC.close cppStrm;                  PrintAsC.close cppStrm;
741                  PrintAsCL.close clppStrm;                  PrintAsCL.close clppStrm;
742                  TextIO.closeOut cOutS;                  TextIO.closeOut cOutS;
743                  TextIO.closeOut clOutS                  TextIO.closeOut clOutS
744                end                end
745    
746        (* 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.  
        *)  
747          fun generate (basename, prog as Prog{double, parallel, debug, ...}) = let          fun generate (basename, prog as Prog{double, parallel, debug, ...}) = let
748                fun condCons (true, x, xs) = x::xs                fun condCons (true, x, xs) = x::xs
749                  | condCons (false, _, xs) = xs                  | condCons (false, _, xs) = xs
# Line 999  Line 769 
769                  RunCC.link (basename, ldOpts)                  RunCC.link (basename, ldOpts)
770                end                end
771    
772        end        end (* Program *)
773    
774    (* strands *)    (* strands *)
775      structure Strand =      structure Strand =
# Line 1028  Line 798 
798          fun init (Strand{name, tyName, code,init_code, ...}, params, init) = let          fun init (Strand{name, tyName, code,init_code, ...}, params, init) = let
799                val fName = RN.strandInit name                val fName = RN.strandInit name
800                val params =                val params =
801                      CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfOut") ::                      clParam ("",CL.T_Ptr(CL.T_Named tyName), "selfOut") ::
802                        List.map (fn (ToCL.V(ty, x)) => CL.PARAM([], ty, x)) params                        List.map (fn (ToCL.V(ty, x)) => CL.PARAM([], ty, x)) params
803                val initFn = CL.D_Func([], CL.voidTy, fName, params, init)                val initFn = CL.D_Func([], CL.voidTy, fName, params, init)
804                in                in
# Line 1037  Line 807 
807    
808        (* register a strand method *)        (* register a strand method *)
809          fun method (Strand{name, tyName, code,...}, methName, body) = let          fun method (Strand{name, tyName, code,...}, methName, body) = let
810                val fName = concat[name, "_", methName]                val fName = concat[name, "_", MethodName.toString methName]
811                val params = [                val params = [
812                        CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfIn"),                        clParam ("",CL.T_Ptr(CL.T_Named tyName), "selfIn"),
813                        CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfOut")                        clParam ("",CL.T_Ptr(CL.T_Named tyName), "selfOut"),
814                          globalParam (CL.T_Ptr(CL.T_Named (RN.globalsTy)), RN.globalsVarName),
815                          CL.PARAM([],CL.T_Named(RN.imageDataType),RN.globalImageDataName)
816                      ]                      ]
817                val methFn = CL.D_Func([], CL.int32, fName, params, body)                val resTy = (case methName
818                         of MethodName.Update => CL.T_Named "StrandStatus_t"
819                          | MethodName.Stabilize => CL.voidTy
820                        (* end case *))
821                  val methFn = CL.D_Func([], resTy, fName, params, body)
822                in                in
823                  code := methFn :: !code                  code := methFn :: !code
824                end                end

Legend:
Removed from v.1285  
changed lines
  Added in v.1460

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