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 1264, Thu Jun 2 06:08:07 2011 UTC revision 1425, Fri Jul 1 18:05:07 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 12  Line 12 
12      structure Ty = IL.Ty      structure Ty = IL.Ty
13      structure CL = CLang      structure CL = CLang
14      structure RN = RuntimeNames      structure RN = RuntimeNames
15      structure ToC = TreeToCL      structure ToCL = TreeToCL
16        structure N = CNames
17    
18      type var = ToC.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 *)
89        val clIntTy = CL.T_Named "cl_int"
90        val clProgramTy = CL.T_Named "cl_program"
91        val clKernelTy  = CL.T_Named "cl_kernel"
92        val clCmdQueueTy = CL.T_Named "cl_command_queue"
93        val clContextTy = CL.T_Named "cl_context"
94        val clDeviceIdTy = CL.T_Named "cl_device_id"
95        val clPlatformIdTy = CL.T_Named "cl_platform_id"
96        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
115        }        }
116    
117      datatype program = Prog of {      datatype program = Prog of {
118            name : string,                  (* stem of source file *)
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 56  Line 145 
145        | StrandScope of TreeIL.var list  (* strand initialization *)        | StrandScope of TreeIL.var list  (* strand initialization *)
146        | MethodScope of TreeIL.var list  (* method body; vars are state variables *)        | MethodScope of TreeIL.var list  (* method body; vars are state variables *)
147    
148    (* the supprted widths of vectors of reals on the target.  For the GNU vector extensions,    (* the supprted widths of vectors of reals on the target. *)
149     * the supported sizes are powers of two, but float2 is broken.  (* FIXME: for OpenCL 1.1, 3 is also valid *)
150     * NOTE: we should also consider the AVX vector hardware, which has 256-bit registers.      fun vectorWidths () = [2, 4, 8, 16]
    *)  
     fun vectorWidths () = if !RuntimeNames.doublePrecision  
           then [2, 4, 8]  
           else [4, 8]  
151    
152    (* tests for whether various expression forms can appear inline *)    (* tests for whether various expression forms can appear inline *)
153      fun inlineCons n = (n < 2)          (* vectors are inline, but not matrices *)      fun inlineCons n = (n < 2)          (* vectors are inline, but not matrices *)
# Line 72  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) = ToC.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) => ToC.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 => ToC.trBlock (vMap, saveState "StrandScope" stateVars, blk)                  case scope
180                  | MethodScope stateVars => ToC.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                     of StrandScope stateVars =>
182                          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)                  | _ => ToC.trBlock (vMap, fn (_, _, stm) => [stm], blk)
187                (* end case *))                  (* end case *)
188          fun exp (ENV{vMap, ...}, e) = ToC.trExp(vMap, e)                end
189            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 name (ToC.V(_, name)) = name          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
203           fun global (Prog{globals,imgGlobals, ...}, name, ty) = let           fun global (Prog{globals,imgGlobals, ...}, name, ty) = let
204                val ty' = ToC.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               ToC.V(ty', name)                  isImgGlobal (ty, name);
211                    ToCL.V(#gpuTy x, name)
212                end                end
213          fun param x = ToC.V(ToC.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' = ToC.trType(V.ty x)                val ty = V.ty x
216                val x' = ToC.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 141  Line 246 
246    (* programs *)    (* programs *)
247      structure Program =      structure Program =
248        struct        struct
249          fun new {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,
254                    double = double, parallel = parallel, debug = debug,                    double = double, parallel = parallel, debug = debug,
255                    globals = ref [                    globals = ref [],
                     CL.D_Verbatim[  
                         if double  
                           then "#define DIDEROT_DOUBLE_PRECISION"  
                           else "#define DIDEROT_SINGLE_PRECISION",  
                         "#include \"Diderot/opencl_types.h\""  
                       ]],  
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.E_Var RN.globalsVarName,globalVar),CL.E_Var 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 178  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 208  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 216  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, (ToC.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.E_Int(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.E_Int(IntInf.fromInt nDims, CL.int32))))                        CL.mkDecl(worldTy, wrld,
327                            SOME(CL.I_Exp(CL.E_Apply(RN.allocInitially, [
328                                CL.mkVar "ProgramName",
329                                CL.mkUnOp(CL.%&, CL.E_Var(N.strandDesc name)),
330                                CL.E_Bool isArray,
331                                CL.E_Int(IntInf.fromInt nDims, CL.int32),
332                                CL.E_Var "base",
333                                CL.E_Var "size"
334                              ]))))
335                            ]                            ]
336                (* create the loop nest for the initially iterations
337                                    val numStrandsLoopBody = CL.mkExpStm(CL.mkAssignOp(CL.E_Var numStrandsVar, CL.*=,CL.mkSubscript(CL.E_Var "size",CL.E_Var "i")))                val indexVar = "ix"
338                  val strandTy = CL.T_Ptr(CL.T_Named(N.strandTy name))
339                  fun mkLoopNest [] = CL.mkBlock(createPrefix @ [
340                                    val numStrandsLoop =  CL.mkFor([(CL.intTy, "i", CL.E_Int(0,CL.intTy))],                        CL.mkDecl(strandTy, "sp",
341                                                           CL.mkBinOp(CL.E_Var "i", CL.#<, CL.E_Var "numDims"),                          SOME(CL.I_Exp(
342                                                           [CL.mkPostOp(CL.E_Var "i", CL.^++)], numStrandsLoopBody)                            CL.E_Cast(strandTy,
343                              CL.E_Apply(N.inState, [CL.E_Var "wrld", CL.E_Var indexVar]))))),
344                          CL.mkCall(N.strandInit name,
345                            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                        ])
348                    | mkLoopNest ((CL.V(ty, param), lo, hi)::iters) = let
349                        val body = mkLoopNest iters
350                        in
351                          CL.mkFor(
352                            [(ty, param, lo)],
353                            CL.mkBinOp(CL.E_Var param, CL.#<=, hi),
354                            [CL.mkPostOp(CL.E_Var param, CL.^++)],
355                            body)
356                        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                                  numDims := nDims;                                  numDims := nDims;
369                                  initially := allocCode @ [numStrandsLoop]                  initially := initFn
   
370                    end                    end
371    
   
372        (***** OUTPUT *****)        (***** OUTPUT *****)
373          fun genStrandPrint (Strand{name, tyName, state, output, code},nDims) = let          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.intTy), "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.E_Var "x"), x)  
                         else if nDims = 2 then  
                                 CL.mkSelect(CL.mkSubscript(CL.mkVar "self",  
                                    CL.mkBinOp(CL.mkBinOp(CL.E_Var "x",CL.#*,CL.E_Var "width"),CL.#+,CL.E_Var "y")), x)  
   
                         else CL.mkSelect(CL.mkVar "self",x)  
   
385                      val prArgs = (case ty                      val prArgs = (case ty
386                             of Ty.IVecTy 1 => [CL.E_Str(!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.E_Str(                                  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 => ToC.ivecIndex(outState, d, i))                                  val args = List.tabulate (d, fn i => ToC.ivecIndex(outState, d, i))
392                                  in                                  in
393                                    fmt :: args                                    fmt :: args
394                                  end                                  end
395                              | Ty.TensorTy[] => [CL.E_Str "%f\n", outState]                              | Ty.TensorTy[] => [CL.mkStr "%f\n", outState]
396                              | Ty.TensorTy[d] => let                              | Ty.TensorTy[d] => let
397                                  val fmt = CL.E_Str(                                  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 => ToC.vecIndex(outState, d, i))                                  val args = List.tabulate (d, fn i => ToC.vecIndex(outState, d, i))
# Line 283  Line 403 
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.E_Int(0,CL.intTy))],  
                                                 CL.mkBinOp(CL.E_Var param, CL.#<=, CL.mkSubscript(CL.E_Var "sizes",CL.E_Int(count,CL.intTy))),  
                                                 [CL.mkPostOp(CL.E_Var 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 ToC.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 globalBufferDecl =  CL.mkDecl(CL.clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)                val globalBuffErr = "error creating OpenCL global buffer"
430                  val globalBuffer = CL.mkAssign(CL.E_Var(concat[RN.globalsVarName,"_cl"]), CL.mkApply("clCreateBuffer",                fun errorFn msg = CL.mkIfThen(CL.mkBinOp(CL.E_Var errVar, CL.#!=, CL.E_Var "CL_SUCCESS"),
431                                                                  [CL.E_Var contextVar,                      CL.mkBlock([CL.mkCall("fprintf",[CL.E_Var "stderr", CL.E_Str msg]),
432                                                                  CL.E_Var "CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR | CL_MEM_COPY_HOST_PTR",                      CL.mkCall("exit",[CL.mkInt 1])]))
433                                                                  CL.mkApply("sizeof",[CL.E_Var RN.globalsTy]),                val shadowTypeDecl =
434                                                                  CL.E_Var RN.globalsVarName,                      CL.mkDecl(CL.T_Named(RN.shadowGlobalsTy), RN.shadowGlaobalsName, NONE)
435                                                                  CL.E_UnOp(CL.%&,CL.E_Var errVar)]))                val globalToShadowStms = List.map (fn (x:mirror_var) => #hToS x ) globals
436                  val globalBufferDecl = CL.mkDecl(clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)
437                  val globalBuffer = CL.mkAssign(CL.mkVar(concat[RN.globalsVarName,"_cl"]),
438                        CL.mkApply("clCreateBuffer", [
439                            CL.mkVar contextVar,
440                            CL.mkBinOp(CL.mkVar "CL_MEM_READ_ONLY", CL.#|, CL.mkVar "CL_MEM_COPY_HOST_PTR"),
441                            CL.mkSizeof(CL.T_Named RN.shadowGlobalsTy),
442                            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")
449                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.E_Var "float"]), CL.#*,                      in
450                                           CL.mkIndirect(CL.E_Var var, "size[0]"))                        CL.mkDecl(clMemoryTy, RN.addBufferSuffixData var ,NONE) ::
451                                          else if nDims = 2 then                        CL.mkAssign(CL.mkVar(RN.addBufferSuffixData var),
452                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.E_Var "float"]), CL.#*,                          CL.mkApply("clCreateBuffer", [
453                                            CL.mkIndirect(CL.E_Var var, concat["size[0]", " * ", var, "->size[1]"]))                              CL.mkVar contextVar,
454                                          else                              CL.mkVar "CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR",
                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.E_Var "float"]), CL.#*,  
                                           CL.mkIndirect(CL.E_Var var,concat["size[0]", " * ", var, "->size[1] * ", var, "->size[2]"]))  
   
                  in  
                    CL.mkDecl(CL.clMemoryTy,RN.addBufferSuffix var ,NONE)::  
                    CL.mkDecl(CL.clMemoryTy,RN.addBufferSuffixData var ,NONE)::  
                    CL.mkAssign(CL.E_Var(RN.addBufferSuffix var), CL.mkApply("clCreateBuffer",  
                                                                 [CL.E_Var contextVar,  
                                                                 CL.E_Var "CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR | CL_MEM_COPY_HOST_PTR",  
                                                                 CL.mkApply("sizeof",[CL.E_Var (RN.imageTy nDims)]),  
                                                                 CL.E_Var var,  
                                                                 CL.E_UnOp(CL.%&,CL.E_Var errVar)])) ::  
                         CL.mkAssign(CL.E_Var(RN.addBufferSuffixData var), CL.mkApply("clCreateBuffer",  
                                                                 [CL.E_Var contextVar,  
                                                                 CL.E_Var "CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR | CL_MEM_COPY_HOST_PTR",  
455                                                                  size,                                                                  size,
456                                                                  CL.mkIndirect(CL.E_Var var,"data"),                              CL.mkIndirect(hostVar, "data"),
457                                                                  CL.E_UnOp(CL.%&,CL.E_Var 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.E_Var errVar,CL.|=,CL.mkApply("clSetKernelArg",                val globalArgErr = "error creating OpenCL global argument"
471                                                                  [CL.E_Var kernelVar,                fun errorFn msg = CL.mkIfThen(CL.mkBinOp(CL.E_Var errVar, CL.#!=, CL.E_Var "CL_SUCCESS"),
472                                                                   CL.E_Int(count,CL.intTy),                      CL.mkBlock([CL.mkCall("fprintf",[CL.E_Var "stderr", CL.E_Str msg]),
473                                                                   CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),                      CL.mkCall("exit",[CL.mkInt 1])]))
474                                                                   CL.E_UnOp(CL.%&,CL.E_Var(concat[RN.globalsVarName,"_cl"]))])))                val globalArgument = CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.&=,
475                        CL.mkApply("clSetKernelArg",
476          fun genDataArguments([],_,_,_) = []                        [CL.mkVar kernelVar,
477            | genDataArguments((var,nDims)::globals,count,kernelVar,errVar) =                         CL.mkPostOp(CL.E_Var count, CL.^++),
478                           CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
479                  CL.mkExpStm(CL.mkAssignOp(CL.E_Var errVar,CL.|=, CL.mkApply("clSetKernelArg",                         CL.mkUnOp(CL.%&,CL.mkVar(concat[RN.globalsVarName,"_cl"]))])))
480                                                                  [CL.E_Var kernelVar,                fun genDataArguments ([],_,_,_,_) = []
481                                                                   CL.E_Int(count,CL.intTy),                  | genDataArguments ((var,nDims)::globals,count,kernelVar,errVar,errFn) =
482                                                                   CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),                      CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.$=,
483                                                                   CL.E_UnOp(CL.%&,CL.E_Var(RN.addBufferSuffix var))])))::                        CL.mkApply("clSetKernelArg",
484                            [CL.mkVar kernelVar,
485                          CL.mkExpStm(CL.mkAssignOp(CL.E_Var errVar,CL.|=,CL.mkApply("clSetKernelArg",                           CL.mkPostOp(CL.E_Var count, CL.^++),
486                                                                  [CL.E_Var kernelVar,                           CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
487                                                                   CL.E_Int((count + 1),CL.intTy),                           CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffixData var))]))) ::
488                                                                   CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),                           errFn(concat["error in creating ",RN.addBufferSuffixData var, " argument"]) ::
489                                                                   CL.E_UnOp(CL.%&,CL.E_Var(RN.addBufferSuffixData var))]))):: genDataArguments (globals, count + 2,kernelVar,errVar)                      genDataArguments (globals,count,kernelVar,errVar,errFn)
   
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.E_Var 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.E_Var RN.globalsTy])]))))  
                 val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.E_Var RN.globalsVarName])  
                 val returnStm = [CL.mkReturn(SOME(CL.E_Int(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.E_Var errVar, CL.#==, CL.E_Var "CL_SUCCESS")])  
499                  val params = [                  val params = [
500                           CL.PARAM([],CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName)                        CL.PARAM([],CL.T_Named("cl_context"), "context"),
501                          CL.PARAM([],CL.T_Named("cl_kernel"), "kernel"),
502                          CL.PARAM([],CL.T_Named("cl_command_queue"), "cmdQ"),
503                          CL.PARAM([],CL.T_Named("int"), "argStart")
504                           ]                           ]
505                  val delcarations = [CL.mkDecl(CL.clProgramTy, programVar, NONE),                val clGlobalBuffers = getGlobalDataBuffers(globals,!imgGlobals, "context", errVar)
506                            CL.mkDecl(CL.clKernelTy, kernelVar, NONE),                val clGlobalArguments = genGlobalArguments(!imgGlobals, "argStart", "kernel", errVar)
                           CL.mkDecl(CL.clCmdQueueTy, cmdVar, NONE),  
                           CL.mkDecl(CL.clContextTy, contextVar, NONE),  
                           CL.mkDecl(CL.intTy, errVar, NONE),  
                           CL.mkDecl(CL.intTy, numStrandsVar, NONE),  
                           CL.mkDecl(CL.intTy, stateSizeVar, NONE),  
                           CL.mkDecl(CL.intTy, "width", NONE),  
                           CL.mkDecl(CL.intTy, imgDataSizeVar, NONE),  
                           CL.mkDecl(CL.clDeviceIdTy, deviceVar, NONE),  
                           CL.mkDecl(CL.T_Ptr(CL.T_Named tyName), inStateVar,NONE),  
                           CL.mkDecl(CL.clMemoryTy,clInstateVar,NONE),  
                           CL.mkDecl(CL.clMemoryTy,clOutStateVar,NONE),  
                           CL.mkDecl(CL.T_Ptr(CL.T_Named tyName), outStateVar,NONE),  
                           CL.mkDecl(CL.charPtr, clFNVar,SOME(CL.I_Exp(CL.E_Str filename))),  
                           CL.mkDecl(CL.charPtr, headerFNVar,SOME(CL.I_Exp(CL.E_Str "Diderot/opencl_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.E_Int(~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.E_Int(~1,CL.intTy))))]  
   
                 (* Retrieve the platforms *)  
                 val platformStm = [CL.mkAssign(CL.E_Var errVar, CL.mkApply("clGetPlatformIDs",  
                                                   [CL.E_Int(10,CL.intTy),  
                                                    CL.E_Var platformsVar,  
                                                    CL.E_UnOp(CL.%&,CL.E_Var numPlatformsVar)])),  
                                                    assertStm]  
   
                 val devicesStm = [CL.mkAssign(CL.E_Var errVar, CL.mkApply("clGetDeviceIDs",  
                                                   [CL.mkSubscript(CL.E_Var platformsVar,CL.E_Int(0,CL.intTy)),  
                                                    CL.E_Var "CL_DEVICE_TYPE_GPU",  
                                                    CL.E_Int(1,CL.intTy),  
                                                    CL.E_UnOp(CL.%&,CL.E_Var deviceVar),  
                                                    CL.E_UnOp(CL.%&,CL.E_Var numDevicesVar)])),  
                                                    assertStm]  
   
                 (* Create Context *)  
                 val contextStm = [CL.mkAssign(CL.E_Var contextVar, CL.mkApply("clCreateContext",  
                                                   [CL.E_Int(0,CL.intTy),  
                                                   CL.E_Int(1,CL.intTy),  
                                                   CL.E_UnOp(CL.%&,CL.E_Var deviceVar),  
                                                   CL.E_Var "NULL",  
                                                   CL.E_Var "NULL",  
                                                   CL.E_UnOp(CL.%&,CL.E_Var errVar)])),  
                                                   assertStm]  
   
                 (* Create Command Queue *)  
                 val commandStm = [CL.mkAssign(CL.E_Var cmdVar, CL.mkApply("clCreateCommandQueue",  
                                                   [CL.E_Var contextVar,  
                                                   CL.E_Var deviceVar,  
                                                   CL.E_Int(0,CL.intTy),  
                                                   CL.E_UnOp(CL.%&,CL.E_Var errVar)])),  
                                                   assertStm]  
   
   
                 (*Create Program/Build/Kernel with Source statement *)  
                 val createProgStm = CL.mkAssign(CL.E_Var programVar, CL.mkApply("clCreateProgramWithSource",  
                                                                                                                 [CL.E_Var contextVar,  
                                                                                                                  CL.E_Int(2,CL.intTy),  
                                                                                                                  CL.E_Cast(CL.T_Ptr(CL.T_Named("const char *")),CL.E_UnOp(CL.%&,CL.E_Var sourcesVar)),  
                                                                                                                  CL.E_Var "NULL",  
                                                                                                                  CL.E_UnOp(CL.%&,CL.E_Var errVar)]))  
   
                 (* FIXME: Remove after testing purposes, Build Log for OpenCL*)  
                 val buildLog = [CL.mkAssign(CL.E_Var errVar, CL.mkApply("clBuildProgram",  
                                                                                                                 [CL.E_Var programVar,  
                                                                                                                  CL.E_Int(0,CL.intTy),  
                                                                                                                  CL.E_Var "NULL",  
                                                                                                                  CL.E_Var "NULL",  
                                                                                                                  CL.E_Var "NULL",  
                                                                                                                  CL.E_Var "NULL"])),  
                                           CL.mkDecl(CL.T_Array(CL.charTy,SOME(2048)), "build", NONE),  
                                            CL.mkAssign(CL.E_Var errVar, CL.mkApply("clGetProgramBuildInfo",  
                                                                                                                 [CL.E_Var programVar,  
                                                                                                                 CL.E_Var deviceVar,  
                                                                                                                  CL.E_Var "CL_PROGRAM_BUILD_LOG",  
                                                                                                                  CL.E_Int (2048,CL.intTy),  
                                                                                                                  CL.E_Var "build",  
                                                                                                                  CL.E_Var "NULL"])),  
                                                 CL.mkCall("printf",[CL.E_Str ( "Build Log:" ^ "\n" ^ "%s" ^ "\n"), CL.E_Var "build"])]  
   
   
   
   
                 val createKernel = CL.mkAssign(CL.E_Var kernelVar, CL.mkApply("clCreateKernel",  
                                                                                                                 [CL.E_Var programVar,  
                                                                                                                  CL.E_Str RN.kernelFuncName,  
                                                                                                                  CL.E_UnOp(CL.%&,CL.E_Var errVar)]))  
   
   
                 val create_build_stms = [createProgStm,assertStm] @ buildLog @ [assertStm,createKernel,assertStm]  
   
   
   
                 (* Create Memory Buffers for Strand States and Globals *)  
                 val strandSize = CL.mkAssign(CL.E_Var stateSizeVar,CL.mkBinOp(CL.mkApply("sizeof",  
                                                                         [CL.E_Var tyName]), CL.#*,CL.E_Var numStrandsVar))  
                 val strandObjects = [CL.mkAssign(CL.E_Var inStateVar, CL.mkApply("malloc",  
                                                                                 [CL.E_Var stateSizeVar])),  
                                                         CL.mkAssign(CL.E_Var outStateVar, CL.mkApply("malloc",  
                                                                                 [CL.E_Var stateSizeVar]))]  
   
                 val clStrandObjects = [CL.mkAssign(CL.E_Var clInstateVar, CL.mkApply("clCreateBuffer",  
                                                                 [CL.E_Var contextVar,  
                                                                 CL.E_Var "CL_MEM_READ_WRITE",  
                                                                 CL.E_Var stateSizeVar,  
                                                                 CL.E_Var "NULL",  
                                                                 CL.E_UnOp(CL.%&,CL.E_Var errVar)])),  
                                                          CL.mkAssign(CL.E_Var clOutStateVar, CL.mkApply("clCreateBuffer",  
                                                                 [CL.E_Var contextVar,  
                                                                 CL.E_Var "CL_MEM_READ_WRITE",  
                                                                 CL.E_Var stateSizeVar,  
                                                                 CL.E_Var "NULL",  
                                                                 CL.E_UnOp(CL.%&,CL.E_Var errVar)]))]  
   
             val clGlobalBuffers = getGlobalDataBuffers(!imgGlobals,3,contextVar,errVar)  
   
   
                 (* Load the Kernel and Header Files *)  
                 val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.E_Var sourcesVar,CL.E_Int(0,CL.intTy)),  
                                                                           CL.mkApply(RN.clLoaderFN, [CL.E_Var clFNVar])),  
                                                   CL.mkAssign(CL.mkSubscript(CL.E_Var sourcesVar,CL.E_Int(1,CL.intTy)),  
                                                                           CL.mkApply(RN.clLoaderFN, [CL.E_Var headerFNVar]))]  
   
                 (* Created Enqueue Statements *)  
 (* FIXME: simplify this code by function abstraction *)  
         val enqueueStm = if nDims = 1  
                         then [CL.mkAssign(CL.E_Var errVar,  
                                                           CL.mkApply("clEnqueueNDRangeKernel",  
                                                                                                 [CL.E_Var cmdVar,  
                                                                                                  CL.E_Var kernelVar,  
                                                                                                  CL.E_Int(1,CL.intTy),  
                                                                                                  CL.E_Var "NULL",  
                                                                                                  CL.E_Var globalVar,  
                                                                                                  CL.E_Var localVar,  
                                                                                                  CL.E_Int(0,CL.intTy),  
                                                                                                  CL.E_Var "NULL",  
                                                                                                  CL.E_Var "NULL"])),CL.mkCall("clFinish",[CL.E_Var cmdVar])]  
                         else if nDims = 2  then  
                          [CL.mkAssign(CL.E_Var errVar,  
                                                         CL.mkApply("clEnqueueNDRangeKernel",  
                                                                                                 [CL.E_Var cmdVar,  
                                                                                                  CL.E_Var kernelVar,  
                                                                                                  CL.E_Int(2,CL.intTy),  
                                                                                                  CL.E_Var "NULL",  
                                                                                                  CL.E_Var globalVar,  
                                                                                                  CL.E_Var localVar,  
                                                                                                  CL.E_Int(0,CL.intTy),  
                                                                                                  CL.E_Var "NULL",  
                                                                                                  CL.E_Var "NULL"])),CL.mkCall("clFinish",[CL.E_Var cmdVar])]  
                         else  
                           [CL.mkAssign(CL.E_Var errVar,  
                                                         CL.mkApply("clEnqueueNDRangeKernel",  
                                                                                                 [CL.E_Var cmdVar,  
                                                                                                  CL.E_Var kernelVar,  
                                                                                                  CL.E_Int(3,CL.intTy),  
                                                                                                  CL.E_Var "NULL",  
                                                                                                  CL.E_Var globalVar,  
                                                                                                  CL.E_Var localVar,  
                                                                                                  CL.E_Int(0,CL.intTy),  
                                                                                                  CL.E_Var "NULL",  
                                                                                                  CL.E_Var "NULL"])),CL.mkCall("clFinish",[CL.E_Var cmdVar])]  
   
                 (* Setup up selfOut variable *)  
                 val selfOutStm = CL.mkAssign(CL.E_Var outStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.E_Var numStrandsVar,  
                                                                         CL.#*, CL.mkApply("sizeof",[CL.E_Var tyName]))]))  
   
                 (* Initialize Width Parameter *)  
                 val widthDel = if nDims = 2 then  
                           CL.mkAssign(CL.E_Var "width",CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(1, CL.intTy)))  
                    else  
                           CL.mkAssign(CL.E_Var "width",CL.E_Int(0,CL.intTy))  
   
                 (* Setup Global and Local variables *)  
   
                 val globalAndlocalStms = if nDims = 1 then  
                         [CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.E_Var "size", CL.E_Int(0,CL.intTy))),  
                          CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(0,CL.intTy)),  
                                                                   CL.E_Var "16")]  
   
   
                 else if nDims = 2 then  
                         [CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.E_Var "size", CL.E_Int(0,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(1,CL.intTy)),  
                                                                    CL.mkSubscript(CL.E_Var "size", CL.E_Int(1,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(0,CL.intTy)),  
                                                                   CL.E_Var "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(1,CL.intTy)),  
                                                                   CL.E_Var "16")]  
   
                 else  
                         [CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.E_Var "size", CL.E_Int(0,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(1,CL.intTy)),  
                                                                    CL.mkSubscript(CL.E_Var "size", CL.E_Int(1,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(2,CL.intTy)),  
                                                                    CL.mkSubscript(CL.E_Var "size", CL.E_Int(2,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(0,CL.intTy)),  
                                                                   CL.E_Var "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(1,CL.intTy)),  
                                                                   CL.E_Var "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(2,CL.intTy)),  
                                                                   CL.E_Var "16")]  
   
   
   
                 (* Setup Kernel arguments *)  
                 val kernelArguments = [CL.mkAssign(CL.E_Var errVar,CL.mkApply("clSetKernelArg",  
                                                                 [CL.E_Var kernelVar,  
                                                                  CL.E_Int(0,CL.intTy),  
                                                                  CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),  
                                                                  CL.E_UnOp(CL.%&,CL.E_Var clInstateVar)])),  
                                                             CL.mkExpStm(CL.mkAssignOp(CL.E_Var errVar, CL.|=,CL.mkApply("clSetKernelArg",  
                                                                 [CL.E_Var kernelVar,  
                                                                  CL.E_Int(1,CL.intTy),  
                                                                  CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),  
                                                                  CL.E_UnOp(CL.%&,CL.E_Var clOutStateVar)]))),  
                                                                   CL.mkExpStm(CL.mkAssignOp(CL.E_Var errVar, CL.|=,CL.mkApply("clSetKernelArg",  
                                                                 [CL.E_Var kernelVar,  
                                                                  CL.E_Int(2,CL.intTy),  
                                                                  CL.mkApply("sizeof",[CL.E_Var "int"]),  
                                                                  CL.E_UnOp(CL.%&,CL.E_Var "width")])))]  
   
            val clGlobalArguments = genGlobalArguments(!imgGlobals,3,kernelVar,errVar) @ [assertStm]  
   
                 (* Retrieve output *)  
                 val outputStm = CL.mkAssign(CL.E_Var errVar,  
                                                         CL.mkApply("clEnqueueReadBuffer",  
                                                                                                 [CL.E_Var cmdVar,  
                                                                                                  CL.E_Var clOutStateVar,  
                                                                                                  CL.E_Var "CL_TRUE",  
                                                                                                  CL.E_Int(0,CL.intTy),  
                                                                                                  CL.E_Var stateSizeVar,  
                                                                                                  CL.E_Var outStateVar,  
                                                                                                  CL.E_Int(0,CL.intTy),  
                                                                                                  CL.E_Var "NULL",  
                                                                                                  CL.E_Var "NULL"]))  
   
                 (* Free all the objects *)  
                 val freeStms = [CL.mkCall("clReleaseKernel",[CL.E_Var kernelVar]),  
                                                 CL.mkCall("clReleaseProgram",[CL.E_Var programVar ]),  
                                                 CL.mkCall("clReleaseCommandQueue",[CL.E_Var cmdVar]),  
                                                 CL.mkCall("clReleaseContext",[CL.E_Var contextVar]),  
                                                 CL.mkCall("clReleaseMemObject",[CL.E_Var clInstateVar]),  
                                                 CL.mkCall("clReleaseMemObject",[CL.E_Var 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.E_Str "mip.txt",  
                                                 CL.E_Str "w"])))),  
                                                 CL.mkCall(concat[name, "_print"],  
                                                                         [CL.E_Var "outS",  
                                                                          CL.E_Var "size",  
                                                                          CL.E_Var "width",  
                                                                          CL.E_Var outStateVar])]  
   
   
507                  (* Body put all the statments together *)                  (* Body put all the statments together *)
508                  val body =  delcarations @ platformStm @ devicesStm @ contextStm @ commandStm @ !initially @ [strandSize] @                val body = CL.mkDecl(clIntTy, errVar, SOME(CL.I_Exp(CL.mkInt 0)))
509                                     clStrandObjects @ clGlobalBuffers @ sourceStms  @ [selfOutStm] @ create_build_stms @ globalAndlocalStms @ [widthDel] @                      :: 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(["__global"], CL.T_Ptr(CL.T_Named RN.globalsTy), concat[RN.globalsVarName]) ::                globalParam (CL.T_Ptr(CL.voidTy), RN.addBufferSuffixData name) ::
517                  CL.PARAM(["__global"], CL.T_Ptr(CL.T_Named (RN.imageTy tyname)),RN.addBufferSuffix name) ::                genKeneralGlobalParams rest
518                  CL.PARAM(["__global"], 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.E_Var name, CL.mkIndirect(CL.E_Var RN.globalsVarName, name)) ::  
                                 initGlobalStruct(rest)  
                   | initGlobalStruct ( _::rest) = initGlobalStruct(rest)  
                   | initGlobalStruct([]) = []  
   
                 fun initGlobalImages((name,tyname)::rest) =  
                                 CL.mkAssign(CL.E_Var name, CL.E_Var (RN.addBufferSuffix name)) ::  
                                 CL.mkAssign(CL.mkIndirect(CL.E_Var name,"data"),CL.E_Var (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 sizeParams = if nDims = 1 then
528                            []
529                        else if nDims = 2 then
530                            [CL.PARAM([], CL.intTy, "width")]
531                        else
532                            [CL.PARAM([], CL.intTy, "width"),CL.PARAM([], CL.intTy, "height")]
533               val params = [               val params = [
534                        CL.PARAM(["__global"], CL.T_Ptr(CL.T_Named tyName), "selfIn"),                        globalParam(CL.T_Ptr(CL.T_Named tyName), "selfIn"),
535                        CL.PARAM(["__global"], CL.T_Ptr(CL.T_Named tyName), "selfOut"),                        globalParam(CL.T_Ptr(CL.T_Named tyName), "selfOut"),
536                        CL.PARAM(["__global"], CL.intTy, "width")                        globalParam(CL.T_Ptr(CL.T_Num(RawTypes.RT_UInt8)), "strandStatus")] @
537                          sizeParams @
538                         [ globalParam(globPtrTy, RN.globalsVarName)
539                      ] @ genKeneralGlobalParams(!imgGlobals)                      ] @ genKeneralGlobalParams(!imgGlobals)
540                    val thread_ids = if nDims = 1                    val thread_ids = if nDims = 1
541                          then [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy)))),                        then [
542                                    CL.mkAssign(CL.E_Var "x",CL.mkApply(RN.getGlobalThreadId,[CL.E_Int(0,CL.intTy)]))]                            CL.mkDecl(CL.intTy, "x",
543                                SOME(CL.I_Exp(CL.mkApply(RN.getGlobalThreadId,[CL.mkInt 0]))))
544                            ]
545                        else if nDims = 2
546                          then [
547                              CL.mkDecl(CL.intTy, "x",
548                                SOME(CL.I_Exp(CL.mkApply(RN.getGlobalThreadId,[CL.mkInt 1])))),
549                              CL.mkDecl(CL.intTy, "y",
550                                SOME(CL.I_Exp(CL.mkApply(RN.getGlobalThreadId,[CL.mkInt 0]))))
551                            ]
552                          else                          else
553                                  [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy)))),                         [
554                                   CL.mkDecl(CL.intTy, "y", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy)))),                            CL.mkDecl(CL.intTy, "x",
555                                    CL.mkAssign(CL.E_Var "x",  CL.mkApply(RN.getGlobalThreadId,[CL.E_Int(0,CL.intTy)])),                              SOME(CL.I_Exp(CL.mkApply(RN.getGlobalThreadId,[CL.mkInt 1])))),
556                                    CL.mkAssign(CL.E_Var "y",CL.mkApply(RN.getGlobalThreadId,[CL.E_Int(1,CL.intTy)]))]                            CL.mkDecl(CL.intTy, "y",
557                                SOME(CL.I_Exp(CL.mkApply(RN.getGlobalThreadId,[CL.mkInt 0])))),
558                    val strandDecl = [CL.mkDecl(CL.T_Named tyName, inState, NONE),                            CL.mkDecl(CL.intTy, "z",
559                                                          CL.mkDecl(CL.T_Named tyName, outState,NONE)]                              SOME(CL.I_Exp(CL.mkApply(RN.getGlobalThreadId,[CL.mkInt 2]))))
560                    val strandObjects  = if nDims = 1                         ]
                         then [CL.mkAssign(CL.mkSubscript(CL.E_Var "selfIn",CL.E_Str "x"),  
                                                                          CL.E_Var inState),  
                                   CL.mkAssign(CL.mkSubscript(CL.E_Var "selfOut",CL.E_Str "x"),  
                                                                          CL.E_Var outState)]  
                         else let  
                                 val index = CL.mkBinOp(CL.mkBinOp(CL.E_Var "x",CL.#*,CL.E_Var "width"),CL.#+,CL.E_Var "y")  
                                 in  
                                         [CL.mkAssign(CL.mkSubscript(CL.E_Var "selfIn",index),  
                                                                         CL.E_Var inState),  
                                          CL.mkAssign(CL.mkSubscript(CL.E_Var "selfOut",index),  
                                                                         CL.E_Var outState)]  
                                 end  
561    
562    
563                    val status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy))))                val strandDecl = [
564                    val strand_init_function = CL.mkCall(RN.strandInit name, [CL.E_UnOp(CL.%&,CL.E_Var inState),CL.E_Var "x", CL.E_Var "y"])                        CL.mkAttrDecl(["__global"], CL.T_Ptr(CL.T_Named tyName), inState, NONE),
565                    val local_vars =  thread_ids @ initKernelGlobals(!globals,!imgGlobals)  @ strandDecl @ strandObjects @ [status,strand_init_function]                        CL.mkAttrDecl(["__global"], CL.T_Ptr(CL.T_Named tyName), outState, NONE),
566                    val while_exp = CL.mkBinOp(CL.mkBinOp(CL.E_Var "status",CL.#!=, CL.E_Var RN.kStabilize),CL.#||,CL.mkBinOp(CL.E_Var "status", CL.#!=, CL.E_Var RN.kDie))                        CL.mkAttrDecl(["__global"], CL.T_Ptr(CL.T_Named tyName), tempVar, NONE)
567                    val while_body = [CL.mkAssign(CL.E_Var "status", CL.mkApply(RN.strandUpdate name,[ CL.E_UnOp(CL.%&,CL.E_Var inState), CL.E_UnOp(CL.%&,CL.E_Var outState)])),                      ]
568                                                          CL.mkCall(RN.strandStabilize name,[ CL.E_UnOp(CL.%&,CL.E_Var inState),  CL.E_UnOp(CL.%&,CL.E_Var outState)])]                val imageDataDecl = CL.mkDecl(CL.T_Named(RN.imageDataType),RN.globalImageDataName,NONE)
569                  val imageDataStms = List.map (fn (x,_) =>
570                      CL.mkAssign(CL.mkSelect(CL.mkVar(RN.globalImageDataName),RN.imageDataName x),
571                                  CL.mkVar(RN.addBufferSuffixData x))) (!imgGlobals)
572                  val barrierCode = CL.mkIfThen(CL.mkBinOp(CL.E_Var "status",CL.#==,CL.E_Var "DIDEROT_ACTIVE"),
573                                     CL.mkBlock ([CL.mkAssign(CL.E_Var tempVar, CL.E_Var inState),
574                                     CL.mkAssign(CL.E_Var inState, CL.E_Var outState),
575                                     CL.mkAssign(CL.E_Var outState, CL.E_Var tempVar)]))
576                  val barrierStm = CL.mkCall("barrier",[CL.E_Var "CLK_LOCAL_MEM_FENCE"])
577                  val index = if nDims = 1 then
578                            CL.mkVar "x"
579                        else if nDims = 2 then
580                            CL.mkBinOp(
581                                CL.mkBinOp(CL.mkVar "y", CL.#*, CL.mkVar "width"), CL.#+, CL.mkVar "x")
582                        else
583                           CL.mkBinOp(CL.mkBinOp(CL.mkBinOp(
584                                CL.mkBinOp(CL.mkVar "z", CL.#*, CL.mkVar "width"),CL.#*, CL.mkVar "height"), CL.#+,
585                                CL.mkBinOp(CL.mkVar "y",CL.#*,CL.mkVar "height")),CL.#+,CL.mkVar "x")
586    
587                  val args = if nDims = 1 then
588                            [CL.mkVar "x"]
589                        else if nDims = 2 then
590                            [CL.mkVar "x", CL.mkVar "y"]
591                        else
592                            [CL.mkVar "x", CL.mkVar "y", CL.mkVar "z"]
593    
                   val whileBlock = [CL.mkWhile(while_exp,CL.mkBlock while_body)]  
594    
595                    val body = CL.mkBlock(local_vars  @ whileBlock)                val strandObjects =
596                         [ CL.mkAssign(CL.mkVar inState,  CL.mkBinOp(CL.mkVar "selfIn",CL.#+,index)),
597                           CL.mkAssign(CL.mkVar outState, CL.mkBinOp(CL.mkVar "selfOut",CL.#+,index))
598                         ]
599    
600                    val stabalizeStm = CL.mkAssign(CL.mkSubscript(CL.mkVar "strandStatus",index),
601                                                                            CL.E_Var "status")
602                  val status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.mkSubscript(CL.mkVar "strandStatus",index))))
603                  val strandInitStm = CL.mkCall(RN.strandInit name, [
604                          CL.mkVar RN.globalsVarName,
605                          CL.mkVar inState] @ args)
606                  val local_vars = thread_ids
607                        @ [imageDataDecl]
608                        @ imageDataStms
609                        @ strandDecl
610                        @ strandObjects
611                        @ [strandInitStm,status]
612                  val while_exp = CL.mkBinOp(CL.mkVar "status",CL.#==, CL.mkVar RN.kActive)
613                  val whileBody = CL.mkBlock ([
614                          CL.mkAssign(CL.mkVar "status",
615                            CL.mkApply(RN.strandUpdate name,
616                              [CL.mkVar inState,
617                               CL.mkVar outState,
618                               CL.mkVar RN.globalsVarName,
619                               CL.mkVar RN.globalImageDataName]))] @ [barrierCode,barrierStm] )
620                  val whileBlock = [CL.mkWhile(while_exp, whileBody)]
621                  val body = CL.mkBlock(local_vars @ whileBlock @ [stabalizeStm])
622                  in                  in
623                     CL.D_Func(["__kernel"], CL.voidTy, fName, params, body)                     CL.D_Func(["__kernel"], CL.voidTy, fName, params, body)
624                  end                  end
625          (* generate a global structure from the globals *)  
626          fun genGlobalStruct(globals) = let        (* generate a global structure type definition from the list of globals *)
627                   fun getGlobals(CL.D_Var(_,ty,globalVar,_)::rest) = (ty,globalVar)::getGlobals(rest)          fun genGlobalStruct (targetTy, globals, tyName) = let
628                     | getGlobals([]) = []                val globs = List.map (fn (x : mirror_var) => (targetTy x, #var x)) globals
629                     | getGlobals(_::rest) = getGlobals(rest)                in
630                    CL.D_StructDef(globs, tyName)
631                  end
632    
633          (* generate a global structure type definition from the image data of the image globals *)
634            fun genImageDataStruct (imgGlobals, tyName) = let
635                  val globs = List.map
636                        (fn (x, _) => (globalPtr CL.voidTy, RN.imageDataName x))
637                          imgGlobals
638                  in
639                    CL.D_StructDef(globs, tyName)
640                  end
641    
642            fun genGlobals (declFn, targetTy, globals) = let
643                  fun doVar (x : mirror_var) = declFn (CL.D_Var([], targetTy x, #var x, NONE))
644                  in
645                    List.app doVar globals
646                  end
647    
648            fun genStrandDesc (Strand{name, output, ...}) = let
649                (* the strand's descriptor object *)
650                  val descI = let
651                        fun fnPtr (ty, f) = CL.I_Exp(CL.mkCast(CL.T_Named ty, CL.mkVar f))
652                        val SOME(outTy, _) = !output
653                        in
654                          CL.I_Struct[
655                              ("name", CL.I_Exp(CL.mkStr name)),
656                              ("stateSzb", CL.I_Exp(CL.mkSizeof(CL.T_Named(N.strandTy name)))),
657    (*
658                              ("outputSzb", CL.I_Exp(CL.mkSizeof(ToC.trTy outTy))),
659    *)
660                              ("update", fnPtr("update_method_t", "0")),
661                              ("print", fnPtr("print_method_t", name ^ "Print"))
662                            ]
663                        end
664                  val desc = CL.D_Var([], CL.T_Named N.strandDescTy, N.strandDesc name, SOME descI)
665                   in                   in
666                          CL.D_StructDef(getGlobals(globals),RN.globalsTy)                  desc
667                    end                    end
668    
669        (* generate the table of strand descriptors *)        (* generate the table of strand descriptors *)
670          fun genStrandTable (ppStrm, strands) = let          fun genStrandTable (declFn, strands) = let
671                val nStrands = length strands                val nStrands = length strands
672                fun genInit (Strand{name, ...}) = CL.I_Exp(CL.mkUnOp(CL.%&, CL.E_Var(RN.strandDesc name)))                fun genInit (Strand{name, ...}) = CL.I_Exp(CL.mkUnOp(CL.%&, CL.E_Var(N.strandDesc name)))
673                fun genInits (_, []) = []                fun genInits (_, []) = []
674                  | 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)  
675                in                in
676                  ppDecl (CL.D_Var([], CL.int32, RN.numStrands,                  declFn (CL.D_Var([], CL.int32, N.numStrands,
677                    SOME(CL.I_Exp(CL.E_Int(IntInf.fromInt nStrands, CL.int32)))));                    SOME(CL.I_Exp(CL.E_Int(IntInf.fromInt nStrands, CL.int32)))));
678                  ppDecl (CL.D_Var([],                  declFn (CL.D_Var([],
679                    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),
680                    RN.strands,                    N.strands,
681                    SOME(CL.I_Array(genInits (0, strands)))))                    SOME(CL.I_Array(genInits (0, strands)))))
682                end                end
683    
684            fun genSrc (baseName, prog) = let
685          fun genSrc (baseName, Prog{double,globals, topDecls, strands, initially,imgGlobals,numDims,...}) = let                val Prog{name,double, globals, topDecls, strands, initially, imgGlobals, numDims, ...} = prog
686                val clFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "cl"}                val clFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "cl"}
687                val cFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "c"}                val cFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "c"}
688                val clOutS = TextIO.openOut clFileName                val clOutS = TextIO.openOut clFileName
689                val cOutS = TextIO.openOut cFileName                val cOutS = TextIO.openOut cFileName
690                val clppStrm = PrintAsC.new clOutS                val clppStrm = PrintAsCL.new clOutS
691                val cppStrm = PrintAsC.new cOutS                val cppStrm = PrintAsC.new cOutS
692                  val progName = name
693                fun cppDecl dcl = PrintAsC.output(cppStrm, dcl)                fun cppDecl dcl = PrintAsC.output(cppStrm, dcl)
694                fun clppDecl dcl = PrintAsC.output(clppStrm, dcl)                fun clppDecl dcl = PrintAsCL.output(clppStrm, dcl)
695                val strands = AtomTable.listItems strands                val strands = AtomTable.listItems strands
696                val single_strand as Strand{name, tyName, code, ...}= hd(strands)                val [strand as Strand{name, tyName, code, init_code, ...}] = strands
697                in                in
   
698              (* Generate the OpenCl file *)              (* Generate the OpenCl file *)
699              List.app clppDecl (List.rev (!globals));                  clppDecl (CL.D_Verbatim([
700              clppDecl (genGlobalStruct (!globals));                      if double
701              clppDecl (genStrandTyDef single_strand);                        then "#define DIDEROT_DOUBLE_PRECISION"
702                          else "#define DIDEROT_SINGLE_PRECISION",
703                        "#define DIDEROT_TARGET_CL",
704                        "#include \"Diderot/cl-diderot.h\""
705                      ]));
706                    clppDecl (genGlobalStruct (#gpuTy, !globals, RN.globalsTy));
707                    clppDecl (genImageDataStruct(!imgGlobals,RN.imageDataType));
708                    clppDecl (genStrandTyDef(#gpuTy, strand));
709                    clppDecl  (!init_code);
710              List.app clppDecl (!code);              List.app clppDecl (!code);
711              clppDecl (genKernelFun (single_strand,!numDims,globals,imgGlobals));                  clppDecl (genKernelFun (strand, !numDims, globals, imgGlobals));
712                  (* Generate the Host C file *)
   
             (* Generate the Host file .c *)  
713              cppDecl (CL.D_Verbatim([              cppDecl (CL.D_Verbatim([
714                          if double                          if double
715                            then "#define DIDEROT_DOUBLE_PRECISION"                            then "#define DIDEROT_DOUBLE_PRECISION"
716                            else "#define DIDEROT_SINGLE_PRECISION",                            else "#define DIDEROT_SINGLE_PRECISION",
717                           "#include \"Diderot/diderot.h\"",                      "#define DIDEROT_TARGET_CL",
718                           "#include <OpenCL/OpenCL.h>",                      "#include \"Diderot/diderot.h\""
                          "#include <sys/sysctl.h>",  
                          "#include <sys/stat.h>",  
                          "#include <assert.h>"  
719                        ]));                        ]));
720                    cppDecl (CL.D_Var(["static"], CL.charPtr, "ProgramName",
721              (* cppDecl (CL.D_Verbatim([ "#include <OpenCL/OpenCL.h>",                    SOME(CL.I_Exp(CL.mkStr progName))));
722                                                                   "#include Diderot/diderot.h"])); *)                  cppDecl (genGlobalStruct (#hostTy, !globals, RN.globalsTy));
723                  List.app cppDecl (List.rev (!globals));                  cppDecl (genGlobalStruct (#shadowTy, !globals, RN.shadowGlobalsTy));
724              cppDecl (genGlobalStruct (!globals));  (* FIXME: does this really need to be a global? *)
725              cppDecl (genStrandTyDef single_strand);                  cppDecl (CL.D_Var(["static"], globPtrTy, RN.globalsVarName, NONE));
726                  cppDecl (genStrandPrint(single_strand,!numDims));                  cppDecl (genStrandTyDef (#hostTy, strand));
727              cppDecl (genKernelLoader());                  cppDecl (genStrandPrint strand);
728              List.app cppDecl (List.rev (!topDecls));              List.app cppDecl (List.rev (!topDecls));
729              cppDecl (genHostSetupFunc (single_strand,clFileName,!numDims,initially,imgGlobals));                  cppDecl (genGlobalBuffersArgs (!globals,imgGlobals));
730              cppDecl (genHostMain());                  List.app (fn strand => cppDecl (genStrandDesc strand)) strands;
731                    genStrandTable (cppDecl, strands);
732                    cppDecl (!initially);
   
                 (*List.app (fn strand => List.app ppDecl (genStrand strand)) strands;  
                  genStrandTable (ppStrm, strands);  
                 ppDecl (!initially);*)  
   
733                  PrintAsC.close cppStrm;                  PrintAsC.close cppStrm;
734                  PrintAsC.close clppStrm;                  PrintAsCL.close clppStrm;
735                  TextIO.closeOut cOutS;                  TextIO.closeOut cOutS;
736                  TextIO.closeOut clOutS                  TextIO.closeOut clOutS
737                end                end
738    
739        (* 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.  
        *)  
740          fun generate (basename, prog as Prog{double, parallel, debug, ...}) = let          fun generate (basename, prog as Prog{double, parallel, debug, ...}) = let
741                fun condCons (true, x, xs) = x::xs                fun condCons (true, x, xs) = x::xs
742                  | condCons (false, _, xs) = xs                  | condCons (false, _, xs) = xs
# Line 903  Line 750 
750              (* generate the loader flags *)              (* generate the loader flags *)
751                val extraLibs = condCons (parallel, #pthread Paths.extraLibs, [])                val extraLibs = condCons (parallel, #pthread Paths.extraLibs, [])
752                val extraLibs = Paths.teemLinkFlags @  #base Paths.extraLibs :: extraLibs                val extraLibs = Paths.teemLinkFlags @  #base Paths.extraLibs :: extraLibs
753                       val extraLibs =  #cl Paths.extraLibs :: extraLibs
754                val rtLib = TargetUtil.runtimeName {                val rtLib = TargetUtil.runtimeName {
755                        target = TargetUtil.TARGET_CL,                        target = TargetUtil.TARGET_CL,
756                        parallel = parallel, double = double, debug = debug                        parallel = parallel, double = double, debug = debug
757                      }                      }
758                val ldOpts = rtLib :: extraLibs                val ldOpts = rtLib :: extraLibs
759                in                in
760                  genSrc (basename, prog)                  genSrc (basename, prog);
761                    RunCC.compile (basename, cflags);
762                    RunCC.link (basename, ldOpts)
763                  end                  end
764    
                 (*RunCC.compile (basename, cflags);  
                 RunCC.link (basename, ldOpts)*)  
   
   
765        end        end
766    
767    (* strands *)    (* strands *)
768      structure Strand =      structure Strand =
769        struct        struct
# Line 927  Line 774 
774                        tyName = RN.strandTy name,                        tyName = RN.strandTy name,
775                        state = ref [],                        state = ref [],
776                        output = ref NONE,                        output = ref NONE,
777                        code = ref []                        code = ref [],
778                          init_code = ref (CL.D_Comment(["no init code"]))
779                      }                      }
780                in                in
781                  AtomTable.insert strands (strandId, strand);                  AtomTable.insert strands (strandId, strand);
# Line 940  Line 788 
788        (* register the strand-state initialization code.  The variables are the strand        (* register the strand-state initialization code.  The variables are the strand
789         * parameters.         * parameters.
790         *)         *)
791          fun init (Strand{name, tyName, code, ...}, params, init) = let          fun init (Strand{name, tyName, code, init_code, ...}, params, init) = let
792                val fName = RN.strandInit name                val fName = RN.strandInit name
793                val params =                val params =
794                      CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfOut") ::                      globalParam (globPtrTy, RN.globalsVarName) ::
795                        List.map (fn (ToC.V(ty, x)) => CL.PARAM([], ty, x)) params                      globalParam (CL.T_Ptr(CL.T_Named tyName), "selfOut") ::
796                          List.map (fn (ToCL.V(ty, x)) => CL.PARAM([], ty, x)) params
797                val initFn = CL.D_Func([], CL.voidTy, fName, params, init)                val initFn = CL.D_Func([], CL.voidTy, fName, params, init)
798                in                in
799                  code := initFn :: !code                  init_code := initFn
800                end                end
801    
802        (* register a strand method *)        (* register a strand method *)
803          fun method (Strand{name, tyName, code, ...}, methName, body) = let          fun method (Strand{name, tyName, code, ...}, methName, body) = let
804                val fName = concat[name, "_", methName]                val fName = concat[name, "_", methName]
805                val params = [                val params = [
806                        CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfIn"),                        globalParam (CL.T_Ptr(CL.T_Named tyName), "selfIn"),
807                        CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfOut")                        globalParam (CL.T_Ptr(CL.T_Named tyName), "selfOut"),
808                          globalParam (CL.T_Ptr(CL.T_Named (RN.globalsTy)), RN.globalsVarName),
809                          CL.PARAM([],CL.T_Named(RN.imageDataType),RN.globalImageDataName)
810                      ]                      ]
811                val methFn = CL.D_Func([], CL.int32, fName, params, body)                val methFn = CL.D_Func([], CL.int32, fName, params, body)
812                in                in
813                  code := methFn :: !code                  code := methFn :: !code
814                end                end
815    
816          fun output (Strand{output, ...}, ty, ToC.V(_, x)) = output := SOME(ty, x)          fun output (Strand{output, ...}, ty, ToCL.V(_, x)) = output := SOME(ty, x)
817    
818        end        end
819    

Legend:
Removed from v.1264  
changed lines
  Added in v.1425

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