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 1281, Mon Jun 6 18:21:50 2011 UTC revision 1421, Thu Jun 30 21:22:28 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      type var = ToCL.var    (* 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(CL.mkSelect(CL.mkVar(RN.shadowGlaobalsName),name),
34                                    CL.mkIndirect(CL.mkVar(RN.globalsVarName), name))
35                | Ty.TensorTy[n]=> CL.mkCall(RN.convertToShadowVec n,[
36                             CL.mkSelect(CL.mkVar(RN.shadowGlaobalsName),name),
37                              CL.mkIndirect(CL.mkVar(RN.globalsVarName), name)])
38                | Ty.ImageTy(ImageInfo.ImgInfo{dim, ...}) =>  CL.mkCall(RN.shadowImageFunc dim, [
39                             CL.mkVar "context",
40                             CL.mkUnOp(CL.%&,CL.mkSelect(CL.mkVar(RN.shadowGlaobalsName),name)),
41                             CL.mkIndirect(CL.mkVar(RN.globalsVarName),name)
42                             ])
43                | Ty.TensorTy[n, m] => CL.mkCall(RN.convertToShadowMat(m,n),[
44                             CL.mkSelect(CL.mkVar(RN.shadowGlaobalsName),name),
45                              CL.mkIndirect(CL.mkVar(RN.globalsVarName), name)])
46                | _ => CL.mkAssign(CL.mkSelect(CL.mkVar(RN.shadowGlaobalsName),name),
47                                    CL.mkIndirect(CL.mkVar(RN.globalsVarName), name))
48               (*end case *))
49    
50      (* helper functions for specifying parameters in various address spaces *)
51        fun clParam (spc, ty, x) = CL.PARAM([spc], ty, x)
52        fun globalParam (ty, x) = CL.PARAM(["__global"], ty, x)
53        fun constantParam (ty, x) = CL.PARAM(["__constant"], ty, x)
54        fun localParam (ty, x) = CL.PARAM(["__local"], ty, x)
55        fun privateParam (ty, x) = CL.PARAM(["__private"], ty, x)
56    
57      (* OpenCL global pointer type *)
58        fun globalPtr ty = CL.T_Qual("__global", CL.T_Ptr ty)
59    
60      (* C variable translation *)
61        structure TrCVar =
62          struct
63            type env = CL.typed_var TreeIL.Var.Map.map
64            fun lookup (env, x) = (case V.Map.find (env, x)
65                   of SOME(CL.V(_, x')) => x'
66                    | NONE => raise Fail(concat["TrCVar.lookup(_, ", V.name x, ")"])
67                  (* end case *))
68          (* translate a variable that occurs in an l-value context (i.e., as the target of an assignment) *)
69            fun lvalueVar (env, x) = (case V.kind x
70                   of IL.VK_Global => CL.mkIndirect(CL.mkVar RN.globalsVarName, lookup(env, x))
71                    | IL.VK_State strand => CL.mkIndirect(CL.mkVar "selfOut", lookup(env, x))
72                    | IL.VK_Local => CL.mkVar(lookup(env, x))
73                  (* end case *))
74          (* translate a variable that occurs in an r-value context *)
75            fun rvalueVar (env, x) = (case V.kind x
76                   of IL.VK_Global => CL.mkIndirect(CL.mkVar RN.globalsVarName, lookup(env, x))
77                    | IL.VK_State strand => CL.mkIndirect(CL.mkVar "selfIn", lookup(env, x))
78                    | IL.VK_Local => CL.mkVar(lookup(env, x))
79                  (* end case *))
80          end
81    
82        structure ToC = TreeToCFn (TrCVar)
83    
84        type var = CL.typed_var
85      type exp = CL.exp      type exp = CL.exp
86      type stm = CL.stm      type stm = CL.stm
87    
88    (* OpenCL specific types *)    (* OpenCL specific types *)
89        val clIntTy = CL.T_Named "cl_int"
90      val clProgramTy = CL.T_Named "cl_program"      val clProgramTy = CL.T_Named "cl_program"
91      val clKernelTy  = CL.T_Named "cl_kernel"      val clKernelTy  = CL.T_Named "cl_kernel"
92      val clCmdQueueTy = CL.T_Named "cl_command_queue"      val clCmdQueueTy = CL.T_Named "cl_command_queue"
# Line 26  Line 94 
94      val clDeviceIdTy = CL.T_Named "cl_device_id"      val clDeviceIdTy = CL.T_Named "cl_device_id"
95      val clPlatformIdTy = CL.T_Named "cl_platform_id"      val clPlatformIdTy = CL.T_Named "cl_platform_id"
96      val clMemoryTy = CL.T_Named "cl_mem"      val clMemoryTy = CL.T_Named "cl_mem"
97        val globPtrTy = CL.T_Ptr(CL.T_Named RN.globalsTy)
98    
99      (* variable or field that is mirrored between host and GPU *)
100        type mirror_var = {
101                hostTy : CL.ty,             (* variable type on Host (i.e., C type) *)
102                shadowTy : CL.ty,           (* host-side shadow type of GPU type *)
103                gpuTy : CL.ty,              (* variable's type on GPU (i.e., OpenCL type) *)
104                hToS: stm,                                  (*the statement that converts the variable to its shadow representation *)
105                var : CL.var                (* variable name *)
106              }
107    
108      datatype strand = Strand of {      datatype strand = Strand of {
109          name : string,          name : string,
110          tyName : string,          tyName : string,
111          state : var list ref,          state : mirror_var list ref,
112          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) *)
113          code : CL.decl list ref,          code : CL.decl list ref,
114          init_code: CL.decl ref          init_code: CL.decl ref
# Line 41  Line 119 
119          double : bool,                  (* true for double-precision support *)          double : bool,                  (* true for double-precision support *)
120          parallel : bool,                (* true for multithreaded (or multi-GPU) target *)          parallel : bool,                (* true for multithreaded (or multi-GPU) target *)
121          debug : bool,                   (* true for debug support in executable *)          debug : bool,                   (* true for debug support in executable *)
122          globals : CL.decl list ref,          globals : mirror_var list ref,
123          topDecls : CL.decl list ref,          topDecls : CL.decl list ref,
124          strands : strand AtomTable.hash_table,          strands : strand AtomTable.hash_table,
125          initially : CL.stm list ref,          initially :  CL.decl ref,
126          numDims: int ref,          numDims: int ref,               (* number of dimensions in initially iteration *)
127          imgGlobals: (string * int) list ref,          imgGlobals: (string * int) list ref,
128          prFn: CL.decl ref          prFn: CL.decl ref
129        }        }
# Line 79  Line 157 
157      structure Tr =      structure Tr =
158        struct        struct
159          fun fragment (ENV{info, vMap, scope}, blk) = let          fun fragment (ENV{info, vMap, scope}, blk) = let
160                val (vMap, stms) = ToCL.trFragment (vMap, blk)                val (vMap, stms) = (case scope
161                         of GlobalScope => ToC.trFragment (vMap, blk)
162    (* NOTE: if we move strand initialization to the GPU, then we'll have to change the following code! *)
163                          | InitiallyScope => ToC.trFragment (vMap, blk)
164                          | _ => ToCL.trFragment (vMap, blk)
165                        (* end case *))
166                in                in
167                  (ENV{info=info, vMap=vMap, scope=scope}, stms)                  (ENV{info=info, vMap=vMap, scope=scope}, stms)
168                end                end
169          fun saveState cxt stateVars (env, args, stm) = (          fun block (ENV{vMap, scope, ...}, blk) = let
170                  fun saveState cxt stateVars trAssign (env, args, stm) = (
171                ListPair.foldrEq                ListPair.foldrEq
172                  (fn (x, e, stms) => ToCL.trAssign(env, x, e)@stms)                        (fn (x, e, stms) => trAssign(env, x, e)@stms)
173                    [stm]                    [stm]
174                      (stateVars, args)                      (stateVars, args)
175                ) handle ListPair.UnequalLengths => (                ) handle ListPair.UnequalLengths => (
176                  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"]);
177                  raise Fail(concat["saveState ", cxt, ": length mismatch"]))                  raise Fail(concat["saveState ", cxt, ": length mismatch"]))
178          fun block (ENV{vMap, scope, ...}, blk) = (case scope                in
179                 of StrandScope stateVars => ToCL.trBlock (vMap, saveState "StrandScope" stateVars, blk)                  case scope
180                  | 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! *)
181                  | _ => ToCL.trBlock (vMap, fn (_, _, stm) => [stm], blk)                   of StrandScope stateVars =>
182                (* end case *))                        ToCL.trBlock (vMap, saveState "StrandScope" stateVars ToCL.trAssign, blk)
183                      | MethodScope stateVars =>
184                          ToCL.trBlock (vMap, saveState "MethodScope" stateVars ToCL.trAssign, blk)
185                      | InitiallyScope => ToCL.trBlock (vMap, fn (_, _, stm) => [stm], blk)
186                      | _ => ToC.trBlock (vMap, fn (_, _, stm) => [stm], blk)
187                    (* end case *)
188                  end
189          fun exp (ENV{vMap, ...}, e) = ToCL.trExp(vMap, e)          fun exp (ENV{vMap, ...}, e) = ToCL.trExp(vMap, e)
190        end        end
191    
192    (* variables *)    (* variables *)
193      structure Var =      structure Var =
194        struct        struct
195            fun mirror (ty, name) = {
196                    hostTy = ToC.trType ty,
197                    shadowTy = shadowTy ty,
198                    gpuTy = ToCL.trType ty,
199                    hToS = convertToShadow(ty,name),
200                    var = name
201                  }
202          fun name (ToCL.V(_, name)) = name          fun name (ToCL.V(_, name)) = name
203          fun global (Prog{globals, imgGlobals, ...}, name, ty) = let          fun global (Prog{globals, imgGlobals, ...}, name, ty) = let
204                val ty' = ToCL.trType ty                val x = mirror (ty, name)
205                fun isImgGlobal (imgGlobals, Ty.ImageTy(ImageInfo.ImgInfo{dim, ...}), name) =  imgGlobals  := (name,dim):: !imgGlobals                fun isImgGlobal (Ty.ImageTy(ImageInfo.ImgInfo{dim, ...}), name) =
206                  | isImgGlobal (imgGlobals, _, _) =  ()                      imgGlobals  := (name,dim) :: !imgGlobals
207                in                  | isImgGlobal _ =  ()
208                  globals := CL.D_Var([], ty', name, NONE) :: !globals;                in
209                  isImgGlobal(imgGlobals,ty,name);                  globals := x :: !globals;
210               ToCL.V(ty', name)                  isImgGlobal (ty, name);
211                    ToCL.V(#gpuTy x, name)
212                end                end
213          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)
214          fun state (Strand{state, ...}, x) = let          fun state (Strand{state, ...}, x) = let
215                val ty' = ToCL.trType(V.ty x)                val ty = V.ty x
216                val x' = ToCL.V(ty', V.name x)                val x' = mirror (ty, V.name x)
217                in                in
218                  state := x' :: !state;                  state := x' :: !state;
219                  x'                  ToCL.V(#gpuTy x', #var x')
220                end                end
221        end        end
222    
# Line 150  Line 248 
248        struct        struct
249          fun new {name, double, parallel, debug} = (          fun new {name, double, parallel, debug} = (
250                RN.initTargetSpec double;                RN.initTargetSpec double;
251                  CNames.initTargetSpec double;
252                Prog{                Prog{
253                    name = name,                    name = name,
254                    double = double, parallel = parallel, debug = debug,                    double = double, parallel = parallel, debug = debug,
255                    globals = ref [],                    globals = ref [],
256                    topDecls = ref [],                    topDecls = ref [],
257                    strands = AtomTable.mkTable (16, Fail "strand table"),                    strands = AtomTable.mkTable (16, Fail "strand table"),
258                    initially = ref([CL.S_Comment["missing initially"]]),                    initially = ref(CL.D_Comment["missing initially"]),
259                                    numDims = ref(0),                    numDims = ref 0,
260                                    imgGlobals = ref[],                                    imgGlobals = ref[],
261                                    prFn = ref(CL.D_Comment(["No Print Function"]))                                    prFn = ref(CL.D_Comment(["No Print Function"]))
262                  })                  })
       (* 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  
263    
264        (* 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 *)
265          fun inputs (Prog{topDecls, ...}, stm) = let          fun inputs (Prog{topDecls, ...}, stm) = let
# Line 182  Line 271 
271                  topDecls := inputsFn :: !topDecls                  topDecls := inputsFn :: !topDecls
272                end                end
273    
274          fun init (Prog{globals, topDecls,...}, CL.S_Block(init)) = let        (* register the global initialization part of a program *)
275                val params = [          fun init (Prog{topDecls, ...}, init) = let
276                        CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName)                val globalsDecl = CL.mkAssign(CL.E_Var RN.globalsVarName,
277                      ]                      CL.mkApply("malloc", [CL.mkSizeof(CL.T_Named RN.globalsTy)]))
278                val body = CL.S_Block(globalIndirects(!globals,init))                val initFn = CL.D_Func(
279                val initFn = CL.D_Func([], CL.voidTy, RN.initGlobals, params, body)                      [], CL.voidTy, RN.initGlobals, [],
280                in                      CL.mkBlock[
281                  topDecls := initFn :: !topDecls              globalsDecl,
282                end              CL.mkCall(RN.initGlobalsHelper, [CL.mkVar RN.globalsVarName])
283            | init (Prog{globals,topDecls,...}, init) = let                ])
284                val params = [                val initHelperFn = CL.D_Func(
285                        CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName)                      [], CL.voidTy, RN.initGlobalsHelper,
286                      ]              [CL.PARAM([], globPtrTy, RN.globalsVarName)],
287                val initFn = CL.D_Func([], CL.voidTy, RN.initGlobals, params, init)                      init)
288                  val shutdownFn = CL.D_Func(
289                        [], CL.voidTy, RN.shutdown,
290                        [CL.PARAM([], CL.T_Ptr(CL.T_Named RN.worldTy), "wrld")],
291                        CL.S_Block[])
292                in                in
293                  topDecls := initFn :: !topDecls                  topDecls := shutdownFn :: initFn :: initHelperFn :: !topDecls
294                end                end
295    
296        (* create and register the initially function for a program *)        (* create and register the initially function for a program *)
297          fun initially {          fun initially {
298                prog = Prog{strands, initially, numDims,...},                prog = Prog{name=progName, strands, initially, numDims, ...},
299                isArray : bool,                isArray : bool,
300                iterPrefix : stm list,                iterPrefix : stm list,
301                iters : (var * exp * exp) list,                iters : (var * exp * exp) list,
# Line 212  Line 305 
305              } = let              } = let
306                val name = Atom.toString strand                val name = Atom.toString strand
307                val nDims = List.length iters                val nDims = List.length iters
308                  val worldTy = CL.T_Ptr(CL.T_Named N.worldTy)
309                fun mapi f xs = let                fun mapi f xs = let
310                      fun mapf (_, []) = []                      fun mapf (_, []) = []
311                        | mapf (i, x::xs) = f(i, x) :: mapf(i+1, xs)                        | mapf (i, x::xs) = f(i, x) :: mapf(i+1, xs)
# Line 220  Line 314 
314                      end                      end
315                val baseInit = mapi (fn (i, (_, e, _)) => (i, CL.I_Exp e)) iters                val baseInit = mapi (fn (i, (_, e, _)) => (i, CL.I_Exp e)) iters
316                val sizeInit = mapi                val sizeInit = mapi
317                      (fn (i, (ToCL.V(ty, _), lo, hi)) =>                      (fn (i, (CL.V(ty, _), lo, hi)) =>
318                          (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))))
319                      ) iters                      ) iters
320                    val numStrandsVar = "numStrandsVar"              (* code to allocate the world and initial strands *)
321                val allocCode = iterPrefix @ [                val wrld = "wrld"
322                  val allocCode = [
323                        CL.mkComment["allocate initial block of strands"],                        CL.mkComment["allocate initial block of strands"],
324                        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)),
325                        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)),
326                        CL.mkDecl(CL.int32,"numDims",SOME(CL.I_Exp(CL.mkInt(IntInf.fromInt nDims, CL.int32))))                        CL.mkDecl(worldTy, wrld,
327                      ]                          SOME(CL.I_Exp(CL.E_Apply(RN.allocInitially, [
328                val numStrandsLoopBody =                              CL.mkVar "ProgramName",
329                      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)),
330                val numStrandsLoop =  CL.mkFor([(CL.intTy, "i", CL.mkInt(0,CL.intTy))],                              CL.E_Bool isArray,
331                      CL.mkBinOp(CL.mkVar "i", CL.#<, CL.mkVar "numDims"),                              CL.E_Int(IntInf.fromInt nDims, CL.int32),
332                      [CL.mkPostOp(CL.mkVar "i", CL.^++)], numStrandsLoopBody)                              CL.E_Var "base",
333                in                              CL.E_Var "size"
334                  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")  
335                      ]                      ]
336                val body = let              (* create the loop nest for the initially iterations
337                      fun loopParams 3 = ["x", "y", "k"]                val indexVar = "ix"
338                        | loopParams 2 = ["x", "y"]                val strandTy = CL.T_Ptr(CL.T_Named(N.strandTy name))
339                        | loopParams 1 = ["x"]                fun mkLoopNest [] = CL.mkBlock(createPrefix @ [
340                        | loopParams _ = raise Fail "genStrandInit: missing size dim"                        CL.mkDecl(strandTy, "sp",
341                      fun mkLoopNest ([], _, nDims) = if nDims = 1                          SOME(CL.I_Exp(
342                            then CL.mkBlock [                            CL.E_Cast(strandTy,
343                                CL.mkCall(RN.strandInit name, [                            CL.E_Apply(N.inState, [CL.E_Var "wrld", CL.E_Var indexVar]))))),
344                                  CL.mkUnOp(CL.%&,CL.mkSubscript(CL.mkVar "strands",CL.mkStr "x")),                        CL.mkCall(N.strandInit name,
345                                                  CL.mkVar "x"])                          CL.E_Var RN.globalsVarName :: CL.E_Var "sp" :: args),
346                              ]                        CL.mkAssign(CL.E_Var indexVar, CL.mkBinOp(CL.E_Var indexVar, CL.#+, CL.E_Int(1, CL.uint32)))
347                            else let                      ])
348                              val index = CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")                  | mkLoopNest ((CL.V(ty, param), lo, hi)::iters) = let
349                              in                      val body = mkLoopNest iters
                               CL.mkBlock([CL.mkCall(RN.strandInit name, [CL.mkUnOp(CL.%&,CL.mkSubscript(CL.mkVar "strands",index)),  
                               CL.mkVar "x", CL.mkVar"y"])])  
                             end  
                       | mkLoopNest (param::rest,count,nDims) = let  
                           val body = mkLoopNest (rest, count + 1,nDims)  
350                            in                            in
351                              CL.mkFor(                              CL.mkFor(
352                                  [(CL.intTy, param, CL.mkInt(0,CL.intTy))],                          [(ty, param, lo)],
353                                  CL.mkBinOp(CL.mkVar param, CL.#<, CL.mkSubscript(CL.mkVar "sizes",CL.mkInt(count,CL.intTy))),                          CL.mkBinOp(CL.E_Var param, CL.#<=, hi),
354                                  [CL.mkPostOp(CL.mkVar param, CL.^++)],                          [CL.mkPostOp(CL.E_Var param, CL.^++)],
355                                  body)                                  body)
356                            end                            end
357                  val iterCode = [
358                          CL.mkComment["initially"],
359                          CL.mkDecl(CL.uint32, indexVar, SOME(CL.I_Exp(CL.E_Int(0, CL.uint32)))),
360                          mkLoopNest iters
361                        ] *)
362                  val body = CL.mkBlock(
363                        iterPrefix @
364                        allocCode @
365                        [CL.mkReturn(SOME(CL.E_Var "wrld"))])
366                  val initFn = CL.D_Func([], worldTy, N.initially, [], body)
367                      in                      in
368                        [mkLoopNest ((loopParams nDims),0,nDims)]                  numDims := nDims;
369                      end                  initially := initFn
                 in  
                   CL.D_Func(["static"], CL.voidTy, RN.strandInitSetup, params,CL.mkBlock(body))  
370                  end                  end
371    
372          fun genStrandPrint (Strand{name, tyName, state, output, code,...},nDims) = let        (***** OUTPUT *****)
373            fun genStrandPrint (Strand{name, tyName, state, output, code,...}) = let
374              (* the print function *)              (* the print function *)
375                val prFnName = concat[name, "_print"]                val prFnName = concat[name, "Print"]
376                val prFn = let                val prFn = let
377                      val params = [                      val params = [
378                            CL.PARAM([], CL.T_Ptr(CL.T_Named "FILE"), "outS"),                            CL.PARAM([], CL.T_Ptr(CL.T_Named "FILE"), "outS"),
379                            CL.PARAM([], CL.T_Ptr(CL.uint32), "sizes" ),                              CL.PARAM([], CL.T_Ptr(CL.T_Num(RawTypes.RT_UInt8)),"status"),
380                            CL.PARAM([], CL.intTy, "width"),                              CL.PARAM([], CL.intTy,"numStrands"),
381                            CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "self")                            CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "self")
382                          ]                          ]
   
383                     val SOME(ty, x) = !output                     val SOME(ty, x) = !output
384                     val outState = if nDims = 1 then                      val outState = CL.mkSelect(CL.mkSubscript(CL.mkVar "self", CL.E_Var "i"), x)
                           CL.mkSelect(CL.mkSubscript(CL.mkVar "self",CL.mkVar "x"), x)  
                         else if nDims = 2 then  
                                 CL.mkSelect(CL.mkSubscript(CL.mkVar "self",  
                                    CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")), x)  
   
                         else CL.mkSelect(CL.mkVar "self",x)  
   
385                      val prArgs = (case ty                      val prArgs = (case ty
386                             of Ty.IVecTy 1 => [CL.mkStr(!RN.gIntFormat ^ "\n"), outState]                             of Ty.IVecTy 1 => [CL.E_Str(!N.gIntFormat ^ "\n"), outState]
387                              | Ty.IVecTy d => let                              | Ty.IVecTy d => let
388                                  val fmt = CL.mkStr(                                  val fmt = CL.mkStr(
389                                        String.concatWith " " (List.tabulate(d, fn _ => !RN.gIntFormat))                                        String.concatWith " " (List.tabulate(d, fn _ => !N.gIntFormat))
390                                        ^ "\n")                                        ^ "\n")
391                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))                                  val args = List.tabulate (d, fn i => ToC.ivecIndex(outState, d, i))
392                                  in                                  in
393                                    fmt :: args                                    fmt :: args
394                                  end                                  end
# Line 316  Line 397 
397                                  val fmt = CL.mkStr(                                  val fmt = CL.mkStr(
398                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))
399                                        ^ "\n")                                        ^ "\n")
400                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))                                  val args = List.tabulate (d, fn i => ToC.vecIndex(outState, d, i))
401                                  in                                  in
402                                    fmt :: args                                    fmt :: args
403                                  end                                  end
404                              | _ => raise Fail("genStrand: unsupported output type " ^ Ty.toString ty)                              | _ => raise Fail("genStrand: unsupported output type " ^ Ty.toString ty)
405                            (* end case *))                            (* end case *))
406                        val forBody = CL.mkIfThen(
407                            val body = let                            CL.mkBinOp(CL.mkSubscript(CL.E_Var "status",CL.E_Var "i"), CL.#==, CL.E_Var "DIDEROT_STABILIZE"),
408                              CL.mkBlock([CL.mkCall("fprintf", CL.mkVar "outS" :: prArgs)]))
409                              fun loopParams (3) =                      val body =  CL.mkFor(
410                                   "x"::"y"::"k"::[]                          [(CL.intTy, "i", CL.mkInt 0)],
411                                | loopParams (2) =                          CL.mkBinOp(CL.E_Var "i", CL.#<, CL.E_Var "numStrands"),
412                                   "x"::"y"::[]                          [CL.mkPostOp(CL.E_Var "i", CL.^++)],
413                                | loopParams (1) =                          forBody)
                                  "x"::[]  
                               | loopParams (_) =  
                                 raise Fail("genStrandPrint: unsupported output type " ^ Ty.toString ty)  
   
                            fun mkLoopNest ([],_) =  
                                                 CL.mkCall("fprintf", CL.mkVar "outS" :: prArgs)  
                                 | mkLoopNest (param::rest,count) = let  
                                         val body = mkLoopNest (rest, count + 1)  
414                                     in                                     in
415                                                  CL.mkFor(                        CL.D_Func(["static"], CL.voidTy, prFnName, params, body)
                                                         [(CL.intTy, param, CL.mkInt(0,CL.intTy))],  
                                                 CL.mkBinOp(CL.mkVar param, CL.#<, CL.mkSubscript(CL.mkVar "sizes",CL.mkInt(count,CL.intTy))),  
                                                 [CL.mkPostOp(CL.mkVar param, CL.^++)],  
                                                 body)  
                                    end  
                         in  
                                 [mkLoopNest ((loopParams nDims),0)]  
                         end  
   
                     in  
                       CL.D_Func(["static"], CL.voidTy, prFnName, params,CL.mkBlock(body))  
416                      end                      end
417                in                in
418                                   prFn                                   prFn
419                end                end
420          fun genStrandTyDef (Strand{tyName, state,...}) =  
421            fun genStrandTyDef (targetTy, Strand{tyName, state,...}) =
422              (* the type declaration for the strand's state struct *)              (* the type declaration for the strand's state struct *)
423                CL.D_StructDef(                CL.D_StructDef(
424                        List.rev (List.map (fn ToCL.V(ty, x) => (ty, x)) (!state)),                  List.rev (List.map (fn x => (targetTy x, #var x)) (!state)),
425                        tyName)                        tyName)
426    
   
         (* generates the load kernel function *)  
 (* FIXME: this code might be part of the runtime system *)  
         fun genKernelLoader() =  
                 CL.D_Verbatim ( ["/* Loads the Kernel from a file */",  
                                                 "char * loadKernel (const char * filename) {",  
                                                 "struct stat statbuf;",  
                                                 "FILE *fh;",  
                                                 "char *source;",  
                                                 "fh = fopen(filename, \"r\");",  
                                                 "if (fh == 0)",  
                                                 "   return 0;",  
                                                 "stat(filename, &statbuf);",  
                                                 "source = (char *) malloc(statbuf.st_size + 1);",  
                                                 "fread(source, statbuf.st_size, 1, fh);",  
                                                 "fread(source, statbuf.st_size, 1, fh);",  
                                                 "return source;",  
                                                 "}"])  
427  (* generates the opencl buffers for the image data *)  (* generates the opencl buffers for the image data *)
428          fun getGlobalDataBuffers(globals,count,contextVar,errVar) = let          fun getGlobalDataBuffers (globals, imgGlobals, contextVar, errVar) = let
429                  val globalBuffErr = "error creating OpenCL global buffer"
430                  fun errorFn msg = CL.mkIfThen(CL.mkBinOp(CL.E_Var errVar, CL.#!=, CL.E_Var "CL_SUCCESS"),
431                        CL.mkBlock([CL.mkCall("fprintf",[CL.E_Var "stderr", CL.E_Str msg]),
432                        CL.mkCall("exit",[CL.mkInt 1])]))
433                  val shadowTypeDecl =
434                        CL.mkDecl(CL.T_Named(RN.shadowGlobalsTy), RN.shadowGlaobalsName, NONE)
435                  val globalToShadowStms = List.map (fn (x:mirror_var) => #hToS x ) globals
436                  val globalBufferDecl =  CL.mkDecl(clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)                  val globalBufferDecl =  CL.mkDecl(clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)
437                  val globalBuffer = CL.mkAssign(CL.mkVar(concat[RN.globalsVarName,"_cl"]), CL.mkApply("clCreateBuffer",                val globalBuffer = CL.mkAssign(CL.mkVar(concat[RN.globalsVarName,"_cl"]),
438                                                                  [CL.mkVar contextVar,                      CL.mkApply("clCreateBuffer", [
439                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",                          CL.mkVar contextVar,
440                                                                  CL.mkApply("sizeof",[CL.mkVar RN.globalsTy]),                          CL.mkBinOp(CL.mkVar "CL_MEM_READ_ONLY", CL.#|, CL.mkVar "CL_MEM_COPY_HOST_PTR"),
441                                                                  CL.mkVar RN.globalsVarName,                          CL.mkSizeof(CL.T_Named RN.shadowGlobalsTy),
442                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))                          CL.mkUnOp(CL.%&,CL.mkVar RN.shadowGlaobalsName),
443                            CL.mkUnOp(CL.%&,CL.mkVar errVar)
444                          ]))
445          fun genDataBuffers([],_,_,_) = []          fun genDataBuffers([],_,_,_) = []
446            | genDataBuffers((var,nDims)::globals,count,contextVar,errVar) = let                  | genDataBuffers ((var,nDims)::globals, contextVar, errVar,errFn) = let
447          (* FIXME: use CL constructors to  build expressions (not strings) *)                      val hostVar = CL.mkIndirect(CL.mkVar RN.globalsVarName, var)
448                    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]"]))  
   
449                   in                   in
                    CL.mkDecl(clMemoryTy, RN.addBufferSuffix var ,NONE)::  
450                     CL.mkDecl(clMemoryTy, RN.addBufferSuffixData var ,NONE)::                     CL.mkDecl(clMemoryTy, RN.addBufferSuffixData var ,NONE)::
451                     CL.mkAssign(CL.mkVar(RN.addBufferSuffix var), CL.mkApply("clCreateBuffer",                        CL.mkAssign(CL.mkVar(RN.addBufferSuffixData var),
452                                                                  [CL.mkVar contextVar,                          CL.mkApply("clCreateBuffer", [
453                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",                              CL.mkVar contextVar,
454                                                                  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",  
455                                                                  size,                                                                  size,
456                                                                  CL.mkIndirect(CL.mkVar var,"data"),                              CL.mkIndirect(hostVar, "data"),
457                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)])):: genDataBuffers(globals,count + 2,contextVar,errVar)                              CL.mkUnOp(CL.%&,CL.mkVar errVar)
458                              ])) ::
459                            errFn(concat["error in creating ",RN.addBufferSuffixData var, " global buffer"]) ::
460                            genDataBuffers(globals,contextVar,errVar,errFn)
461                  end                  end
462          in          in
463                  [globalBufferDecl] @ [globalBuffer] @ genDataBuffers(globals,count + 2,contextVar,errVar)                  [shadowTypeDecl] @ globalToShadowStms
464                    @ [globalBufferDecl, globalBuffer,errorFn(globalBuffErr)]
465                    @ genDataBuffers(imgGlobals,contextVar,errVar,errorFn)
466          end          end
467    
   
468  (* generates the kernel arguments for the image data *)  (* generates the kernel arguments for the image data *)
469          fun genGlobalArguments(globals,count,kernelVar,errVar) = let          fun genGlobalArguments(globals,count,kernelVar,errVar) = let
470          val globalArgument = CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,CL.mkApply("clSetKernelArg",                val globalArgErr = "error creating OpenCL global argument"
471                  fun errorFn msg = CL.mkIfThen(CL.mkBinOp(CL.E_Var errVar, CL.#!=, CL.E_Var "CL_SUCCESS"),
472                        CL.mkBlock([CL.mkCall("fprintf",[CL.E_Var "stderr", CL.E_Str msg]),
473                        CL.mkCall("exit",[CL.mkInt 1])]))
474                  val globalArgument = CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.&=,
475                        CL.mkApply("clSetKernelArg",
476                                                                  [CL.mkVar kernelVar,                                                                  [CL.mkVar kernelVar,
477                                                                   CL.mkInt(count,CL.intTy),                         CL.mkPostOp(CL.E_Var count, CL.^++),
478                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
479                                                                   CL.mkUnOp(CL.%&,CL.mkVar(concat[RN.globalsVarName,"_cl"]))])))                                                                   CL.mkUnOp(CL.%&,CL.mkVar(concat[RN.globalsVarName,"_cl"]))])))
480                  fun genDataArguments ([],_,_,_,_) = []
481          fun genDataArguments([],_,_,_) = []                  | genDataArguments ((var,nDims)::globals,count,kernelVar,errVar,errFn) =
482            | genDataArguments((var,nDims)::globals,count,kernelVar,errVar) =                      CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.$=,
483                          CL.mkApply("clSetKernelArg",
                 CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=, CL.mkApply("clSetKernelArg",  
484                                                                  [CL.mkVar kernelVar,                                                                  [CL.mkVar kernelVar,
485                                                                   CL.mkInt(count,CL.intTy),                           CL.mkPostOp(CL.E_Var count, CL.^++),
486                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
487                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffix var))])))::                           CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffixData var))]))) ::
488                             errFn(concat["error in creating ",RN.addBufferSuffixData var, " argument"]) ::
489                          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)  
   
490          in          in
491                   [globalArgument,errorFn(globalArgErr)] @ genDataArguments(globals, count, kernelVar, errVar,errorFn)
                 [globalArgument] @ genDataArguments(globals,count + 1,kernelVar,errVar)  
   
492          end          end
493    
494          (* generates the main function of host code *)        (* generates the globals buffers and arguments function *)
495          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  
496              (* 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"  
497                val errVar = "err"                val errVar = "err"
498                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")])  
499                val params = [                val params = [
500                        CL.PARAM([],CL.T_Named("cl_device_id"), deviceVar)                        CL.PARAM([],CL.T_Named("cl_context"), "context"),
501                      ]                        CL.PARAM([],CL.T_Named("cl_kernel"), "kernel"),
502                val declarations = [                        CL.PARAM([],CL.T_Named("cl_command_queue"), "cmdQ"),
503                      CL.mkDecl(clProgramTy, programVar, NONE),                        CL.PARAM([],CL.T_Named("int"), "argStart")
                     CL.mkDecl(clKernelTy, kernelVar, NONE),  
                     CL.mkDecl(clCmdQueueTy, cmdVar, NONE),  
                     CL.mkDecl(clContextTy, contextVar, NONE),  
                     CL.mkDecl(CL.intTy, errVar, NONE),  
                     CL.mkDecl(CL.intTy, numStrandsVar, SOME(CL.I_Exp(CL.mkInt(1,CL.intTy)))),  
                     CL.mkDecl(CL.intTy, stateSizeVar, NONE),  
                     CL.mkDecl(CL.intTy, "width", NONE),  
                     CL.mkDecl(CL.intTy, imgDataSizeVar, NONE),  
                     (*CL.mkDecl(clDeviceIdTy, deviceVar, NONE), *)  
                     CL.mkDecl(CL.T_Ptr(CL.T_Named tyName), inStateVar,NONE),  
                     CL.mkDecl(clMemoryTy,clInstateVar,NONE),  
                     CL.mkDecl(clMemoryTy,clOutStateVar,NONE),  
                     CL.mkDecl(CL.T_Ptr(CL.T_Named tyName), outStateVar,NONE),  
                     CL.mkDecl(CL.charPtr, clFNVar,SOME(CL.I_Exp(CL.mkStr filename))),  
 (* FIXME:  use Paths.diderotInclude *)  
                     CL.mkDecl(CL.charPtr, headerFNVar,SOME(CL.I_Exp(CL.mkStr "../src/include/Diderot/cl-types.h"))),  
                     CL.mkDecl(CL.T_Array(CL.charPtr,SOME(2)),sourcesVar,NONE),  
                     CL.mkDecl(CL.T_Array(CL.T_Named "size_t",SOME(nDims)),globalVar,NONE),  
                     CL.mkDecl(CL.T_Array(CL.T_Named "size_t",SOME(nDims)),localVar,NONE),  
                     CL.mkDecl(CL.intTy,numDevicesVar,SOME(CL.I_Exp(CL.mkInt(~1,CL.intTy)))),  
                     CL.mkDecl(CL.T_Array(CL.T_Named "cl_platform_id", SOME(1)), platformsVar, NONE),  
                     CL.mkDecl(CL.intTy,"num_platforms",SOME(CL.I_Exp(CL.mkInt(~1,CL.intTy))))  
504                  ]                  ]
505              (* Setup Global Variables *)                val clGlobalBuffers = getGlobalDataBuffers(globals,!imgGlobals, "context", errVar)
506                val globalsDecl = CL.mkDecl(                val clGlobalArguments = genGlobalArguments(!imgGlobals, "argStart", "kernel", errVar)
                     CL.T_Ptr(CL.T_Named RN.globalsTy),  
                     RN.globalsVarName,  
                     SOME(CL.I_Exp(CL.mkApply("malloc", [CL.mkApply("sizeof",[CL.mkVar RN.globalsTy])]))))  
               val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.mkVar RN.globalsVarName])  
   
                 (* Retrieve the platforms  
                 val platformStm = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetPlatformIDs",  
                                                   [CL.mkInt(10,CL.intTy),  
                                                    CL.mkVar platformsVar,  
                                                    CL.mkUnOp(CL.%&,CL.mkVar numPlatformsVar)])),  
                                                    assertStm]  
   
                 val devicesStm = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetDeviceIDs",  
                                                   [CL.mkSubscript(CL.mkVar platformsVar,CL.mkInt(0,CL.intTy)),  
                                                    CL.mkVar "CL_DEVICE_TYPE_GPU",  
                                                    CL.mkInt(1,CL.intTy),  
                                                    CL.mkUnOp(CL.%&,CL.mkVar deviceVar),  
                                                    CL.mkUnOp(CL.%&,CL.mkVar numDevicesVar)])),  
                                                    assertStm] *)  
   
                 (* Create Context *)  
                 val contextStm = [CL.mkAssign(CL.mkVar contextVar, CL.mkApply("clCreateContext",  
                                                   [CL.mkInt(0,CL.intTy),  
                                                   CL.mkInt(1,CL.intTy),  
                                                   CL.mkUnOp(CL.%&,CL.mkVar deviceVar),  
                                                   CL.mkVar "NULL",  
                                                   CL.mkVar "NULL",  
                                                   CL.mkUnOp(CL.%&,CL.mkVar errVar)])),  
                                                   assertStm]  
   
                 (* Create Command Queue *)  
                 val commandStm = [CL.mkAssign(CL.mkVar cmdVar, CL.mkApply("clCreateCommandQueue",  
                                                   [CL.mkVar contextVar,  
                                                   CL.mkVar deviceVar,  
                                                   CL.mkInt(0,CL.intTy),  
                                                   CL.mkUnOp(CL.%&,CL.mkVar errVar)])),  
                                                   assertStm]  
   
   
                 (*Create Program/Build/Kernel with Source statement *)  
                 val createProgStm = CL.mkAssign(CL.mkVar programVar, CL.mkApply("clCreateProgramWithSource",  
                                                                                                                 [CL.mkVar contextVar,  
                                                                                                                  CL.mkInt(2,CL.intTy),  
                                                                                                                  CL.mkCast(CL.T_Ptr(CL.T_Named("const char *")),CL.mkUnOp(CL.%&,CL.mkVar sourcesVar)),  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))  
   
                 (* FIXME: Remove after testing purposes, Build Log for OpenCL*)  
                 val buildLog = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clBuildProgram",  
                                                                                                                 [CL.mkVar programVar,  
                                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkVar "NULL"])),  
                                           CL.mkDecl(CL.charPtr, "build", NONE),  
                                           CL.mkDecl(CL.T_Named("size_t"),"ret_val_size",NONE),  
                                            CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetProgramBuildInfo",  
                                                                                                                 [CL.mkVar programVar,  
                                                                                                                 CL.mkVar deviceVar,  
                                                                                                                  CL.mkVar "CL_PROGRAM_BUILD_LOG",  
                                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkUnOp(CL.%&,CL.mkVar "ret_val_size")])),  
                                           CL.mkAssign(CL.mkVar "build", CL.mkApply("malloc", [CL.mkVar "ret_val_size"])),  
                                                 CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetProgramBuildInfo",  
                                                                                                                 [CL.mkVar programVar,  
                                                                                                                 CL.mkVar deviceVar,  
                                                                                                                  CL.mkVar "CL_PROGRAM_BUILD_LOG",  
                                                                                                                  CL.mkVar "ret_val_size",  
                                                                                                                  CL.mkVar "build",  
                                                                                                                  CL.mkVar "NULL"])),  
                                                 CL.mkAssign(CL.mkSubscript(CL.mkVar "build",CL.mkVar "ret_val_size"),CL.mkVar ("'\\" ^ "0'")),  
                                                 CL.mkCall("printf",[CL.mkStr ( "Build Log:" ^ "\n" ^ "%s" ^ "\n"), CL.mkVar "build"])]  
   
   
   
   
                 val createKernel = CL.mkAssign(CL.mkVar kernelVar, CL.mkApply("clCreateKernel",  
                                                                                                                 [CL.mkVar programVar,  
                                                                                                                  CL.mkStr RN.kernelFuncName,  
                                                                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))  
   
   
                 val create_build_stms = [createProgStm,assertStm] @ buildLog @ [assertStm,createKernel,assertStm]  
   
   
   
                 (* Create Memory Buffers for Strand States and Globals *)  
                 val strandSize = CL.mkAssign(CL.mkVar stateSizeVar,CL.mkBinOp(CL.mkApply("sizeof",  
                                                                         [CL.mkVar tyName]), CL.#*,CL.mkVar numStrandsVar))  
   
                 val clStrandObjects = [CL.mkAssign(CL.mkVar clInstateVar, CL.mkApply("clCreateBuffer",  
                                                                 [CL.mkVar contextVar,  
                                                                 CL.mkVar "CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR",  
                                                                 CL.mkVar stateSizeVar,  
                                                                 CL.mkVar "NULL",  
                                                                 CL.mkUnOp(CL.%&,CL.mkVar errVar)])),  
                                                          CL.mkAssign(CL.mkVar clOutStateVar, CL.mkApply("clCreateBuffer",  
                                                                 [CL.mkVar contextVar,  
                                                                 CL.mkVar "CL_MEM_READ_WRITE",  
                                                                 CL.mkVar stateSizeVar,  
                                                                 CL.mkVar "NULL",  
                                                                 CL.mkUnOp(CL.%&,CL.mkVar errVar)]))]  
   
   
                 (* Setup up selfOut variable *)  
                 val strandsArrays = [CL.mkAssign(CL.mkVar outStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.mkVar numStrandsVar,  
                                                                         CL.#*, CL.mkApply("sizeof",[CL.mkVar tyName]))])),  
                                                                 CL.mkAssign(CL.mkVar inStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.mkVar numStrandsVar,  
                                                                         CL.#*, CL.mkApply("sizeof",[CL.mkVar tyName]))]))]  
   
   
                 (* Initialize Width Parameter *)  
                 val widthDel = if nDims = 2 then  
                           CL.mkAssign(CL.mkVar "width",CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1, CL.intTy)))  
                    else  
                           CL.mkAssign(CL.mkVar "width",CL.mkInt(0,CL.intTy))  
   
   
                 val strands_init = CL.mkCall(RN.strandInitSetup, [  
                         CL.mkVar "size", CL.mkVar "width", CL.mkVar inStateVar  
                       ])  
   
             val clGlobalBuffers = getGlobalDataBuffers(!imgGlobals,3,contextVar,errVar)  
   
   
                 (* Load the Kernel and Header Files *)  
                 val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(1,CL.intTy)),  
                                                                           CL.mkApply(RN.clLoaderFN, [CL.mkVar clFNVar])),  
            CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(0,CL.intTy)),  
                                                                           CL.mkApply(RN.clLoaderFN, [CL.mkVar headerFNVar]))]  
   
                 (* val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(1,CL.intTy)),  
                                                                           CL.mkApply(RN.clLoaderFN, [CL.mkVar clFNVar]))] *)  
   
   
                 (* Created Enqueue Statements *)  
 (* FIXME: simplify this code by function abstraction *)  
         val enqueueStm = if nDims = 1  
                         then [CL.mkAssign(CL.mkVar errVar,  
                                                           CL.mkApply("clEnqueueNDRangeKernel",  
                                                                                                 [CL.mkVar cmdVar,  
                                                                                                  CL.mkVar kernelVar,  
                                                                                                  CL.mkInt(1,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar globalVar,  
                                                                                                  CL.mkVar localVar,  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]  
                         else if nDims = 2  then  
                          [CL.mkAssign(CL.mkVar errVar,  
                                                         CL.mkApply("clEnqueueNDRangeKernel",  
                                                                                                 [CL.mkVar cmdVar,  
                                                                                                  CL.mkVar kernelVar,  
                                                                                                  CL.mkInt(2,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar globalVar,  
                                                                                                  CL.mkVar localVar,  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]  
                         else  
                           [CL.mkAssign(CL.mkVar errVar,  
                                                         CL.mkApply("clEnqueueNDRangeKernel",  
                                                                                                 [CL.mkVar cmdVar,  
                                                                                                  CL.mkVar kernelVar,  
                                                                                                  CL.mkInt(3,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar globalVar,  
                                                                                                  CL.mkVar localVar,  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]  
   
   
   
                 (* Setup Global and Local variables *)  
   
                 val globalAndlocalStms = if nDims = 1 then  
                         [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),  
                          CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),  
                                                                   CL.mkVar "16")]  
   
   
                 else if nDims = 2 then  
                         [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(1,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),  
                                                                   CL.mkVar "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(1,CL.intTy)),  
                                                                   CL.mkVar "16")]  
   
                 else  
                         [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(1,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(2,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(2,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),  
                                                                   CL.mkVar "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(1,CL.intTy)),  
                                                                   CL.mkVar "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(2,CL.intTy)),  
                                                                   CL.mkVar "16")]  
   
   
   
                 (* Setup Kernel arguments *)  
                 val kernelArguments = [CL.mkAssign(CL.mkVar errVar,CL.mkApply("clSetKernelArg",  
                                                                 [CL.mkVar kernelVar,  
                                                                  CL.mkInt(0,CL.intTy),  
                                                                  CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),  
                                                                  CL.mkUnOp(CL.%&,CL.mkVar clInstateVar)])),  
                                                             CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar, CL.|=,CL.mkApply("clSetKernelArg",  
                                                                 [CL.mkVar kernelVar,  
                                                                  CL.mkInt(1,CL.intTy),  
                                                                  CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),  
                                                                  CL.mkUnOp(CL.%&,CL.mkVar clOutStateVar)]))),  
                                                                   CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar, CL.|=,CL.mkApply("clSetKernelArg",  
                                                                 [CL.mkVar kernelVar,  
                                                                  CL.mkInt(2,CL.intTy),  
                                                                  CL.mkApply("sizeof",[CL.mkVar "int"]),  
                                                                  CL.mkUnOp(CL.%&,CL.mkVar "width")])))]  
   
            val clGlobalArguments = genGlobalArguments(!imgGlobals,3,kernelVar,errVar) @ [assertStm]  
   
                 (* Retrieve output *)  
                 val outputStm = CL.mkAssign(CL.mkVar errVar,  
                                                         CL.mkApply("clEnqueueReadBuffer",  
                                                                                                 [CL.mkVar cmdVar,  
                                                                                                  CL.mkVar clOutStateVar,  
                                                                                                  CL.mkVar "CL_TRUE",  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar stateSizeVar,  
                                                                                                  CL.mkVar outStateVar,  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar "NULL"]))  
   
                 (* Free all the objects *)  
                 val freeStms = [CL.mkCall("clReleaseKernel",[CL.mkVar kernelVar]),  
                                                 CL.mkCall("clReleaseProgram",[CL.mkVar programVar ]),  
                                                 CL.mkCall("clReleaseCommandQueue",[CL.mkVar cmdVar]),  
                                                 CL.mkCall("clReleaseContext",[CL.mkVar contextVar]),  
                                                 CL.mkCall("clReleaseMemObject",[CL.mkVar clInstateVar]),  
                                                 CL.mkCall("clReleaseMemObject",[CL.mkVar clOutStateVar])]  
   
   
                 (*Setup Strand Print Function *)  
                 val outputData = [CL.mkDecl(CL.T_Ptr(CL.T_Named("FILE")), "outS", SOME(CL.I_Exp(CL.mkApply("fopen",  
                                                 [CL.mkStr "mip.txt",  
                                                 CL.mkStr "w"])))),  
                                                 CL.mkCall(concat[name, "_print"],  
                                                                         [CL.mkVar "outS",  
                                                                          CL.mkVar "size",  
                                                                          CL.mkVar "width",  
                                                                          CL.mkVar outStateVar])]  
   
   
   
507                  (* Body put all the statments together *)                  (* Body put all the statments together *)
508                  val body =  declarations @ [globalsDecl,initGlobalsCall] (*@ platformStm @ devicesStm *) @ contextStm @ commandStm @ !initially @ [strandSize] @                val body = CL.mkDecl(clIntTy, errVar, SOME(CL.I_Exp(CL.mkInt 0)))
509                                     strandsArrays @ globalAndlocalStms @ [widthDel,strands_init]  @ clStrandObjects @ clGlobalBuffers @ sourceStms  @ create_build_stms  (*@                      :: clGlobalBuffers @ clGlobalArguments
                                    kernelArguments @ clGlobalArguments @ enqueueStm @  [outputStm] @ freeStms @ outputData *)  
   
510                  in                  in
511                    CL.D_Func([],CL.voidTy,RN.globalsSetupName,params,CL.mkBlock(body))
                 CL.D_Func([],CL.voidTy,RN.setupFName,params,CL.mkBlock(body))  
   
512                  end                  end
513    
514  (* generate the data and global parameters *)  (* generate the data and global parameters *)
515          fun genKeneralGlobalParams ((name,tyname)::rest) =          fun genKeneralGlobalParams ((name,tyname)::rest) =
516                  CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), concat[RN.globalsVarName]) ::                globalParam (CL.T_Ptr(CL.voidTy), RN.addBufferSuffixData name) ::
517                  CL.PARAM([], CL.T_Ptr(CL.T_Named (RN.imageTy tyname)),RN.addBufferSuffix name) ::                genKeneralGlobalParams rest
518                  CL.PARAM([], CL.T_Ptr(CL.voidTy),RN.addBufferSuffixData name) ::            | genKeneralGlobalParams [] = []
                 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([]) = []  
                 in  
                   initGlobalStruct(globals) @ initGlobalImages(imgGlobals)  
                 end  
519    
520          (* generate the main kernel function for the .cl file *)          (* generate the main kernel function for the .cl file *)
521          fun genKernelFun(Strand{name, tyName, state, output, code,...},nDims,globals,imgGlobals) = let          fun genKernelFun (strand, nDims, globals, imgGlobals) = let
522                  val Strand{name, tyName, state, output, code,...} = strand
523                   val fName = RN.kernelFuncName;                   val fName = RN.kernelFuncName;
524                   val inState = "strand_in"                   val inState = "strand_in"
525                   val outState = "strand_out"                   val outState = "strand_out"
526                  val tempVar = "tmp"
527               val params = [               val params = [
528                        CL.PARAM(["__global"], CL.T_Ptr(CL.T_Named tyName), "selfIn"),                        globalParam(CL.T_Ptr(CL.T_Named tyName), "selfIn"),
529                        CL.PARAM(["__global"], CL.T_Ptr(CL.T_Named tyName), "selfOut"),                        globalParam(CL.T_Ptr(CL.T_Named tyName), "selfOut"),
530                        CL.PARAM(["__global"], CL.intTy, "width")                        globalParam(CL.T_Ptr(CL.T_Num(RawTypes.RT_UInt8)), "strandStatus"),
531                          CL.PARAM([], CL.intTy, "width"),
532                          globalParam(globPtrTy, RN.globalsVarName)
533                      ] @ genKeneralGlobalParams(!imgGlobals)                      ] @ genKeneralGlobalParams(!imgGlobals)
534                    val thread_ids = if nDims = 1                    val thread_ids = if nDims = 1
535                          then [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),                        then [
536                                    CL.mkAssign(CL.mkVar "x",CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(0,CL.intTy)]))]                            CL.mkDecl(CL.intTy, "x",
537                                SOME(CL.I_Exp(CL.mkApply(RN.getGlobalThreadId,[CL.mkInt 0]))))
538                            ]
539                        else if nDims = 2
540                          then [
541                              CL.mkDecl(CL.intTy, "x",
542                                SOME(CL.I_Exp(CL.mkApply(RN.getGlobalThreadId,[CL.mkInt 0])))),
543                              CL.mkDecl(CL.intTy, "y",
544                                SOME(CL.I_Exp(CL.mkApply(RN.getGlobalThreadId,[CL.mkInt 1]))))
545                            ]
546                        else raise Fail "nDims > 2"
547                  val strandDecl = [
548                          CL.mkAttrDecl(["__global"], CL.T_Ptr(CL.T_Named tyName), inState, NONE),
549                          CL.mkAttrDecl(["__global"], CL.T_Ptr(CL.T_Named tyName), outState, NONE),
550                          CL.mkAttrDecl(["__global"], CL.T_Ptr(CL.T_Named tyName), tempVar, NONE)
551                        ]
552                  val imageDataDecl = CL.mkDecl(CL.T_Named(RN.imageDataType),RN.globalImageDataName,NONE)
553                  val imageDataStms = List.map (fn (x,_) =>
554                      CL.mkAssign(CL.mkSelect(CL.mkVar(RN.globalImageDataName),RN.imageDataName x),
555                                  CL.mkVar(RN.addBufferSuffixData x))) (!imgGlobals)
556                  val barrierCode = CL.mkIfThen(CL.mkBinOp(CL.E_Var "status",CL.#==,CL.E_Var "DIDEROT_ACTIVE"),
557                                     CL.mkBlock ([CL.mkAssign(CL.E_Var tempVar, CL.E_Var inState),
558                                     CL.mkAssign(CL.E_Var inState, CL.E_Var outState),
559                                     CL.mkAssign(CL.E_Var outState, CL.E_Var tempVar)]))
560                  val barrierStm = CL.mkCall("barrier",[CL.E_Var "CLK_LOCAL_MEM_FENCE"])
561                  val index = if nDims = 1 then
562                            CL.mkStr "x"
563                          else                          else
564                                  [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),                          CL.mkBinOp(
565                                   CL.mkDecl(CL.intTy, "y", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),                              CL.mkBinOp(CL.mkVar "x", CL.#*, CL.mkVar "width"), CL.#+, CL.mkVar "y")
566                                    CL.mkAssign(CL.mkVar "x",  CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(0,CL.intTy)])),  
567                                    CL.mkAssign(CL.mkVar "y",CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(1,CL.intTy)]))]                val strandObjects =
568                         [ CL.mkAssign(CL.mkVar inState,  CL.mkBinOp(CL.mkVar "selfIn",CL.#+,index)),
569                    val strandDecl = [CL.mkDecl(CL.T_Named tyName, inState, NONE),                         CL.mkAssign(CL.mkVar outState, CL.mkBinOp(CL.mkVar "selfOut",CL.#+,index))
570                                                          CL.mkDecl(CL.T_Named tyName, outState,NONE)]                       ]
571                    val strandObjects  = if nDims = 1  
572                          then [CL.mkAssign(CL.mkSubscript(CL.mkVar "selfIn",CL.mkStr "x"),                  val stabalizeStm = CL.mkAssign(CL.mkSubscript(CL.mkVar "strandStatus",index),
573                                                                           CL.mkVar inState),                                                                          CL.E_Var "status")
574                                    CL.mkAssign(CL.mkSubscript(CL.mkVar "selfOut",CL.mkStr "x"),                val status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.mkSubscript(CL.mkVar "strandStatus",index))))
575                                                                           CL.mkVar outState)]                val strandInitStm = CL.mkCall(RN.strandInit name, [
576                          else let                        CL.mkVar RN.globalsVarName,
577                                  val index = CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")                        CL.mkVar inState,
578                                  in                        CL.mkVar "x",
579                                          [CL.mkAssign(CL.mkSubscript(CL.mkVar "selfIn",index),  (* FIXME: if nDims = 1, then "y" is not defined! the arguments to this call should really come from
580                                                                          CL.mkVar inState),   * the initially code!
581                                           CL.mkAssign(CL.mkSubscript(CL.mkVar "selfOut",index),   *)
582                                                                          CL.mkVar outState)]                        CL.mkVar "y"])
583                  val local_vars = thread_ids
584                        @ [imageDataDecl]
585                        @ imageDataStms
586                        @ strandDecl
587                        @ strandObjects
588                        @ [strandInitStm,status]
589                  val while_exp = CL.mkBinOp(CL.mkVar "status",CL.#==, CL.mkVar RN.kActive)
590                  val whileBody = CL.mkBlock ([
591                          CL.mkAssign(CL.mkVar "status",
592                            CL.mkApply(RN.strandUpdate name,
593                              [CL.mkVar inState,
594                               CL.mkVar outState,
595                               CL.mkVar RN.globalsVarName,
596                               CL.mkVar RN.globalImageDataName]))] @ [barrierCode,barrierStm] )
597                  val whileBlock = [CL.mkWhile(while_exp, whileBody)]
598                  val body = CL.mkBlock(local_vars @ whileBlock @ [stabalizeStm])
599                  in
600                    CL.D_Func(["__kernel"], CL.voidTy, fName, params, body)
601                                  end                                  end
602    
603          (* generate a global structure type definition from the list of globals *)
604            fun genGlobalStruct (targetTy, globals, tyName) = let
605                  val globs = List.map (fn (x : mirror_var) => (targetTy x, #var x)) globals
606                  in
607                    CL.D_StructDef(globs, tyName)
608                  end
609    
610                    val status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy))))        (* generate a global structure type definition from the image data of the image globals *)
611                    val local_vars =  thread_ids @ initKernelGlobals(!globals,!imgGlobals)  @ strandDecl @ strandObjects @ [status]          fun genImageDataStruct (imgGlobals, tyName) = let
612                    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 globs = List.map
613                    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)])),                      (fn (x, _) => (globalPtr CL.voidTy, RN.imageDataName x))
614                                                          CL.mkCall(RN.strandStabilize name,[ CL.mkUnOp(CL.%&,CL.mkVar inState),  CL.mkUnOp(CL.%&,CL.mkVar outState)])]                        imgGlobals
615                  in
616                    CL.D_StructDef(globs, tyName)
617                  end
618    
619                    val whileBlock = [CL.mkWhile(while_exp,CL.mkBlock while_body)]          fun genGlobals (declFn, targetTy, globals) = let
620                  fun doVar (x : mirror_var) = declFn (CL.D_Var([], targetTy x, #var x, NONE))
621                  in
622                    List.app doVar globals
623                  end
624    
625                    val body = CL.mkBlock(local_vars  @ whileBlock)          fun genStrandDesc (Strand{name, output, ...}) = let
626                (* the strand's descriptor object *)
627                  val descI = let
628                        fun fnPtr (ty, f) = CL.I_Exp(CL.mkCast(CL.T_Named ty, CL.mkVar f))
629                        val SOME(outTy, _) = !output
630                  in                  in
631                     CL.D_Func(["__kernel"], CL.voidTy, fName, params, body)                        CL.I_Struct[
632                              ("name", CL.I_Exp(CL.mkStr name)),
633                              ("stateSzb", CL.I_Exp(CL.mkSizeof(CL.T_Named(N.strandTy name)))),
634    (*
635                              ("outputSzb", CL.I_Exp(CL.mkSizeof(ToC.trTy outTy))),
636    *)
637                              ("update", fnPtr("update_method_t", "0")),
638                              ("print", fnPtr("print_method_t", name ^ "Print"))
639                            ]
640                  end                  end
641          (* 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)  
642                   in                   in
643                          CL.D_StructDef(getGlobals(globals),RN.globalsTy)                  desc
644                    end                    end
645    
646        (* generate the table of strand descriptors *)        (* generate the table of strand descriptors *)
647          fun genStrandTable (ppStrm, strands) = let          fun genStrandTable (declFn, strands) = let
648                val nStrands = length strands                val nStrands = length strands
649                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)))
650                fun genInits (_, []) = []                fun genInits (_, []) = []
651                  | 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)  
652                in                in
653                  ppDecl (CL.D_Var([], CL.int32, RN.numStrands,                  declFn (CL.D_Var([], CL.int32, N.numStrands,
654                    SOME(CL.I_Exp(CL.mkInt(IntInf.fromInt nStrands, CL.int32)))));                    SOME(CL.I_Exp(CL.E_Int(IntInf.fromInt nStrands, CL.int32)))));
655                  ppDecl (CL.D_Var([],                  declFn (CL.D_Var([],
656                    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),
657                    RN.strands,                    N.strands,
658                    SOME(CL.I_Array(genInits (0, strands)))))                    SOME(CL.I_Array(genInits (0, strands)))))
659                end                end
660    
661            fun genSrc (baseName, prog) = let
662          fun genSrc (baseName, Prog{double,globals, topDecls, strands, initially,imgGlobals,numDims,...}) = let                val Prog{name,double, globals, topDecls, strands, initially, imgGlobals, numDims, ...} = prog
663                val clFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "cl"}                val clFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "cl"}
664                val cFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "c"}                val cFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "c"}
665                val clOutS = TextIO.openOut clFileName                val clOutS = TextIO.openOut clFileName
666                val cOutS = TextIO.openOut cFileName                val cOutS = TextIO.openOut cFileName
667  (* FIXME: need to use PrintAsC and PrintAsCL *)                val clppStrm = PrintAsCL.new clOutS
               val clppStrm = PrintAsC.new clOutS  
668                val cppStrm = PrintAsC.new cOutS                val cppStrm = PrintAsC.new cOutS
669                  val progName = name
670                fun cppDecl dcl = PrintAsC.output(cppStrm, dcl)                fun cppDecl dcl = PrintAsC.output(cppStrm, dcl)
671                fun clppDecl dcl = PrintAsC.output(clppStrm, dcl)                fun clppDecl dcl = PrintAsCL.output(clppStrm, dcl)
672                val strands = AtomTable.listItems strands                val strands = AtomTable.listItems strands
673                val [strand as Strand{name, tyName, code,init_code, ...}] = strands                val [strand as Strand{name, tyName, code,init_code, ...}] = strands
674                in                in
# Line 921  Line 678 
678                        then "#define DIDEROT_DOUBLE_PRECISION"                        then "#define DIDEROT_DOUBLE_PRECISION"
679                        else "#define DIDEROT_SINGLE_PRECISION",                        else "#define DIDEROT_SINGLE_PRECISION",
680                      "#define DIDEROT_TARGET_CL",                      "#define DIDEROT_TARGET_CL",
681                      "#include \"Diderot/cl-types.h\""                      "#include \"Diderot/cl-diderot.h\""
682                    ]));                    ]));
683                  List.app clppDecl (List.rev (!globals));                  clppDecl (genGlobalStruct (#gpuTy, !globals, RN.globalsTy));
684                  clppDecl (genGlobalStruct (!globals));                  clppDecl (genImageDataStruct(!imgGlobals,RN.imageDataType));
685                  clppDecl (genStrandTyDef strand);                  clppDecl (genStrandTyDef(#gpuTy, strand));
686                    clppDecl  (!init_code);
687                  List.app clppDecl (!code);                  List.app clppDecl (!code);
688                  clppDecl (genKernelFun (strand,!numDims,globals,imgGlobals));                  clppDecl (genKernelFun (strand,!numDims,globals,imgGlobals));
689                (* Generate the Host file .c *)                (* Generate the Host C file *)
690                  cppDecl (CL.D_Verbatim([                  cppDecl (CL.D_Verbatim([
691                      if double                      if double
692                        then "#define DIDEROT_DOUBLE_PRECISION"                        then "#define DIDEROT_DOUBLE_PRECISION"
# Line 936  Line 694 
694                      "#define DIDEROT_TARGET_CL",                      "#define DIDEROT_TARGET_CL",
695                      "#include \"Diderot/diderot.h\""                      "#include \"Diderot/diderot.h\""
696                    ]));                    ]));
697                  List.app cppDecl (List.rev (!globals));                  cppDecl (CL.D_Var(["static"], CL.charPtr, "ProgramName",
698                  cppDecl (genGlobalStruct (!globals));                    SOME(CL.I_Exp(CL.mkStr progName))));
699                  cppDecl (genStrandTyDef strand);                  cppDecl (genGlobalStruct (#hostTy, !globals, RN.globalsTy));
700                  cppDecl  (!init_code);                  cppDecl (genGlobalStruct (#shadowTy, !globals, RN.shadowGlobalsTy));
701                  cppDecl (genStrandInit(strand,!numDims));  (* FIXME: does this really need to be a global? *)
702                  cppDecl (genStrandPrint(strand,!numDims));                  cppDecl (CL.D_Var(["static"], globPtrTy, RN.globalsVarName, NONE));
703                  (* cppDecl (genKernelLoader());*)                  cppDecl (genStrandTyDef (#hostTy, strand));
704                    cppDecl (genStrandPrint strand);
705                  List.app cppDecl (List.rev (!topDecls));                  List.app cppDecl (List.rev (!topDecls));
706                  cppDecl (genHostSetupFunc (strand, clFileName, !numDims, initially, imgGlobals));                  cppDecl (genGlobalBuffersArgs (!globals,imgGlobals));
707                    List.app (fn strand => cppDecl (genStrandDesc strand)) strands;
708                    genStrandTable (cppDecl, strands);
709                    cppDecl (!initially);
710                  PrintAsC.close cppStrm;                  PrintAsC.close cppStrm;
711                  PrintAsC.close clppStrm;                  PrintAsCL.close clppStrm;
712                  TextIO.closeOut cOutS;                  TextIO.closeOut cOutS;
713                  TextIO.closeOut clOutS                  TextIO.closeOut clOutS
714                end                end
715    
716        (* 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.  
        *)  
717          fun generate (basename, prog as Prog{double, parallel, debug, ...}) = let          fun generate (basename, prog as Prog{double, parallel, debug, ...}) = let
718                fun condCons (true, x, xs) = x::xs                fun condCons (true, x, xs) = x::xs
719                  | condCons (false, _, xs) = xs                  | condCons (false, _, xs) = xs
# Line 1008  Line 768 
768          fun init (Strand{name, tyName, code,init_code, ...}, params, init) = let          fun init (Strand{name, tyName, code,init_code, ...}, params, init) = let
769                val fName = RN.strandInit name                val fName = RN.strandInit name
770                val params =                val params =
771                      CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfOut") ::                      globalParam (globPtrTy, RN.globalsVarName) ::
772                        globalParam (CL.T_Ptr(CL.T_Named tyName), "selfOut") ::
773                        List.map (fn (ToCL.V(ty, x)) => CL.PARAM([], ty, x)) params                        List.map (fn (ToCL.V(ty, x)) => CL.PARAM([], ty, x)) params
774                val initFn = CL.D_Func([], CL.voidTy, fName, params, init)                val initFn = CL.D_Func([], CL.voidTy, fName, params, init)
775                in                in
# Line 1019  Line 780 
780          fun method (Strand{name, tyName, code,...}, methName, body) = let          fun method (Strand{name, tyName, code,...}, methName, body) = let
781                val fName = concat[name, "_", methName]                val fName = concat[name, "_", methName]
782                val params = [                val params = [
783                        CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfIn"),                        globalParam (CL.T_Ptr(CL.T_Named tyName), "selfIn"),
784                        CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfOut")                        globalParam (CL.T_Ptr(CL.T_Named tyName), "selfOut"),
785                          globalParam (CL.T_Ptr(CL.T_Named (RN.globalsTy)), RN.globalsVarName),
786                          CL.PARAM([],CL.T_Named(RN.imageDataType),RN.globalImageDataName)
787                      ]                      ]
788                val methFn = CL.D_Func([], CL.int32, fName, params, body)                val methFn = CL.D_Func([], CL.int32, fName, params, body)
789                in                in

Legend:
Removed from v.1281  
changed lines
  Added in v.1421

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