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 1256, Tue May 24 17:58:28 2011 UTC revision 1286, Tue Jun 7 10:54:18 2011 UTC
# Line 12  Line 12 
12      structure Ty = IL.Ty      structure Ty = IL.Ty
13      structure CL = CLang      structure CL = CLang
14      structure RN = RuntimeNames      structure RN = RuntimeNames
15      structure ToC = TreeToCL      structure ToCL = TreeToCL
16    
17      type var = ToC.var    (* C variable translation *)
18        structure TrCVar =
19          struct
20            type env = CL.typed_var TreeIL.Var.Map.map
21            fun lookup (env, x) = (case V.Map.find (env, x)
22                   of SOME(CL.V(_, x')) => x'
23                    | NONE => raise Fail(concat["TrCVar.lookup(_, ", V.name x, ")"])
24                  (* end case *))
25          (* translate a variable that occurs in an l-value context (i.e., as the target of an assignment) *)
26            fun lvalueVar (env, x) = (case V.kind x
27                   of IL.VK_Global => CL.mkIndirect(CL.mkVar RN.globalsVarName, lookup(env, x))
28                    | IL.VK_State strand => raise Fail "unexpected strand context"
29                    | IL.VK_Local => CL.mkVar(lookup(env, x))
30                  (* end case *))
31          (* translate a variable that occurs in an r-value context *)
32            val rvalueVar = lvalueVar
33          end
34    
35        structure ToC = TreeToCFn (TrCVar)
36    
37        type var = CL.typed_var
38      type exp = CL.exp      type exp = CL.exp
39      type stm = CL.stm      type stm = CL.stm
40    
41      (* OpenCL specific types *)
42        val clProgramTy = CL.T_Named "cl_program"
43        val clKernelTy  = CL.T_Named "cl_kernel"
44        val clCmdQueueTy = CL.T_Named "cl_command_queue"
45        val clContextTy = CL.T_Named "cl_context"
46        val clDeviceIdTy = CL.T_Named "cl_device_id"
47        val clPlatformIdTy = CL.T_Named "cl_platform_id"
48        val clMemoryTy = CL.T_Named "cl_mem"
49    
50      datatype strand = Strand of {      datatype strand = Strand of {
51          name : string,          name : string,
52          tyName : string,          tyName : string,
53          state : var list ref,          state : var list ref,
54          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) *)
55          code : CL.decl list ref          code : CL.decl list ref,
56            init_code: CL.decl ref
57        }        }
58    
59      datatype program = Prog of {      datatype program = Prog of {
60            name : string,                  (* stem of source file *)
61          double : bool,                  (* true for double-precision support *)          double : bool,                  (* true for double-precision support *)
62          parallel : bool,                (* true for multithreaded (or multi-GPU) target *)          parallel : bool,                (* true for multithreaded (or multi-GPU) target *)
63          debug : bool,                   (* true for debug support in executable *)          debug : bool,                   (* true for debug support in executable *)
# Line 36  Line 67 
67          initially : CL.stm list ref,          initially : CL.stm list ref,
68          numDims: int ref,          numDims: int ref,
69          imgGlobals: (string * int) list ref,          imgGlobals: (string * int) list ref,
70          oneDim: CL.exp ref,          prFn: CL.decl ref
         twoDim: CL.exp ref,  
         thirdDim: CL.exp ref  
71      }      }
72    
73      datatype env = ENV of {      datatype env = ENV of {
# Line 58  Line 87 
87        | StrandScope of TreeIL.var list  (* strand initialization *)        | StrandScope of TreeIL.var list  (* strand initialization *)
88        | MethodScope of TreeIL.var list  (* method body; vars are state variables *)        | MethodScope of TreeIL.var list  (* method body; vars are state variables *)
89    
90    (* the supprted widths of vectors of reals on the target.  For the GNU vector extensions,    (* the supprted widths of vectors of reals on the target. *)
91     * the supported sizes are powers of two, but float2 is broken.  (* FIXME: for OpenCL 1.1, 3 is also valid *)
92     * NOTE: we should also consider the AVX vector hardware, which has 256-bit registers.      fun vectorWidths () = [2, 4, 8, 16]
    *)  
     fun vectorWidths () = if !RuntimeNames.doublePrecision  
           then [2, 4, 8]  
           else [4, 8]  
93    
94    (* tests for whether various expression forms can appear inline *)    (* tests for whether various expression forms can appear inline *)
95      fun inlineCons n = (n < 2)          (* vectors are inline, but not matrices *)      fun inlineCons n = (n < 2)          (* vectors are inline, but not matrices *)
# Line 73  Line 98 
98    (* TreeIL to target translations *)    (* TreeIL to target translations *)
99      structure Tr =      structure Tr =
100        struct        struct
101          (* this function is used for the initially clause, so it generates OpenCL *)
102          fun fragment (ENV{info, vMap, scope}, blk) = let          fun fragment (ENV{info, vMap, scope}, blk) = let
103                val (vMap, stms) = ToC.trFragment (vMap, blk)                val (vMap, stms) = ToCL.trFragment (vMap, blk)
104                in                in
105                  (ENV{info=info, vMap=vMap, scope=scope}, stms)                  (ENV{info=info, vMap=vMap, scope=scope}, stms)
106                end                end
107          fun saveState cxt stateVars (env, args, stm) = (          fun saveState cxt stateVars (env, args, stm) = (
108                ListPair.foldrEq                ListPair.foldrEq
109                  (fn (x, e, stms) => ToC.trAssign(env, x, e)@stms)                  (fn (x, e, stms) => ToCL.trAssign(env, x, e)@stms)
110                    [stm]                    [stm]
111                      (stateVars, args)                      (stateVars, args)
112                ) handle ListPair.UnequalLengths => (                ) handle ListPair.UnequalLengths => (
113                  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"]);
114                  raise Fail(concat["saveState ", cxt, ": length mismatch"]))                  raise Fail(concat["saveState ", cxt, ": length mismatch"]))
115          fun block (ENV{vMap, scope, ...}, blk) = (case scope          fun block (ENV{vMap, scope, ...}, blk) = (case scope
116                 of StrandScope stateVars => ToC.trBlock (vMap, saveState "StrandScope" stateVars, blk)                 of StrandScope stateVars => ToCL.trBlock (vMap, saveState "StrandScope" stateVars, blk)
117                  | MethodScope stateVars => ToC.trBlock (vMap, saveState "MethodScope" stateVars, blk)                  | MethodScope stateVars => ToCL.trBlock (vMap, saveState "MethodScope" stateVars, blk)
118                    | InitiallyScope => ToCL.trBlock (vMap, fn (_, _, stm) => [stm], blk)
119                  | _ => ToC.trBlock (vMap, fn (_, _, stm) => [stm], blk)                  | _ => ToC.trBlock (vMap, fn (_, _, stm) => [stm], blk)
120                (* end case *))                (* end case *))
121          fun exp (ENV{vMap, ...}, e) = ToC.trExp(vMap, e)          fun exp (ENV{vMap, ...}, e) = ToCL.trExp(vMap, e)
122        end        end
123    
124    (* variables *)    (* variables *)
125      structure Var =      structure Var =
126        struct        struct
127          fun name (ToC.V(_, name)) = name          fun name (ToCL.V(_, name)) = name
128           fun global (Prog{globals,imgGlobals, ...}, name, ty) = let           fun global (Prog{globals,imgGlobals, ...}, name, ty) = let
129                val ty' = ToC.trType ty                val ty' = ToCL.trType ty
130                fun isImgGlobal (imgGlobals, Ty.ImageTy(ImageInfo.ImgInfo{dim, ...}), name) =  imgGlobals  := (name,dim):: !imgGlobals                fun isImgGlobal (imgGlobals, Ty.ImageTy(ImageInfo.ImgInfo{dim, ...}), name) =  imgGlobals  := (name,dim):: !imgGlobals
131                  | isImgGlobal (imgGlobals, _, _) =  ()                  | isImgGlobal (imgGlobals, _, _) =  ()
132                in                in
133                  globals := CL.D_Var([], ty', name, NONE) :: !globals;                  globals := CL.D_Var([], ty', name, NONE) :: !globals;
134                  isImgGlobal(imgGlobals,ty,name);                  isImgGlobal(imgGlobals,ty,name);
135               ToC.V(ty', name)                  ToCL.V(ty', name)
136                end                end
137          fun param x = ToC.V(ToC.trType(V.ty x), V.name x)          fun param x = ToCL.V(ToCL.trType(V.ty x), V.name x)
138          fun state (Strand{state, ...}, x) = let          fun state (Strand{state, ...}, x) = let
139                val ty' = ToC.trType(V.ty x)                val ty' = ToCL.trType(V.ty x)
140                val x' = ToC.V(ty', V.name x)                val x' = ToCL.V(ty', V.name x)
141                in                in
142                  state := x' :: !state;                  state := x' :: !state;
143                  x'                  x'
# Line 143  Line 170 
170    (* programs *)    (* programs *)
171      structure Program =      structure Program =
172        struct        struct
173          fun new {double, parallel, debug} = (          fun new {name, double, parallel, debug} = (
174                RN.initTargetSpec double;                RN.initTargetSpec double;
175                  CNames.initTargetSpec double;
176                Prog{                Prog{
177                      name = name,
178                    double = double, parallel = parallel, debug = debug,                    double = double, parallel = parallel, debug = debug,
179                    globals = ref [                    globals = ref [],
                     CL.D_Verbatim[  
                         if double  
                           then "#define DIDEROT_DOUBLE_PRECISION"  
                           else "#define DIDEROT_SINGLE_PRECISION",  
                         "#include \"Diderot/opencl_types.h\""  
                       ]],  
180                    topDecls = ref [],                    topDecls = ref [],
181                    strands = AtomTable.mkTable (16, Fail "strand table"),                    strands = AtomTable.mkTable (16, Fail "strand table"),
182                    initially = ref([CL.S_Comment["missing initially"]]),                    initially = ref([CL.S_Comment["missing initially"]]),
183                    numDims = ref(0),                    numDims = ref(0),
184                    imgGlobals = ref[],                    imgGlobals = ref[],
185                    oneDim = ref(CL.E_Str "did not initalize dim"),                                    prFn = ref(CL.D_Comment(["No Print Function"]))
                   twoDim = ref(CL.E_Str "did not initalize dim"),  
                   thirdDim = ref(CL.E_Str "did not initalize dim")  
186                  })                  })
187        (* register the global initialization part of a program *)        (* register the global initialization part of a program *)
188          fun globalIndirects (globals,stms) = let          fun globalIndirects (globals,stms) = let
189                   fun getGlobals(CL.D_Var(_,_,globalVar,_)::rest) = CL.mkAssign(CL.mkIndirect(CL.E_Var RN.globalsVarName,globalVar),CL.E_Var globalVar)::getGlobals(rest)                  fun getGlobals (CL.D_Var(_,_,globalVar,_)::rest) =
190                     | getGlobals([]) = []                        CL.mkAssign(CL.mkIndirect(CL.mkVar RN.globalsVarName,globalVar),CL.mkVar globalVar)
191                     | getGlobals(_::rest) = getGlobals(rest)                          ::getGlobals rest
192                      | getGlobals [] = []
193                      | getGlobals (_::rest) = getGlobals rest
194                    in
195                      stms @ getGlobals globals
196                    end
197    
198          (* register the code that is used to register command-line options for input variables *)
199            fun inputs (Prog{topDecls, ...}, stm) = let
200                  val inputsFn = CL.D_Func(
201                        [], CL.voidTy, RN.registerOpts,
202                        [CL.PARAM([], CL.T_Ptr(CL.T_Named RN.optionsTy), "opts")],
203                        stm)
204                  in                  in
205                          stms @ getGlobals(globals)                  topDecls := inputsFn :: !topDecls
206                  end                  end
207    
208          fun init (Prog{globals,topDecls,...}, CL.S_Block(init)) = let        (* register the global initialization part of a program *)
209                val params = [          fun init (Prog{topDecls, ...}, init) = let
210                            CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName)                val globPtrTy = CL.T_Ptr(CL.T_Named RN.globalsTy)
211                          ]                val initFn = CL.D_Func(
212                     val body = CL.S_Block(globalIndirects(!globals,init))                      [], CL.voidTy, RN.initGlobals, [CL.PARAM([], globPtrTy, RN.globalsVarName)],
213                     val initFn = CL.D_Func([], CL.voidTy, RN.initGlobals, params, body)                      init)
214                  val shutdownFn = CL.D_Func(
215                in                      [], CL.voidTy, RN.shutdown,
216                          topDecls := initFn :: !topDecls                      [CL.PARAM([], CL.T_Ptr(CL.T_Named RN.worldTy), "wrld")],
217                end                      CL.S_Block[])
   
           | init (Prog{globals,topDecls,...}, init) = let  
               val params = [  
                           CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName)  
                         ]  
                    val initFn = CL.D_Func([], CL.voidTy, RN.initGlobals, params, init)  
   
218                in                in
219                  topDecls := initFn :: !topDecls                  topDecls := shutdownFn :: initFn :: !topDecls
220                end                end
221    
222        (* create and register the initially function for a program *)        (* create and register the initially function for a program *)
223          fun initially {          fun initially {
224                prog = Prog{strands, initially,numDims,oneDim,twoDim,thirdDim,...},                prog = Prog{strands, initially, numDims,...},
225                isArray : bool,                isArray : bool,
226                iterPrefix : stm list,                iterPrefix : stm list,
227                iters : (var * exp * exp) list,                iters : (var * exp * exp) list,
# Line 213  Line 239 
239                      end                      end
240                val baseInit = mapi (fn (i, (_, e, _)) => (i, CL.I_Exp e)) iters                val baseInit = mapi (fn (i, (_, e, _)) => (i, CL.I_Exp e)) iters
241                val sizeInit = mapi                val sizeInit = mapi
242                      (fn (i, (ToC.V(ty, _), lo, hi)) =>                      (fn (i, (ToCL.V(ty, _), lo, hi)) =>
243                          (i, CL.I_Exp(CL.mkBinOp(CL.mkBinOp(hi, CL.#-, lo), CL.#+, CL.E_Int(1, ty))))                          (i, CL.I_Exp(CL.mkBinOp(CL.mkBinOp(hi, CL.#-, lo), CL.#+, CL.mkInt(1, ty))))
244                      ) iters                      ) iters
245                    val numStrandsVar = "numStrandsVar"                    val numStrandsVar = "numStrandsVar"
246                val allocCode = iterPrefix @ [                val allocCode = iterPrefix @ [
247                        CL.mkComment["allocate initial block of strands"],                        CL.mkComment["allocate initial block of strands"],
248                        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)),
249                        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)),
250                        CL.mkDecl(CL.int32,"numDims",SOME(CL.I_Exp(CL.E_Int(IntInf.fromInt nDims, CL.int32))))                        CL.mkDecl(CL.int32,"numDims",SOME(CL.I_Exp(CL.mkInt(IntInf.fromInt nDims, CL.int32))))
251                            ]                            ]
252                  val numStrandsLoopBody =
253              fun mkLoopNest ([],_,_,_,_) = ()                      CL.mkExpStm(CL.mkAssignOp(CL.mkVar numStrandsVar, CL.*=,CL.mkSubscript(CL.mkVar "size",CL.mkVar "i")))
254            | mkLoopNest ((ToC.V(ty, param), lo, hi)::iters, oneDim,twoDim,thirdDim, 3) =                val numStrandsLoop =  CL.mkFor([(CL.intTy, "i", CL.mkInt(0,CL.intTy))],
255                                  (oneDim := hi; mkLoopNest (iters,oneDim,twoDim,thirdDim, 2))                      CL.mkBinOp(CL.mkVar "i", CL.#<, CL.mkVar "numDims"),
256            | mkLoopNest ((ToC.V(ty, param), lo, hi)::iters, oneDim,twoDim,thirdDim, 2) =                      [CL.mkPostOp(CL.mkVar "i", CL.^++)], numStrandsLoopBody)
                                 (twoDim := hi; mkLoopNest (iters,oneDim,twoDim,thirdDim, 1))  
           | mkLoopNest ((ToC.V(ty, param), lo, hi)::iters, oneDim,twoDim,thirdDim, 1) =  
                                  (thirdDim := hi; mkLoopNest (iters,oneDim,twoDim,thirdDim, 0))  
           | mkLoopNest ((ToC.V(ty, param), lo, hi)::iters,_,_,_,_) = ()  
   
   
   
                   val numStrandsLoopBody = CL.mkExpStm(CL.mkAssignOp(CL.E_Var numStrandsVar, CL.*=,CL.mkSubscript(CL.E_Var "size",CL.E_Var "i")))  
   
   
                   val numStrandsLoop =  CL.mkFor([(CL.intTy, "i", CL.E_Int(0,CL.intTy))],  
                                                                                    CL.mkBinOp(CL.E_Var "i", CL.#<, CL.E_Var "numDims"),  
                                                                                    [CL.mkPostOp(CL.E_Var "i", CL.^++)], numStrandsLoopBody)  
257                in                in
258                    numDims := nDims;                    numDims := nDims;
259                    initially := allocCode @ [numStrandsLoop];                  initially := allocCode @ [numStrandsLoop]
                   mkLoopNest (iters,oneDim, twoDim, thirdDim, nDims)  
   
260                end                end
261    
262    
263        (***** OUTPUT *****)        (***** OUTPUT *****)
264      fun genStrand (Strand{name, tyName, state, output, code}) = let          fun genStrandInit (Strand{name,tyName,state,output,code,...}, nDims) = let
265                  val params = [
266                          CL.PARAM([], CL.T_Ptr(CL.uint32), "sizes" ),
267                          CL.PARAM([], CL.intTy, "width"),
268                          CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "strands")
269                        ]
270                  val body = let
271                        fun loopParams 3 = ["x", "y", "k"]
272                          | loopParams 2 = ["x", "y"]
273                          | loopParams 1 = ["x"]
274                          | loopParams _ = raise Fail "genStrandInit: missing size dim"
275                        fun mkLoopNest ([], _, nDims) = if nDims = 1
276                              then CL.mkBlock [
277                                  CL.mkCall(RN.strandInit name, [
278                                    CL.mkUnOp(CL.%&,CL.mkSubscript(CL.mkVar "strands",CL.mkStr "x")),
279                                                    CL.mkVar "x"])
280                                ]
281                              else let
282                                val index = CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")
283                                in
284                                  CL.mkBlock([CL.mkCall(RN.strandInit name, [CL.mkUnOp(CL.%&,CL.mkSubscript(CL.mkVar "strands",index)),
285                                  CL.mkVar "x", CL.mkVar"y"])])
286                                end
287                          | mkLoopNest (param::rest,count,nDims) = let
288                              val body = mkLoopNest (rest, count + 1,nDims)
289                              in
290                                CL.mkFor(
291                                    [(CL.intTy, param, CL.mkInt(0,CL.intTy))],
292                                    CL.mkBinOp(CL.mkVar param, CL.#<, CL.mkSubscript(CL.mkVar "sizes",CL.mkInt(count,CL.intTy))),
293                                    [CL.mkPostOp(CL.mkVar param, CL.^++)],
294                                    body)
295                              end
296                        in
297                          [mkLoopNest ((loopParams nDims),0,nDims)]
298                        end
299                    in
300                      CL.D_Func(["static"], CL.voidTy, RN.strandInitSetup, params,CL.mkBlock(body))
301                    end
302    
303            fun genStrandPrint (Strand{name, tyName, state, output, code,...},nDims) = let
304              (* the print function *)              (* the print function *)
305                val prFnName = concat[name, "_print"]                val prFnName = concat[name, "_print"]
306                val prFn = let                val prFn = let
307                      val params = [                      val params = [
308                            CL.PARAM([], CL.T_Ptr(CL.T_Named "FILE"), "outS"),                            CL.PARAM([], CL.T_Ptr(CL.T_Named "FILE"), "outS"),
309                              CL.PARAM([], CL.T_Ptr(CL.uint32), "sizes" ),
310                              CL.PARAM([], CL.intTy, "width"),
311                            CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "self")                            CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "self")
312                          ]                          ]
313    
314                      val SOME(ty, x) = !output                      val SOME(ty, x) = !output
315                      val outState = CL.mkIndirect(CL.mkVar "self", x)                     val outState = if nDims = 1 then
316                              CL.mkSelect(CL.mkSubscript(CL.mkVar "self",CL.mkVar "x"), x)
317                            else if nDims = 2 then
318                                    CL.mkSelect(CL.mkSubscript(CL.mkVar "self",
319                                       CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")), x)
320    
321                            else CL.mkSelect(CL.mkVar "self",x)
322    
323                      val prArgs = (case ty                      val prArgs = (case ty
324                             of Ty.IVecTy 1 => [CL.E_Str(!RN.gIntFormat ^ "\n"), outState]                             of Ty.IVecTy 1 => [CL.mkStr(!RN.gIntFormat ^ "\n"), outState]
325                              | Ty.IVecTy d => let                              | Ty.IVecTy d => let
326                                  val fmt = CL.E_Str(                                  val fmt = CL.mkStr(
327                                        String.concatWith " " (List.tabulate(d, fn _ => !RN.gIntFormat))                                        String.concatWith " " (List.tabulate(d, fn _ => !RN.gIntFormat))
328                                        ^ "\n")                                        ^ "\n")
329                                  val args = List.tabulate (d, fn i => ToC.ivecIndex(outState, d, i))                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))
330                                  in                                  in
331                                    fmt :: args                                    fmt :: args
332                                  end                                  end
333                              | Ty.TensorTy[] => [CL.E_Str "%f\n", outState]                              | Ty.TensorTy[] => [CL.mkStr "%f\n", outState]
334                              | Ty.TensorTy[d] => let                              | Ty.TensorTy[d] => let
335                                  val fmt = CL.E_Str(                                  val fmt = CL.mkStr(
336                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))
337                                        ^ "\n")                                        ^ "\n")
338                                  val args = List.tabulate (d, fn i => ToC.vecIndex(outState, d, i))                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))
339                                  in                                  in
340                                    fmt :: args                                    fmt :: args
341                                  end                                  end
342                              | _ => raise Fail("genStrand: unsupported output type " ^ Ty.toString ty)                              | _ => raise Fail("genStrand: unsupported output type " ^ Ty.toString ty)
343                            (* end case *))                            (* end case *))
344    
345                              val body = let
346    
347                                fun loopParams (3) =
348                                     "x"::"y"::"k"::[]
349                                  | loopParams (2) =
350                                     "x"::"y"::[]
351                                  | loopParams (1) =
352                                     "x"::[]
353                                  | loopParams (_) =
354                                    raise Fail("genStrandPrint: unsupported output type " ^ Ty.toString ty)
355    
356                               fun mkLoopNest ([],_) =
357                                                    CL.mkCall("fprintf", CL.mkVar "outS" :: prArgs)
358                                    | mkLoopNest (param::rest,count) = let
359                                            val body = mkLoopNest (rest, count + 1)
360                                       in
361                                                    CL.mkFor(
362                                                            [(CL.intTy, param, CL.mkInt(0,CL.intTy))],
363                                                    CL.mkBinOp(CL.mkVar param, CL.#<, CL.mkSubscript(CL.mkVar "sizes",CL.mkInt(count,CL.intTy))),
364                                                    [CL.mkPostOp(CL.mkVar param, CL.^++)],
365                                                    body)
366                                       end
367                            in
368                                    [mkLoopNest ((loopParams nDims),0)]
369                            end
370    
371                      in                      in
372                        CL.D_Func(["static"], CL.voidTy, prFnName, params,                        CL.D_Func(["static"], CL.voidTy, prFnName, params,CL.mkBlock(body))
                         CL.mkCall("fprintf", CL.mkVar "outS" :: prArgs))  
373                      end                      end
374                in                in
375                                   List.rev (prFn :: !code)                                   prFn
376                end                end
377          fun genStrandTyDef (Strand{tyName, state,...}) =          fun genStrandTyDef (Strand{tyName, state,...}) =
378              (* the type declaration for the strand's state struct *)              (* the type declaration for the strand's state struct *)
379                CL.D_StructDef(                CL.D_StructDef(
380                        List.rev (List.map (fn ToC.V(ty, x) => (ty, x)) (!state)),                        List.rev (List.map (fn ToCL.V(ty, x) => (ty, x)) (!state)),
381                        tyName)                        tyName)
382    
383    
384          (* generates the load kernel function *)          (* generates the load kernel function *)
385    (* FIXME: this code might be part of the runtime system *)
386          fun genKernelLoader() =          fun genKernelLoader() =
387                  CL.D_Verbatim ( ["/* Loads the Kernel from a file */",                  CL.D_Verbatim ( ["/* Loads the Kernel from a file */",
388                                                  "char * loadKernel (const char * filename) {",                                                  "char * loadKernel (const char * filename) {",
# Line 312  Line 400 
400                                                  "}"])                                                  "}"])
401          (* generates the opencl buffers for the image data *)          (* generates the opencl buffers for the image data *)
402          fun getGlobalDataBuffers(globals,count,contextVar,errVar) = let          fun getGlobalDataBuffers(globals,count,contextVar,errVar) = let
403          val globalBufferDecl =  CL.mkDecl(CL.clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)                  val globalBufferDecl =  CL.mkDecl(clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)
404          val globalBuffer = CL.mkAssign(CL.E_Var(concat[RN.globalsVarName,"_cl"]), CL.mkApply("clCreateBuffer",                  val globalBuffer = CL.mkAssign(CL.mkVar(concat[RN.globalsVarName,"_cl"]), CL.mkApply("clCreateBuffer",
405                                                                  [CL.E_Var contextVar,                                                                  [CL.mkVar contextVar,
406                                                                  CL.E_Var "CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR | CL_MEM_COPY_HOST_PTR",                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",
407                                                                  CL.mkApply("sizeof",[CL.E_Var RN.globalsTy]),                                                                  CL.mkApply("sizeof",[CL.mkVar RN.globalsTy]),
408                                                                  CL.E_Var RN.globalsVarName,                                                                  CL.mkVar RN.globalsVarName,
409                                                                  CL.E_UnOp(CL.%&,CL.E_Var errVar)]))                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))
410    
411          fun genDataBuffers([],_,_,_) = []          fun genDataBuffers([],_,_,_) = []
412            | genDataBuffers((var,nDims)::globals,count,contextVar,errVar) = let            | genDataBuffers((var,nDims)::globals,count,contextVar,errVar) = let
413            (* FIXME: use CL constructors to  build expressions (not strings) *)
414                     val size = if nDims = 1 then                     val size = if nDims = 1 then
415                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.E_Var "float"]), CL.#*,                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,
416                                           CL.mkIndirect(CL.E_Var var, "size[0]"))                                           CL.mkIndirect(CL.mkVar var, "size[0]"))
417                                          else if nDims = 2 then                                          else if nDims = 2 then
418                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.E_Var "float"]), CL.#*,                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,
419                                            CL.mkIndirect(CL.E_Var var, concat["size[0]", " * ", var, "->size[1]"]))                                            CL.mkIndirect(CL.mkVar var, concat["size[0]", " * ", var, "->size[1]"]))
420                                          else                                          else
421                                           CL.mkBinOp(CL.mkApply("sizeof",[CL.E_Var "float"]), CL.#*,                                           CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,
422                                            CL.mkIndirect(CL.E_Var var,concat["size[0]", " * ", var, "->size[1] * ", var, "->size[2]"]))                                            CL.mkIndirect(CL.mkVar var,concat["size[0]", " * ", var, "->size[1] * ", var, "->size[2]"]))
423    
424                   in                   in
425                     CL.mkDecl(CL.clMemoryTy,RN.addBufferSuffix var ,NONE)::                     CL.mkDecl(clMemoryTy, RN.addBufferSuffix var ,NONE)::
426                     CL.mkDecl(CL.clMemoryTy,RN.addBufferSuffixData var ,NONE)::                     CL.mkDecl(clMemoryTy, RN.addBufferSuffixData var ,NONE)::
427                     CL.mkAssign(CL.E_Var(RN.addBufferSuffix var), CL.mkApply("clCreateBuffer",                     CL.mkAssign(CL.mkVar(RN.addBufferSuffix var), CL.mkApply("clCreateBuffer",
428                                                                  [CL.E_Var contextVar,                                                                  [CL.mkVar contextVar,
429                                                                  CL.E_Var "CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR | CL_MEM_COPY_HOST_PTR",                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",
430                                                                  CL.mkApply("sizeof",[CL.E_Var (RN.imageTy nDims)]),                                                                  CL.mkApply("sizeof",[CL.mkVar (RN.imageTy nDims)]),
431                                                                  CL.E_Var var,                                                                  CL.mkVar var,
432                                                                  CL.E_UnOp(CL.%&,CL.E_Var errVar)])) ::                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)])) ::
433                          CL.mkAssign(CL.E_Var(RN.addBufferSuffixData var), CL.mkApply("clCreateBuffer",                          CL.mkAssign(CL.mkVar(RN.addBufferSuffixData var), CL.mkApply("clCreateBuffer",
434                                                                  [CL.E_Var contextVar,                                                                  [CL.mkVar contextVar,
435                                                                  CL.E_Var "CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR | CL_MEM_COPY_HOST_PTR",                                                                   CL.mkVar "CL_MEM_COPY_HOST_PTR",
436                                                                  size,                                                                  size,
437                                                                  CL.mkIndirect(CL.E_Var var,"data"),                                                                  CL.mkIndirect(CL.mkVar var,"data"),
438                                                                  CL.E_UnOp(CL.%&,CL.E_Var errVar)])):: genDataBuffers(globals,count + 2,contextVar,errVar)                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)])):: genDataBuffers(globals,count + 2,contextVar,errVar)
439                  end                  end
440          in          in
441                  [globalBufferDecl] @ [globalBuffer] @ genDataBuffers(globals,count + 2,contextVar,errVar)                  [globalBufferDecl] @ [globalBuffer] @ genDataBuffers(globals,count + 2,contextVar,errVar)
   
442          end          end
443    
444    
445          (* generates the kernel arguments for the image data *)          (* generates the kernel arguments for the image data *)
446          fun genGlobalArguments(globals,count,kernelVar,errVar) = let          fun genGlobalArguments(globals,count,kernelVar,errVar) = let
447          val globalArgument = CL.mkAssign(CL.E_Var errVar,CL.mkApply("clSetKernelArg",          val globalArgument = CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,CL.mkApply("clSetKernelArg",
448                                                                  [CL.E_Var kernelVar,                                                                  [CL.mkVar kernelVar,
449                                                                   CL.E_Int(count,CL.intTy),                                                                   CL.mkInt(count,CL.intTy),
450                                                                   CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
451                                                                   CL.E_UnOp(CL.%&,CL.E_Var(concat[RN.globalsVarName,"_cl"]))]))                                                                   CL.mkUnOp(CL.%&,CL.mkVar(concat[RN.globalsVarName,"_cl"]))])))
452    
453          fun genDataArguments([],_,_,_) = []          fun genDataArguments([],_,_,_) = []
454            | genDataArguments((var,nDims)::globals,count,kernelVar,errVar) =            | genDataArguments((var,nDims)::globals,count,kernelVar,errVar) =
455    
456                  CL.mkAssign(CL.E_Var errVar,CL.mkApply("clSetKernelArg",                  CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=, CL.mkApply("clSetKernelArg",
457                                                                  [CL.E_Var kernelVar,                                                                  [CL.mkVar kernelVar,
458                                                                   CL.E_Int(count,CL.intTy),                                                                   CL.mkInt(count,CL.intTy),
459                                                                   CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
460                                                                   CL.E_UnOp(CL.%&,CL.E_Var(concat[var,"_cl"]))]))::                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffix var))])))::
461    
462                          CL.mkAssign(CL.E_Var errVar,CL.mkApply("clSetKernelArg",                          CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,CL.mkApply("clSetKernelArg",
463                                                                  [CL.E_Var kernelVar,                                                                  [CL.mkVar kernelVar,
464                                                                   CL.E_Int((count + 1),CL.intTy),                                                                   CL.mkInt((count + 1),CL.intTy),
465                                                                   CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
466                                                                   CL.E_UnOp(CL.%&,CL.E_Var(concat[var,"_cl", IntegerLit.toString (count + 1)]))])):: genDataArguments (globals, count + 2,kernelVar,errVar)                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffixData var))]))):: genDataArguments (globals, count + 2,kernelVar,errVar)
467    
468          in          in
469    
470                  [globalArgument] @ genDataArguments(globals,count + 1,kernelVar,errVar)                  [globalArgument] @ genDataArguments(globals,count + 1,kernelVar,errVar)
471    
472          end          end
473    
474          (* generates the main function of host code *)          (* generates the main function of host code *)
475          fun genHostMain() = let          fun genHostMain() = let
476                  val setupCall = [CL.mkCall(RN.setupFName,[CL.E_Var RN.globalsVarName])]                val setupCall = [CL.mkCall(RN.setupFName,[CL.mkVar RN.globalsVarName])]
477                  val globalsDecl = CL.mkDecl(CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName,SOME(CL.I_Exp(CL.mkApply("malloc",                val globalsDecl = CL.mkDecl(
478                                                                          [CL.mkApply("sizeof",[CL.E_Var RN.globalsTy])]))))                      CL.T_Ptr(CL.T_Named RN.globalsTy),
479                  val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.E_Var RN.globalsVarName])                      RN.globalsVarName,
480                  val returnStm = [CL.mkReturn(SOME(CL.E_Int(0,CL.intTy)))]                      SOME(CL.I_Exp(CL.mkApply("malloc", [CL.mkApply("sizeof",[CL.mkVar RN.globalsTy])]))))
481                  val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.mkVar RN.globalsVarName])
482                  val returnStm = [CL.mkReturn(SOME(CL.mkInt(0,CL.intTy)))]
483                  val params = [                  val params = [
484                           CL.PARAM([],CL.intTy, "argc"),                           CL.PARAM([],CL.intTy, "argc"),
485                           CL.PARAM([],CL.charArrayPtr,"argv")                           CL.PARAM([],CL.charArrayPtr,"argv")
# Line 396  Line 488 
488                  in                  in
489                    CL.D_Func([],CL.intTy,"main",params,body)                    CL.D_Func([],CL.intTy,"main",params,body)
490                  end                  end
491    
492          (* generates the host-side setup function *)          (* generates the host-side setup function *)
493          fun genHostSetupFunc(strand as Strand{name,tyName,...}, filename, nDims, initially, imgGlobals, oneDim, twoDim, thirdDim) = let          fun genHostSetupFunc (strand as Strand{name,tyName,...}, filename, nDims, initially, imgGlobals) = let
494                  (*Delcare opencl setup objects *)                  (*Delcare opencl setup objects *)
495                  val programVar= "program"                  val programVar= "program"
496                  val kernelVar = "kernel"                  val kernelVar = "kernel"
# Line 421  Line 514 
514                  val platformsVar = "platforms"                  val platformsVar = "platforms"
515                  val numPlatformsVar = "num_platforms"                  val numPlatformsVar = "num_platforms"
516                  val numDevicesVar = "num_devices"                  val numDevicesVar = "num_devices"
517                  val assertStm = CL.mkCall("assert",[CL.mkBinOp(CL.E_Var errVar, CL.#==, CL.E_Var "CL_SUCCESS")])                val assertStm = CL.mkCall("assert",[CL.mkBinOp(CL.mkVar errVar, CL.#==, CL.mkVar "CL_SUCCESS")])
518                  val params = [                  val params = [
519                           CL.PARAM([],CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName)                        CL.PARAM([],CL.T_Named("cl_device_id"), deviceVar)
520                           ]                           ]
521                  val delcarations = [CL.mkDecl(CL.clProgramTy, programVar, NONE),                val declarations = [
522                            CL.mkDecl(CL.clKernelTy, kernelVar, NONE),                      CL.mkDecl(clProgramTy, programVar, NONE),
523                            CL.mkDecl(CL.clCmdQueueTy, cmdVar, NONE),                      CL.mkDecl(clKernelTy, kernelVar, NONE),
524                            CL.mkDecl(CL.clContextTy, contextVar, NONE),                      CL.mkDecl(clCmdQueueTy, cmdVar, NONE),
525                        CL.mkDecl(clContextTy, contextVar, NONE),
526                            CL.mkDecl(CL.intTy, errVar, NONE),                            CL.mkDecl(CL.intTy, errVar, NONE),
527                            CL.mkDecl(CL.intTy, numStrandsVar, NONE),                      CL.mkDecl(CL.intTy, numStrandsVar, SOME(CL.I_Exp(CL.mkInt(1,CL.intTy)))),
                           CL.mkDecl(CL.intTy, numPlatformsVar, NONE),  
528                            CL.mkDecl(CL.intTy, stateSizeVar, NONE),                            CL.mkDecl(CL.intTy, stateSizeVar, NONE),
529                        CL.mkDecl(CL.intTy, "width", NONE),
530                            CL.mkDecl(CL.intTy, imgDataSizeVar, NONE),                            CL.mkDecl(CL.intTy, imgDataSizeVar, NONE),
531                            CL.mkDecl(CL.clDeviceIdTy, deviceVar, NONE),                      (*CL.mkDecl(clDeviceIdTy, deviceVar, NONE), *)
532                            CL.mkDecl(CL.T_Ptr(CL.T_Named tyName), inStateVar,NONE),                            CL.mkDecl(CL.T_Ptr(CL.T_Named tyName), inStateVar,NONE),
533                            CL.mkDecl(CL.clMemoryTy,clInstateVar,NONE),                      CL.mkDecl(clMemoryTy,clInstateVar,NONE),
534                            CL.mkDecl(CL.clMemoryTy,clOutStateVar,NONE),                      CL.mkDecl(clMemoryTy,clOutStateVar,NONE),
535                            CL.mkDecl(CL.T_Ptr(CL.T_Named tyName), outStateVar,NONE),                            CL.mkDecl(CL.T_Ptr(CL.T_Named tyName), outStateVar,NONE),
536                            CL.mkDecl(CL.charPtr, clFNVar,SOME(CL.I_Exp(CL.E_Str filename))),                      CL.mkDecl(CL.charPtr, clFNVar,SOME(CL.I_Exp(CL.mkStr filename))),
537                            CL.mkDecl(CL.charPtr, headerFNVar,SOME(CL.I_Exp(CL.E_Str "Diderot/opencl_types.h"))),  (* FIXME:  use Paths.diderotInclude *)
538                        CL.mkDecl(CL.charPtr, headerFNVar,SOME(CL.I_Exp(CL.mkStr "../src/include/Diderot/cl-types.h"))),
539                            CL.mkDecl(CL.T_Array(CL.charPtr,SOME(2)),sourcesVar,NONE),                            CL.mkDecl(CL.T_Array(CL.charPtr,SOME(2)),sourcesVar,NONE),
540                            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)),globalVar,NONE),
541                            CL.mkDecl(CL.T_Array(CL.T_Named "size_t",SOME(nDims)),localVar,NONE),                            CL.mkDecl(CL.T_Array(CL.T_Named "size_t",SOME(nDims)),localVar,NONE),
542                            CL.mkDecl(CL.intTy,numDevicesVar,SOME(CL.I_Exp(CL.E_Int(~1,CL.intTy)))),                      CL.mkDecl(CL.intTy,numDevicesVar,SOME(CL.I_Exp(CL.mkInt(~1,CL.intTy)))),
543                            CL.mkDecl(CL.T_Array(CL.clDeviceIdTy, SOME(1)), platformsVar, NONE),                      CL.mkDecl(CL.T_Array(CL.T_Named "cl_platform_id", SOME(1)), platformsVar, NONE),
544                            CL.mkDecl(CL.intTy,"num_platforms",SOME(CL.I_Exp(CL.E_Int(~1,CL.intTy))))]                      CL.mkDecl(CL.intTy,"num_platforms",SOME(CL.I_Exp(CL.mkInt(~1,CL.intTy))))
545                    ]
546                  (* Retrieve the platforms *)              (* Setup Global Variables *)
547                  val platformStm = [CL.mkAssign(CL.E_Var errVar, CL.mkApply("clGetPlatformIDs",                val globalsDecl = CL.mkDecl(
548                                                    [CL.E_Int(10,CL.intTy),                      CL.T_Ptr(CL.T_Named RN.globalsTy),
549                                                     CL.E_UnOp(CL.%&,CL.E_Var platformsVar),                      RN.globalsVarName,
550                                                     CL.E_UnOp(CL.%&,CL.E_Var numDevicesVar)])),                      SOME(CL.I_Exp(CL.mkApply("malloc", [CL.mkApply("sizeof",[CL.mkVar RN.globalsTy])]))))
551                  val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.mkVar RN.globalsVarName])
552    
553                    (* Retrieve the platforms
554                    val platformStm = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetPlatformIDs",
555                                                      [CL.mkInt(10,CL.intTy),
556                                                       CL.mkVar platformsVar,
557                                                       CL.mkUnOp(CL.%&,CL.mkVar numPlatformsVar)])),
558                                                     assertStm]                                                     assertStm]
559    
560                  val devicesStm = [CL.mkAssign(CL.E_Var errVar, CL.mkApply("clGetDeviceIDs",                  val devicesStm = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetDeviceIDs",
561                                                    [CL.mkSubscript(CL.E_Var platformsVar,CL.E_Int(0,CL.intTy)),                                                    [CL.mkSubscript(CL.mkVar platformsVar,CL.mkInt(0,CL.intTy)),
562                                                     CL.E_Var "CL_DEVICE_TYPE_GPU",                                                     CL.mkVar "CL_DEVICE_TYPE_GPU",
563                                                     CL.E_Int(1,CL.intTy),                                                     CL.mkInt(1,CL.intTy),
564                                                     CL.E_UnOp(CL.%&,CL.E_Var deviceVar),                                                     CL.mkUnOp(CL.%&,CL.mkVar deviceVar),
565                                                     CL.E_UnOp(CL.%&,CL.E_Var numDevicesVar)])),                                                     CL.mkUnOp(CL.%&,CL.mkVar numDevicesVar)])),
566                                                     assertStm]                                                     assertStm] *)
567    
568                  (* Create Context *)                  (* Create Context *)
569                  val contextStm = [CL.mkAssign(CL.E_Var contextVar, CL.mkApply("clCreateContext",                  val contextStm = [CL.mkAssign(CL.mkVar contextVar, CL.mkApply("clCreateContext",
570                                                    [CL.E_Int(0,CL.intTy),                                                    [CL.mkInt(0,CL.intTy),
571                                                    CL.E_Int(1,CL.intTy),                                                    CL.mkInt(1,CL.intTy),
572                                                    CL.E_UnOp(CL.%&,CL.E_Var deviceVar),                                                    CL.mkUnOp(CL.%&,CL.mkVar deviceVar),
573                                                    CL.E_Var "NULL",                                                    CL.mkVar "NULL",
574                                                    CL.E_Var "NULL",                                                    CL.mkVar "NULL",
575                                                    CL.E_UnOp(CL.%&,CL.E_Var errVar)])),                                                    CL.mkUnOp(CL.%&,CL.mkVar errVar)])),
576                                                    assertStm]                                                    assertStm]
577    
578                  (* Create Command Queue *)                  (* Create Command Queue *)
579                  val commandStm = [CL.mkAssign(CL.E_Var cmdVar, CL.mkApply("clCreateCommandQueue",                  val commandStm = [CL.mkAssign(CL.mkVar cmdVar, CL.mkApply("clCreateCommandQueue",
580                                                    [CL.E_Var contextVar,                                                    [CL.mkVar contextVar,
581                                                    CL.E_Var deviceVar,                                                    CL.mkVar deviceVar,
582                                                    CL.E_Int(0,CL.intTy),                                                    CL.mkInt(0,CL.intTy),
583                                                    CL.E_UnOp(CL.%&,CL.E_Var errVar)])),                                                    CL.mkUnOp(CL.%&,CL.mkVar errVar)])),
584                                                    assertStm]                                                    assertStm]
585    
586    
587                    (*Create Program/Build/Kernel with Source statement *)
588                    val createProgStm = CL.mkAssign(CL.mkVar programVar, CL.mkApply("clCreateProgramWithSource",
589                                                                                                                    [CL.mkVar contextVar,
590                                                                                                                     CL.mkInt(2,CL.intTy),
591                                                                                                                     CL.mkCast(CL.T_Ptr(CL.T_Named("const char *")),CL.mkUnOp(CL.%&,CL.mkVar sourcesVar)),
592                                                                                                                     CL.mkVar "NULL",
593                                                                                                                     CL.mkUnOp(CL.%&,CL.mkVar errVar)]))
594    
595                    (* FIXME: Remove after testing purposes, Build Log for OpenCL*)
596                    val buildLog = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clBuildProgram",
597                                                                                                                    [CL.mkVar programVar,
598                                                                                                                     CL.mkInt(0,CL.intTy),
599                                                                                                                     CL.mkVar "NULL",
600                                                                                                                     CL.mkVar "NULL",
601                                                                                                                     CL.mkVar "NULL",
602                                                                                                                     CL.mkVar "NULL"])),
603                                              CL.mkDecl(CL.charPtr, "build", NONE),
604                                              CL.mkDecl(CL.T_Named("size_t"),"ret_val_size",NONE),
605                                               CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetProgramBuildInfo",
606                                                                                                                    [CL.mkVar programVar,
607                                                                                                                    CL.mkVar deviceVar,
608                                                                                                                     CL.mkVar "CL_PROGRAM_BUILD_LOG",
609                                                                                                                     CL.mkInt(0,CL.intTy),
610                                                                                                                     CL.mkVar "NULL",
611                                                                                                                     CL.mkUnOp(CL.%&,CL.mkVar "ret_val_size")])),
612                                              CL.mkAssign(CL.mkVar "build", CL.mkApply("malloc", [CL.mkVar "ret_val_size"])),
613                                                    CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetProgramBuildInfo",
614                                                                                                                    [CL.mkVar programVar,
615                                                                                                                    CL.mkVar deviceVar,
616                                                                                                                     CL.mkVar "CL_PROGRAM_BUILD_LOG",
617                                                                                                                     CL.mkVar "ret_val_size",
618                                                                                                                     CL.mkVar "build",
619                                                                                                                     CL.mkVar "NULL"])),
620                                                    CL.mkAssign(CL.mkSubscript(CL.mkVar "build",CL.mkVar "ret_val_size"),CL.mkVar ("'\\" ^ "0'")),
621                                                    CL.mkCall("printf",[CL.mkStr ( "Build Log:" ^ "\n" ^ "%s" ^ "\n"), CL.mkVar "build"])]
622    
623    
624    
625    
626                    val createKernel = CL.mkAssign(CL.mkVar kernelVar, CL.mkApply("clCreateKernel",
627                                                                                                                    [CL.mkVar programVar,
628                                                                                                                     CL.mkStr RN.kernelFuncName,
629                                                                                                                     CL.mkUnOp(CL.%&,CL.mkVar errVar)]))
630    
631    
632                    val create_build_stms = [createProgStm,assertStm] @ buildLog @ [assertStm,createKernel,assertStm]
633    
634    
635    
636                  (* Create Memory Buffers for Strand States and Globals *)                  (* Create Memory Buffers for Strand States and Globals *)
637                  val strandSize = CL.mkAssign(CL.E_Var stateSizeVar,CL.mkBinOp(CL.mkApply("sizeof",                  val strandSize = CL.mkAssign(CL.mkVar stateSizeVar,CL.mkBinOp(CL.mkApply("sizeof",
638                                                                          [CL.E_Var tyName]), CL.#*,CL.E_Var numStrandsVar))                                                                          [CL.mkVar tyName]), CL.#*,CL.mkVar numStrandsVar))
639                  val strandObjects = [CL.mkAssign(CL.E_Var inStateVar, CL.mkApply("malloc",  
640                                                                                  [CL.E_Var stateSizeVar])),                  val clStrandObjects = [CL.mkAssign(CL.mkVar clInstateVar, CL.mkApply("clCreateBuffer",
641                                                          CL.mkAssign(CL.E_Var outStateVar, CL.mkApply("malloc",                                                                  [CL.mkVar contextVar,
642                                                                                  [CL.E_Var stateSizeVar]))]                                                                  CL.mkVar "CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR",
643                                                                    CL.mkVar stateSizeVar,
644                  val clStrandObjects = [CL.mkAssign(CL.E_Var clInstateVar, CL.mkApply("clCreateBuffer",                                                                  CL.mkVar "NULL",
645                                                                  [CL.E_Var contextVar,                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)])),
646                                                                  CL.E_Var "CL_MEM_READ_WRITE",                                                           CL.mkAssign(CL.mkVar clOutStateVar, CL.mkApply("clCreateBuffer",
647                                                                  CL.E_Var stateSizeVar,                                                                  [CL.mkVar contextVar,
648                                                                  CL.E_Var "NULL",                                                                  CL.mkVar "CL_MEM_READ_WRITE",
649                                                                  CL.E_UnOp(CL.%&,CL.E_Var errVar)])),                                                                  CL.mkVar stateSizeVar,
650                                                           CL.mkAssign(CL.E_Var clOutStateVar, CL.mkApply("clCreateBuffer",                                                                  CL.mkVar "NULL",
651                                                                  [CL.E_Var contextVar,                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))]
652                                                                  CL.E_Var "CL_MEM_READ_WRITE",  
653                                                                  CL.E_Var stateSizeVar,  
654                                                                  CL.E_Var "NULL",                  (* Setup up selfOut variable *)
655                                                                  CL.E_UnOp(CL.%&,CL.E_Var errVar)]))]                  val strandsArrays = [CL.mkAssign(CL.mkVar outStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.mkVar numStrandsVar,
656                                                                            CL.#*, CL.mkApply("sizeof",[CL.mkVar tyName]))])),
657                                                                    CL.mkAssign(CL.mkVar inStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.mkVar numStrandsVar,
658                                                                            CL.#*, CL.mkApply("sizeof",[CL.mkVar tyName]))]))]
659    
660    
661                    (* Initialize Width Parameter *)
662                    val widthDel = if nDims = 2 then
663                              CL.mkAssign(CL.mkVar "width",CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1, CL.intTy)))
664                       else
665                              CL.mkAssign(CL.mkVar "width",CL.mkInt(0,CL.intTy))
666    
667    
668                    val strands_init = CL.mkCall(RN.strandInitSetup, [
669                            CL.mkVar "size", CL.mkVar "width", CL.mkVar inStateVar
670                          ])
671    
672              val clGlobalBuffers = getGlobalDataBuffers(!imgGlobals,3,contextVar,errVar)              val clGlobalBuffers = getGlobalDataBuffers(!imgGlobals,3,contextVar,errVar)
673    
674    
675                  (* Load the Kernel and Header Files *)                  (* Load the Kernel and Header Files *)
676                  val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.E_Var sourcesVar,CL.E_Int(0,CL.intTy)),                  val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(1,CL.intTy)),
677                                                                            CL.mkApply(RN.clLoaderFN, [CL.E_Var clFNVar])),                                                                            CL.mkApply(RN.clLoaderFN, [CL.mkVar clFNVar])),
678                                                    CL.mkAssign(CL.mkSubscript(CL.E_Var sourcesVar,CL.E_Int(1,CL.intTy)),             CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(0,CL.intTy)),
679                                                                            CL.mkApply(RN.clLoaderFN, [CL.E_Var headerFNVar]))]                                                                            CL.mkApply(RN.clLoaderFN, [CL.mkVar headerFNVar]))]
680    
681                    (* val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(1,CL.intTy)),
682                                                                              CL.mkApply(RN.clLoaderFN, [CL.mkVar clFNVar]))] *)
683    
684    
685                  (* Created Enqueue Statements *)                  (* Created Enqueue Statements *)
686    (* FIXME: simplify this code by function abstraction *)
687                  val enqueueStm = if nDims = 1                  val enqueueStm = if nDims = 1
688                          then [CL.mkAssign(CL.E_Var errVar,                          then [CL.mkAssign(CL.mkVar errVar,
689                                                            CL.mkApply("clEnqueueNDRangeKernel",                                                            CL.mkApply("clEnqueueNDRangeKernel",
690                                                                                                  [CL.E_Var cmdVar,                                                                                                  [CL.mkVar cmdVar,
691                                                                                                   CL.E_Var kernelVar,                                                                                                   CL.mkVar kernelVar,
692                                                                                                   CL.E_Int(1,CL.intTy),                                                                                                   CL.mkInt(1,CL.intTy),
693                                                                                                   CL.E_Var "NULL",                                                                                                   CL.mkVar "NULL",
694                                                                                                   CL.E_Var globalVar,                                                                                                   CL.mkVar globalVar,
695                                                                                                   CL.E_Var localVar,                                                                                                   CL.mkVar localVar,
696                                                                                                   CL.E_Int(0,CL.intTy),                                                                                                   CL.mkInt(0,CL.intTy),
697                                                                                                   CL.E_Var "NULL",                                                                                                   CL.mkVar "NULL",
698                                                                                                   CL.E_Var "NULL"])),CL.mkCall("clFinish",[CL.E_Var cmdVar])]                                                                                                   CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]
699                          else if nDims = 2  then                          else if nDims = 2  then
700                           [CL.mkAssign(CL.E_Var errVar,                           [CL.mkAssign(CL.mkVar errVar,
701                                                          CL.mkApply("clEnqueueNDRangeKernel",                                                          CL.mkApply("clEnqueueNDRangeKernel",
702                                                                                                  [CL.E_Var cmdVar,                                                                                                  [CL.mkVar cmdVar,
703                                                                                                   CL.E_Var kernelVar,                                                                                                   CL.mkVar kernelVar,
704                                                                                                   CL.E_Int(2,CL.intTy),                                                                                                   CL.mkInt(2,CL.intTy),
705                                                                                                   CL.E_Var "NULL",                                                                                                   CL.mkVar "NULL",
706                                                                                                   CL.E_Var globalVar,                                                                                                   CL.mkVar globalVar,
707                                                                                                   CL.E_Var localVar,                                                                                                   CL.mkVar localVar,
708                                                                                                   CL.E_Int(0,CL.intTy),                                                                                                   CL.mkInt(0,CL.intTy),
709                                                                                                   CL.E_Var "NULL",                                                                                                   CL.mkVar "NULL",
710                                                                                                   CL.E_Var "NULL"])),CL.mkCall("clFinish",[CL.E_Var cmdVar])]                                                                                                   CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]
711                          else                          else
712                            [CL.mkAssign(CL.E_Var errVar,                            [CL.mkAssign(CL.mkVar errVar,
713                                                          CL.mkApply("clEnqueueNDRangeKernel",                                                          CL.mkApply("clEnqueueNDRangeKernel",
714                                                                                                  [CL.E_Var cmdVar,                                                                                                  [CL.mkVar cmdVar,
715                                                                                                   CL.E_Var kernelVar,                                                                                                   CL.mkVar kernelVar,
716                                                                                                   CL.E_Int(3,CL.intTy),                                                                                                   CL.mkInt(3,CL.intTy),
717                                                                                                   CL.E_Var "NULL",                                                                                                   CL.mkVar "NULL",
718                                                                                                   CL.E_Var globalVar,                                                                                                   CL.mkVar globalVar,
719                                                                                                   CL.E_Var localVar,                                                                                                   CL.mkVar localVar,
720                                                                                                   CL.E_Int(0,CL.intTy),                                                                                                   CL.mkInt(0,CL.intTy),
721                                                                                                   CL.E_Var "NULL",                                                                                                   CL.mkVar "NULL",
722                                                                                                   CL.E_Var "NULL"])),CL.mkCall("clFinish",[CL.E_Var cmdVar])]                                                                                                   CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]
723    
724    
                 (* Setup up selfOut variable *)  
                 val selfOutStm = CL.mkAssign(CL.E_Var outStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.E_Var numStrandsVar,  
                                                                         CL.#*, CL.mkApply("sizeof",[CL.E_Var tyName]))]))  
725    
726                  (* Setup Global and Local variables *)                  (* Setup Global and Local variables *)
727    
728                  val globalAndlocalStms = if nDims = 1 then                  val globalAndlocalStms = if nDims = 1 then
729                          [CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(0,CL.intTy)),                          [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),
730                                                                     CL.mkSubscript(CL.E_Var "size", CL.E_Int(0,CL.intTy))),                                                                     CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),
731                           CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(0,CL.intTy)),                           CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),
732                                                                    CL.E_Var "16")]                                                                    CL.mkVar "16")]
733    
734    
735                  else if nDims = 2 then                  else if nDims = 2 then
736                          [CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(0,CL.intTy)),                          [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),
737                                                                     CL.mkSubscript(CL.E_Var "size", CL.E_Int(0,CL.intTy))),                                                                     CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),
738                          CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(1,CL.intTy)),                          CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1,CL.intTy)),
739                                                                     CL.mkSubscript(CL.E_Var "sizes", CL.E_Int(1,CL.intTy))),                                                                     CL.mkSubscript(CL.mkVar "size", CL.mkInt(1,CL.intTy))),
740                          CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(0,CL.intTy)),                          CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),
741                                                                    CL.E_Var "16"),                                                                    CL.mkVar "16"),
742                          CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(1,CL.intTy)),                          CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(1,CL.intTy)),
743                                                                    CL.E_Var "16")]                                                                    CL.mkVar "16")]
744    
745                  else                  else
746                          [CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(0,CL.intTy)),                          [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),
747                                                                     CL.mkSubscript(CL.E_Var "size", CL.E_Int(0,CL.intTy))),                                                                     CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),
748                          CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(1,CL.intTy)),                          CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1,CL.intTy)),
749                                                                     CL.mkSubscript(CL.E_Var "size", CL.E_Int(1,CL.intTy))),                                                                     CL.mkSubscript(CL.mkVar "size", CL.mkInt(1,CL.intTy))),
750                          CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(2,CL.intTy)),                          CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(2,CL.intTy)),
751                                                                     CL.mkSubscript(CL.E_Var "size", CL.E_Int(2,CL.intTy))),                                                                     CL.mkSubscript(CL.mkVar "size", CL.mkInt(2,CL.intTy))),
752                          CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(0,CL.intTy)),                          CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),
753                                                                    CL.E_Var "16"),                                                                    CL.mkVar "16"),
754                          CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(1,CL.intTy)),                          CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(1,CL.intTy)),
755                                                                    CL.E_Var "16"),                                                                    CL.mkVar "16"),
756                          CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(2,CL.intTy)),                          CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(2,CL.intTy)),
757                                                                    CL.E_Var "16")]                                                                    CL.mkVar "16")]
758    
759    
760    
761                  (* Setup Kernel arguments *)                  (* Setup Kernel arguments *)
762                  val kernelArguments = [CL.mkAssign(CL.E_Var errVar,CL.mkApply("clSetKernelArg",                  val kernelArguments = [CL.mkAssign(CL.mkVar errVar,CL.mkApply("clSetKernelArg",
763                                                                  [CL.E_Var kernelVar,                                                                  [CL.mkVar kernelVar,
764                                                                   CL.E_Int(0,CL.intTy),                                                                   CL.mkInt(0,CL.intTy),
765                                                                   CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
766                                                                   CL.E_UnOp(CL.%&,CL.E_Var clInstateVar)])),                                                                   CL.mkUnOp(CL.%&,CL.mkVar clInstateVar)])),
767                                                              CL.mkExpStm(CL.mkAssignOp(CL.E_Var errVar, CL.|=,CL.mkApply("clSetKernelArg",                                                              CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar, CL.|=,CL.mkApply("clSetKernelArg",
768                                                                  [CL.E_Var kernelVar,                                                                  [CL.mkVar kernelVar,
769                                                                   CL.E_Int(1,CL.intTy),                                                                   CL.mkInt(1,CL.intTy),
770                                                                   CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
771                                                                   CL.E_UnOp(CL.%&,CL.E_Var clOutStateVar)]))),                                                                   CL.mkUnOp(CL.%&,CL.mkVar clOutStateVar)]))),
772                                                                    CL.mkExpStm(CL.mkAssignOp(CL.E_Var errVar, CL.|=,CL.mkApply("clSetKernelArg",                                                                    CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar, CL.|=,CL.mkApply("clSetKernelArg",
773                                                                  [CL.E_Var kernelVar,                                                                  [CL.mkVar kernelVar,
774                                                                   CL.E_Int(2,CL.intTy),                                                                   CL.mkInt(2,CL.intTy),
775                                                                   CL.mkApply("sizeof",[CL.E_Var "int"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "int"]),
776                                                                   CL.E_UnOp(CL.%&,CL.E_Var "width")])))]                                                                   CL.mkUnOp(CL.%&,CL.mkVar "width")])))]
777    
778             val clGlobalArguments = genGlobalArguments(!imgGlobals,3,kernelVar,errVar)             val clGlobalArguments = genGlobalArguments(!imgGlobals,3,kernelVar,errVar) @ [assertStm]
779    
780                  (* Retrieve output *)                  (* Retrieve output *)
781                  val outputStm = CL.mkAssign(CL.E_Var errVar,                  val outputStm = CL.mkAssign(CL.mkVar errVar,
782                                                          CL.mkApply("clEnqueueReadBuffer",                                                          CL.mkApply("clEnqueueReadBuffer",
783                                                                                                  [CL.E_Var cmdVar,                                                                                                  [CL.mkVar cmdVar,
784                                                                                                   CL.E_Var clOutStateVar,                                                                                                   CL.mkVar clOutStateVar,
785                                                                                                   CL.E_Var "CL_TRUE",                                                                                                   CL.mkVar "CL_TRUE",
786                                                                                                   CL.E_Int(0,CL.intTy),                                                                                                   CL.mkInt(0,CL.intTy),
787                                                                                                   CL.E_Var stateSizeVar,                                                                                                   CL.mkVar stateSizeVar,
788                                                                                                   CL.E_Var outStateVar,                                                                                                   CL.mkVar outStateVar,
789                                                                                                   CL.E_Int(0,CL.intTy),                                                                                                   CL.mkInt(0,CL.intTy),
790                                                                                                   CL.E_Var "NULL",                                                                                                   CL.mkVar "NULL",
791                                                                                                   CL.E_Var "NULL"]))                                                                                                   CL.mkVar "NULL"]))
792    
793                  (* Free all the objects *)                  (* Free all the objects *)
794                  val freeStms = [CL.mkCall("clReleaseKernel",[CL.E_Var kernelVar]),                  val freeStms = [CL.mkCall("clReleaseKernel",[CL.mkVar kernelVar]),
795                                                  CL.mkCall("clReleaseProgram",[CL.E_Var programVar ]),                                                  CL.mkCall("clReleaseProgram",[CL.mkVar programVar ]),
796                                                  CL.mkCall("clReleaseCommandQueue",[CL.E_Var cmdVar]),                                                  CL.mkCall("clReleaseCommandQueue",[CL.mkVar cmdVar]),
797                                                  CL.mkCall("clReleaseContext",[CL.E_Var contextVar]),                                                  CL.mkCall("clReleaseContext",[CL.mkVar contextVar]),
798                                                  CL.mkCall("clReleaseMemObject",[CL.E_Var clInstateVar]),                                                  CL.mkCall("clReleaseMemObject",[CL.mkVar clInstateVar]),
799                                                  CL.mkCall("clReleaseMemObject",[CL.E_Var clOutStateVar])]                                                  CL.mkCall("clReleaseMemObject",[CL.mkVar clOutStateVar])]
800    
801    
802                    (*Setup Strand Print Function *)
803                    val outputData = [CL.mkDecl(CL.T_Ptr(CL.T_Named("FILE")), "outS", SOME(CL.I_Exp(CL.mkApply("fopen",
804                                                    [CL.mkStr "mip.txt",
805                                                    CL.mkStr "w"])))),
806                                                    CL.mkCall(concat[name, "_print"],
807                                                                            [CL.mkVar "outS",
808                                                                             CL.mkVar "size",
809                                                                             CL.mkVar "width",
810                                                                             CL.mkVar outStateVar])]
811    
812    
813    
814                  (* Body put all the statments together *)                  (* Body put all the statments together *)
815                  val body =  delcarations @ platformStm @ devicesStm @ contextStm @ commandStm @ !initially @ [strandSize] @                  val body =  declarations @ [globalsDecl,initGlobalsCall] (*@ platformStm @ devicesStm *) @ contextStm @ commandStm @ !initially @ [strandSize] @
816                                     clStrandObjects @ clGlobalBuffers @ sourceStms  @ [selfOutStm] @ globalAndlocalStms @                                     strandsArrays @ globalAndlocalStms @ [widthDel,strands_init]  @ clStrandObjects @ clGlobalBuffers @ sourceStms  @ create_build_stms  (*@
817                                     kernelArguments @ clGlobalArguments @ enqueueStm @  [outputStm] @ freeStms                                     kernelArguments @ clGlobalArguments @ enqueueStm @  [outputStm] @ freeStms @ outputData *)
818    
819                  in                  in
820    
821                  CL.D_Func([],CL.voidTy,RN.setupFName,params,CL.mkBlock(body))                  CL.D_Func([],CL.voidTy,RN.setupFName,params,CL.mkBlock(body))
822    
823                  end                  end
824    (* generate the data and global parameters *)
825            fun genKeneralGlobalParams ((name,tyname)::rest) =
826                    CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), concat[RN.globalsVarName]) ::
827                    CL.PARAM([], CL.T_Ptr(CL.T_Named (RN.imageTy tyname)),RN.addBufferSuffix name) ::
828                    CL.PARAM([], CL.T_Ptr(CL.voidTy),RN.addBufferSuffixData name) ::
829                    genKeneralGlobalParams(rest)
830    
831              | genKeneralGlobalParams ([]) = []
832    
833            (*generate code for intilizing kernel global data *)
834            fun initKernelGlobals (globals,imgGlobals) = let
835                    fun initGlobalStruct (CL.D_Var(_, _ , name, _)::rest) =
836                                    CL.mkAssign(CL.mkVar name, CL.mkIndirect(CL.mkVar RN.globalsVarName, name)) ::
837                                    initGlobalStruct(rest)
838                      | initGlobalStruct ( _::rest) = initGlobalStruct(rest)
839                      | initGlobalStruct([]) = []
840    
841                    fun initGlobalImages((name,tyname)::rest) =
842                                    CL.mkAssign(CL.mkVar name, CL.mkVar (RN.addBufferSuffix name)) ::
843                                    CL.mkAssign(CL.mkIndirect(CL.mkVar name,"data"),CL.mkVar (RN.addBufferSuffixData name)) ::
844                                    initGlobalImages(rest)
845                      | initGlobalImages([]) = []
846                    in
847                      initGlobalStruct(globals) @ initGlobalImages(imgGlobals)
848                    end
849    
850          (* generate the main kernel function for the .cl file *)          (* generate the main kernel function for the .cl file *)
851          fun genKernelFun(Strand{name, tyName, state, output, code},nDims) = let          fun genKernelFun(Strand{name, tyName, state, output, code,...},nDims,globals,imgGlobals) = let
852                   val fName = RN.kernelFuncName;                   val fName = RN.kernelFuncName;
853                   val inState = "strand_in"                   val inState = "strand_in"
854                   val outState = "strand_out"                   val outState = "strand_out"
# Line 649  Line 856 
856                        CL.PARAM(["__global"], CL.T_Ptr(CL.T_Named tyName), "selfIn"),                        CL.PARAM(["__global"], CL.T_Ptr(CL.T_Named tyName), "selfIn"),
857                        CL.PARAM(["__global"], CL.T_Ptr(CL.T_Named tyName), "selfOut"),                        CL.PARAM(["__global"], CL.T_Ptr(CL.T_Named tyName), "selfOut"),
858                        CL.PARAM(["__global"], CL.intTy, "width")                        CL.PARAM(["__global"], CL.intTy, "width")
859                      ]                      ] @ genKeneralGlobalParams(!imgGlobals)
860                    val thread_ids = if nDims = 1                    val thread_ids = if nDims = 1
861                          then [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy)))),                          then [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),
862                                    CL.mkAssign(CL.E_Var "x",CL.mkApply(RN.getGlobalThreadId,[CL.E_Int(0,CL.intTy)]))]                                    CL.mkAssign(CL.mkVar "x",CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(0,CL.intTy)]))]
863                          else                          else
864                                  [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy)))),                                  [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),
865                                   CL.mkDecl(CL.intTy, "y", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy)))),                                   CL.mkDecl(CL.intTy, "y", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),
866                                    CL.mkAssign(CL.E_Var "x",  CL.mkApply(RN.getGlobalThreadId,[CL.E_Int(0,CL.intTy)])),                                    CL.mkAssign(CL.mkVar "x",  CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(0,CL.intTy)])),
867                                    CL.mkAssign(CL.E_Var "y",CL.mkApply(RN.getGlobalThreadId,[CL.E_Int(1,CL.intTy)]))]                                    CL.mkAssign(CL.mkVar "y",CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(1,CL.intTy)]))]
868    
869                    val strandDecl = [CL.mkDecl(CL.T_Named tyName, inState, NONE),                    val strandDecl = [CL.mkDecl(CL.T_Named tyName, inState, NONE),
870                                                          CL.mkDecl(CL.T_Named tyName, outState,NONE)]                                                          CL.mkDecl(CL.T_Named tyName, outState,NONE)]
871                    val strandObjects  = if nDims = 1                    val strandObjects  = if nDims = 1
872                          then [CL.mkAssign(CL.mkSubscript(CL.E_Var "selfIn",CL.E_Str "x"),                          then [CL.mkAssign(CL.mkSubscript(CL.mkVar "selfIn",CL.mkStr "x"),
873                                                                           CL.E_Var inState),                                                                           CL.mkVar inState),
874                                    CL.mkAssign(CL.mkSubscript(CL.E_Var "selfOut",CL.E_Str "x"),                                    CL.mkAssign(CL.mkSubscript(CL.mkVar "selfOut",CL.mkStr "x"),
875                                                                           CL.E_Var outState)]                                                                           CL.mkVar outState)]
876                          else let                          else let
877                                  val index = CL.mkBinOp(CL.mkBinOp(CL.E_Var "y",CL.#*,CL.E_Var "width"),CL.#+,CL.E_Var "x")                                  val index = CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")
878                                  in                                  in
879                                          [CL.mkAssign(CL.mkSubscript(CL.E_Var "selfIn",index),                                          [CL.mkAssign(CL.mkSubscript(CL.mkVar "selfIn",index),
880                                                                          CL.E_Var inState),                                                                          CL.mkVar inState),
881                                           CL.mkAssign(CL.mkSubscript(CL.E_Var "selfOut",index),                                           CL.mkAssign(CL.mkSubscript(CL.mkVar "selfOut",index),
882                                                                          CL.E_Var outState)]                                                                          CL.mkVar outState)]
883                                  end                                  end
884                    val status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy))))  
885                    val strand_init_function = CL.mkCall(RN.strandInit name, [CL.E_Var inState])  
886                    val local_vars =  thread_ids @ strandDecl @ strandObjects @ [status,strand_init_function]                    val status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy))))
887                    val while_exp = CL.mkBinOp(CL.E_Var "status",CL.#!=, CL.E_Var RN.kStabilize)                    val local_vars =  thread_ids @ initKernelGlobals(!globals,!imgGlobals)  @ strandDecl @ strandObjects @ [status]
888                    val while_body = [CL.mkAssign(CL.E_Var "status", CL.mkApply(RN.strandUpdate name,[CL.E_Var inState,CL.E_Var outState])),                    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))
889                                                          CL.mkCall(RN.strandStabilize name,[CL.E_Var inState,CL.E_Var outState]),                    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)])),
890                                                          CL.mkIfThen(CL.mkBinOp(CL.E_Var "status",CL.#==, CL.E_Var RN.kStabilize),CL.mkBreak)]                                                          CL.mkCall(RN.strandStabilize name,[ CL.mkUnOp(CL.%&,CL.mkVar inState),  CL.mkUnOp(CL.%&,CL.mkVar outState)])]
891    
892                    val whileBlock = [CL.mkWhile(while_exp,CL.mkBlock while_body)]                    val whileBlock = [CL.mkWhile(while_exp,CL.mkBlock while_body)]
893    
# Line 696  Line 903 
903                   in                   in
904                          CL.D_StructDef(getGlobals(globals),RN.globalsTy)                          CL.D_StructDef(getGlobals(globals),RN.globalsTy)
905                    end                    end
906    
907        (* generate the table of strand descriptors *)        (* generate the table of strand descriptors *)
908          fun genStrandTable (ppStrm, strands) = let          fun genStrandTable (ppStrm, strands) = let
909                val nStrands = length strands                val nStrands = length strands
910                fun genInit (Strand{name, ...}) = CL.I_Exp(CL.mkUnOp(CL.%&, CL.E_Var(RN.strandDesc name)))                fun genInit (Strand{name, ...}) = CL.I_Exp(CL.mkUnOp(CL.%&, CL.mkVar(RN.strandDesc name)))
911                fun genInits (_, []) = []                fun genInits (_, []) = []
912                  | genInits (i, s::ss) = (i, genInit s) :: genInits(i+1, ss)                  | genInits (i, s::ss) = (i, genInit s) :: genInits(i+1, ss)
913                fun ppDecl dcl = PrintAsC.output(ppStrm, dcl)                fun ppDecl dcl = PrintAsC.output(ppStrm, dcl)
914                in                in
915                  ppDecl (CL.D_Var([], CL.int32, RN.numStrands,                  ppDecl (CL.D_Var([], CL.int32, RN.numStrands,
916                    SOME(CL.I_Exp(CL.E_Int(IntInf.fromInt nStrands, CL.int32)))));                    SOME(CL.I_Exp(CL.mkInt(IntInf.fromInt nStrands, CL.int32)))));
917                  ppDecl (CL.D_Var([],                  ppDecl (CL.D_Var([],
918                    CL.T_Array(CL.T_Ptr(CL.T_Named RN.strandDescTy), SOME nStrands),                    CL.T_Array(CL.T_Ptr(CL.T_Named RN.strandDescTy), SOME nStrands),
919                    RN.strands,                    RN.strands,
920                    SOME(CL.I_Array(genInits (0, strands)))))                    SOME(CL.I_Array(genInits (0, strands)))))
921                end                end
922    
923          fun genSrc (baseName, Prog{globals, topDecls, strands, initially,imgGlobals,numDims,oneDim,twoDim,thirdDim,...}) = let  
924            fun genSrc (baseName, Prog{double,globals, topDecls, strands, initially,imgGlobals,numDims,...}) = let
925                val clFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "cl"}                val clFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "cl"}
926                val cFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "c"}                val cFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "c"}
927                val clOutS = TextIO.openOut clFileName                val clOutS = TextIO.openOut clFileName
928                val cOutS = TextIO.openOut cFileName                val cOutS = TextIO.openOut cFileName
929    (* FIXME: need to use PrintAsC and PrintAsCL *)
930                val clppStrm = PrintAsC.new clOutS                val clppStrm = PrintAsC.new clOutS
931                val cppStrm = PrintAsC.new cOutS                val cppStrm = PrintAsC.new cOutS
932                fun cppDecl dcl = PrintAsC.output(cppStrm, dcl)                fun cppDecl dcl = PrintAsC.output(cppStrm, dcl)
933                fun clppDecl dcl = PrintAsC.output(clppStrm, dcl)                fun clppDecl dcl = PrintAsCL.output(clppStrm, dcl)
934                val strands = AtomTable.listItems strands                val strands = AtomTable.listItems strands
935                val single_strand as Strand{name, tyName, code, ...}= hd(strands)                val [strand as Strand{name, tyName, code,init_code, ...}] = strands
936                in                in
937                  (* Generate the OpenCl file *)
938                    clppDecl (CL.D_Verbatim([
939                        if double
940                          then "#define DIDEROT_DOUBLE_PRECISION"
941                          else "#define DIDEROT_SINGLE_PRECISION",
942                        "#define DIDEROT_TARGET_CL",
943                        "#include \"Diderot/cl-types.h\""
944                      ]));
945                    List.app clppDecl (List.rev (!globals));
946                    clppDecl (genGlobalStruct (!globals));
947                    clppDecl (genStrandTyDef strand);
948                    List.app clppDecl (!code);
949                    clppDecl (genKernelFun (strand,!numDims,globals,imgGlobals));
950              (* Generate the Host file .c *)              (* Generate the Host file .c *)
951              cppDecl (CL.D_Verbatim([ "#include <OpenCL/OpenCL.h>",                  cppDecl (CL.D_Verbatim([
952                                                                   "#include Diderot/diderot.h"]));                      if double
953                          then "#define DIDEROT_DOUBLE_PRECISION"
954                          else "#define DIDEROT_SINGLE_PRECISION",
955                        "#define DIDEROT_TARGET_CL",
956                        "#include \"Diderot/diderot.h\""
957                      ]));
958                  List.app cppDecl (List.rev (!globals));                  List.app cppDecl (List.rev (!globals));
959              cppDecl (genGlobalStruct (!globals));              cppDecl (genGlobalStruct (!globals));
960              cppDecl (genStrandTyDef single_strand);                  cppDecl (genStrandTyDef strand);
961              cppDecl (genKernelLoader());                  cppDecl  (!init_code);
962                    cppDecl (genStrandInit(strand,!numDims));
963                    cppDecl (genStrandPrint(strand,!numDims));
964                    (* cppDecl (genKernelLoader());*)
965              List.app cppDecl (List.rev (!topDecls));              List.app cppDecl (List.rev (!topDecls));
966              cppDecl (genHostSetupFunc (single_strand,clFileName,!numDims,initially,imgGlobals,oneDim,twoDim,thirdDim));                  cppDecl (genHostSetupFunc (strand, clFileName, !numDims, initially, imgGlobals));
             cppDecl (genHostMain());  
   
             (* Generate the OpenCl file *)  
             clppDecl (genGlobalStruct (!globals));  
             clppDecl (genStrandTyDef single_strand);  
             List.app clppDecl (!code);  
             clppDecl (genKernelFun (single_strand,!numDims));  
   
                 (*List.app (fn strand => List.app ppDecl (genStrand strand)) strands;  
                  genStrandTable (ppStrm, strands);  
                 ppDecl (!initially);*)  
   
967                  PrintAsC.close cppStrm;                  PrintAsC.close cppStrm;
968                  PrintAsC.close clppStrm;                  PrintAsCL.close clppStrm;
969                  TextIO.closeOut cOutS;                  TextIO.closeOut cOutS;
970                  TextIO.closeOut clOutS                  TextIO.closeOut clOutS
971                end                end
# Line 767  Line 986 
986              (* generate the loader flags *)              (* generate the loader flags *)
987                val extraLibs = condCons (parallel, #pthread Paths.extraLibs, [])                val extraLibs = condCons (parallel, #pthread Paths.extraLibs, [])
988                val extraLibs = Paths.teemLinkFlags @  #base Paths.extraLibs :: extraLibs                val extraLibs = Paths.teemLinkFlags @  #base Paths.extraLibs :: extraLibs
989                       val extraLibs =  #cl Paths.extraLibs :: extraLibs
990                val rtLib = TargetUtil.runtimeName {                val rtLib = TargetUtil.runtimeName {
991                        target = TargetUtil.TARGET_CL,                        target = TargetUtil.TARGET_CL,
992                        parallel = parallel, double = double, debug = debug                        parallel = parallel, double = double, debug = debug
993                      }                      }
994                val ldOpts = rtLib :: extraLibs                val ldOpts = rtLib :: extraLibs
995                in                in
996                  genSrc (basename, prog)                  genSrc (basename, prog);
997                    RunCC.compile (basename, cflags);
998                    RunCC.link (basename, ldOpts)
999                  end                  end
1000    
                 (*RunCC.compile (basename, cflags);  
                 RunCC.link (basename, ldOpts)*)  
   
   
1001        end        end
1002    
1003    (* strands *)    (* strands *)
# Line 792  Line 1010 
1010                        tyName = RN.strandTy name,                        tyName = RN.strandTy name,
1011                        state = ref [],                        state = ref [],
1012                        output = ref NONE,                        output = ref NONE,
1013                        code = ref []                        code = ref [],
1014                          init_code = ref (CL.D_Comment(["no init code"]))
1015                      }                      }
1016                in                in
1017                  AtomTable.insert strands (strandId, strand);                  AtomTable.insert strands (strandId, strand);
# Line 805  Line 1024 
1024        (* register the strand-state initialization code.  The variables are the strand        (* register the strand-state initialization code.  The variables are the strand
1025         * parameters.         * parameters.
1026         *)         *)
1027          fun init (Strand{name, tyName, code, ...}, params, init) = let          fun init (Strand{name, tyName, code,init_code, ...}, params, init) = let
1028                val fName = RN.strandInit name                val fName = RN.strandInit name
1029                val params =                val params =
1030                      CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfOut") ::                      CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "selfOut") ::
1031                        List.map (fn (ToC.V(ty, x)) => CL.PARAM([], ty, x)) params                        List.map (fn (ToCL.V(ty, x)) => CL.PARAM([], ty, x)) params
1032                val initFn = CL.D_Func([], CL.voidTy, fName, params, init)                val initFn = CL.D_Func([], CL.voidTy, fName, params, init)
1033                in                in
1034                  code := initFn :: !code                  init_code := initFn
1035                end                end
1036    
1037        (* register a strand method *)        (* register a strand method *)
# Line 827  Line 1046 
1046                  code := methFn :: !code                  code := methFn :: !code
1047                end                end
1048    
1049          fun output (Strand{output, ...}, ty, ToC.V(_, x)) = output := SOME(ty, x)          fun output (Strand{output, ...}, ty, ToCL.V(_, x)) = output := SOME(ty, x)
1050    
1051        end        end
1052    

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

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