Home My Page Projects Code Snippets Project Openings diderot
Summary Activity Tracker Tasks SCM

SCM Repository

[diderot] Diff of /branches/pure-cfg/src/compiler/cl-target/cl-target.sml
ViewVC logotype

Diff of /branches/pure-cfg/src/compiler/cl-target/cl-target.sml

Parent Directory Parent Directory | Revision Log Revision Log | View Patch Patch

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

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

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