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

SCM Repository

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

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

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

revision 1286, Tue Jun 7 10:54:18 2011 UTC revision 1309, Sat Jun 11 14:25:12 2011 UTC
# Line 13  Line 13 
13      structure CL = CLang      structure CL = CLang
14      structure RN = RuntimeNames      structure RN = RuntimeNames
15      structure ToCL = TreeToCL      structure ToCL = TreeToCL
16        structure N = CNames
17    
18      (* variable translation *)
19        structure TrVar =
20          struct
21            type env = CL.typed_var TreeIL.Var.Map.map
22            fun lookup (env, x) = (case V.Map.find (env, x)
23                   of SOME(CL.V(_, x')) => x'
24                    | NONE => raise Fail(concat["lookup(_, ", V.name x, ")"])
25                  (* end case *))
26          (* translate a variable that occurs in an l-value context (i.e., as the target of an assignment) *)
27            fun lvalueVar (env, x) = (case V.kind x
28                   of IL.VK_Global => CL.mkVar(lookup(env, x))
29                    | IL.VK_State strand => CL.mkIndirect(CL.mkVar "selfOut", lookup(env, x))
30                    | IL.VK_Local => CL.mkVar(lookup(env, x))
31                  (* end case *))
32          (* translate a variable that occurs in an r-value context *)
33            fun rvalueVar (env, x) = (case V.kind x
34                   of IL.VK_Global => CL.mkVar(lookup(env, x))
35                    | IL.VK_State strand => CL.mkIndirect(CL.mkVar "selfIn", lookup(env, x))
36                    | IL.VK_Local => CL.mkVar(lookup(env, x))
37                  (* end case *))
38          end
39    
40            structure ToC = TreeToCFn (TrVar)
41    
42    (* C variable translation *)    (* C variable translation *)
43      structure TrCVar =      structure TrCVar =
# Line 47  Line 72 
72      val clPlatformIdTy = CL.T_Named "cl_platform_id"      val clPlatformIdTy = CL.T_Named "cl_platform_id"
73      val clMemoryTy = CL.T_Named "cl_mem"      val clMemoryTy = CL.T_Named "cl_mem"
74    
75      (* variable or field that is mirrored between host and GPU *)
76        type mirror_var = {
77                hostTy : CL.ty,             (* variable type on Host (i.e., C type) *)
78                gpuTy : CL.ty,              (* variable's type on GPU (i.e., OpenCL type) *)
79                var : CL.var                (* variable name *)
80              }
81    
82      datatype strand = Strand of {      datatype strand = Strand of {
83          name : string,          name : string,
84          tyName : string,          tyName : string,
85          state : var list ref,          state : mirror_var list ref,
86          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) *)
87          code : CL.decl list ref,          code : CL.decl list ref,
88          init_code: CL.decl ref          init_code: CL.decl ref
# Line 61  Line 93 
93          double : bool,                  (* true for double-precision support *)          double : bool,                  (* true for double-precision support *)
94          parallel : bool,                (* true for multithreaded (or multi-GPU) target *)          parallel : bool,                (* true for multithreaded (or multi-GPU) target *)
95          debug : bool,                   (* true for debug support in executable *)          debug : bool,                   (* true for debug support in executable *)
96          globals : CL.decl list ref,          globals : mirror_var list ref,
97          topDecls : CL.decl list ref,          topDecls : CL.decl list ref,
98          strands : strand AtomTable.hash_table,          strands : strand AtomTable.hash_table,
99          initially : CL.stm list ref,          initially :  CL.decl ref,
100          numDims: int ref,          numDims: int ref,
101          imgGlobals: (string * int) list ref,          imgGlobals: (string * int) list ref,
102          prFn: CL.decl ref          prFn: CL.decl ref
# Line 98  Line 130 
130    (* TreeIL to target translations *)    (* TreeIL to target translations *)
131      structure Tr =      structure Tr =
132        struct        struct
       (* this function is used for the initially clause, so it generates OpenCL *)  
133          fun fragment (ENV{info, vMap, scope}, blk) = let          fun fragment (ENV{info, vMap, scope}, blk) = let
134                val (vMap, stms) = ToCL.trFragment (vMap, blk)                val (vMap, stms) = (case scope
135                         of GlobalScope => ToC.trFragment (vMap, blk)
136                          | _ => ToCL.trFragment (vMap, blk)
137                        (* end case *))
138                in                in
139                  (ENV{info=info, vMap=vMap, scope=scope}, stms)                  (ENV{info=info, vMap=vMap, scope=scope}, stms)
140                end                end
# Line 126  Line 160 
160        struct        struct
161          fun name (ToCL.V(_, name)) = name          fun name (ToCL.V(_, name)) = name
162          fun global (Prog{globals, imgGlobals, ...}, name, ty) = let          fun global (Prog{globals, imgGlobals, ...}, name, ty) = let
163                val ty' = ToCL.trType ty                val x = {hostTy = ToC.trType ty, gpuTy = ToCL.trType ty, var = name}
164                fun isImgGlobal (imgGlobals, Ty.ImageTy(ImageInfo.ImgInfo{dim, ...}), name) =  imgGlobals  := (name,dim):: !imgGlobals                fun isImgGlobal (Ty.ImageTy(ImageInfo.ImgInfo{dim, ...}), name) =
165                  | isImgGlobal (imgGlobals, _, _) =  ()                      imgGlobals  := (name,dim) :: !imgGlobals
166                in                  | isImgGlobal _ =  ()
167                  globals := CL.D_Var([], ty', name, NONE) :: !globals;                in
168                  isImgGlobal(imgGlobals,ty,name);                  globals := x :: !globals;
169                  ToCL.V(ty', name)                  isImgGlobal (ty, name);
170                    ToCL.V(#gpuTy x, name)
171                end                end
172          fun param x = ToCL.V(ToCL.trType(V.ty x), V.name x)          fun param x = ToCL.V(ToCL.trType(V.ty x), V.name x)
173          fun state (Strand{state, ...}, x) = let          fun state (Strand{state, ...}, x) = let
174                val ty' = ToCL.trType(V.ty x)                val ty = V.ty x
175                val x' = ToCL.V(ty', V.name x)                val x' = {hostTy = ToC.trType ty, gpuTy = ToCL.trType ty, var = V.name x}
176                in                in
177                  state := x' :: !state;                  state := x' :: !state;
178                  x'                  ToCL.V(#gpuTy x', #var x')
179                end                end
180        end        end
181    
# Line 179  Line 214 
214                    globals = ref [],                    globals = ref [],
215                    topDecls = ref [],                    topDecls = ref [],
216                    strands = AtomTable.mkTable (16, Fail "strand table"),                    strands = AtomTable.mkTable (16, Fail "strand table"),
217                    initially = ref([CL.S_Comment["missing initially"]]),                    initially = ref(CL.D_Comment["missing initially"]),
218                                    numDims = ref(0),                                    numDims = ref(0),
219                                    imgGlobals = ref[],                                    imgGlobals = ref[],
220                                    prFn = ref(CL.D_Comment(["No Print Function"]))                                    prFn = ref(CL.D_Comment(["No Print Function"]))
221                  })                  })
222        (* register the global initialization part of a program *)        (* register the global initialization part of a program *)
223    (* FIXME: unused code; can this be removed??
224            fun globalIndirects (globals,stms) = let            fun globalIndirects (globals,stms) = let
225                  fun getGlobals (CL.D_Var(_,_,globalVar,_)::rest) =                  fun getGlobals ({name,target as TargetUtil.TARGET_CL}::rest) =
226                        CL.mkAssign(CL.mkIndirect(CL.mkVar RN.globalsVarName,globalVar),CL.mkVar globalVar)                        CL.mkAssign(CL.mkIndirect(CL.mkVar RN.globalsVarName,name),CL.mkVar name)
227                          ::getGlobals rest                          ::getGlobals rest
228                    | getGlobals [] = []                    | getGlobals [] = []
229                    | getGlobals (_::rest) = getGlobals rest                    | getGlobals (_::rest) = getGlobals rest
230                  in                  in
231                    stms @ getGlobals globals                    stms @ getGlobals globals
232                  end                  end
233    *)
234        (* register the code that is used to register command-line options for input variables *)        (* register the code that is used to register command-line options for input variables *)
235          fun inputs (Prog{topDecls, ...}, stm) = let          fun inputs (Prog{topDecls, ...}, stm) = let
236                val inputsFn = CL.D_Func(                val inputsFn = CL.D_Func(
# Line 218  Line 254 
254                in                in
255                  topDecls := shutdownFn :: initFn :: !topDecls                  topDecls := shutdownFn :: initFn :: !topDecls
256                end                end
   
257        (* create and register the initially function for a program *)        (* create and register the initially function for a program *)
258          fun initially {          fun initially {
259                prog = Prog{strands, initially, numDims,...},                prog = Prog{name=progName, strands, initially, ...},
260                isArray : bool,                isArray : bool,
261                iterPrefix : stm list,                iterPrefix : stm list,
262                iters : (var * exp * exp) list,                iters : (var * exp * exp) list,
# Line 231  Line 266 
266              } = let              } = let
267                val name = Atom.toString strand                val name = Atom.toString strand
268                val nDims = List.length iters                val nDims = List.length iters
269                  val worldTy = CL.T_Ptr(CL.T_Named N.worldTy)
270                fun mapi f xs = let                fun mapi f xs = let
271                      fun mapf (_, []) = []                      fun mapf (_, []) = []
272                        | mapf (i, x::xs) = f(i, x) :: mapf(i+1, xs)                        | mapf (i, x::xs) = f(i, x) :: mapf(i+1, xs)
# Line 239  Line 275 
275                      end                      end
276                val baseInit = mapi (fn (i, (_, e, _)) => (i, CL.I_Exp e)) iters                val baseInit = mapi (fn (i, (_, e, _)) => (i, CL.I_Exp e)) iters
277                val sizeInit = mapi                val sizeInit = mapi
278                      (fn (i, (ToCL.V(ty, _), lo, hi)) =>                      (fn (i, (CL.V(ty, _), lo, hi)) =>
279                          (i, CL.I_Exp(CL.mkBinOp(CL.mkBinOp(hi, CL.#-, lo), CL.#+, CL.mkInt(1, ty))))                          (i, CL.I_Exp(CL.mkBinOp(CL.mkBinOp(hi, CL.#-, lo), CL.#+, CL.E_Int(1, ty))))
280                      ) iters                      ) iters
281                    val numStrandsVar = "numStrandsVar"              (* code to allocate the world and initial strands *)
282                val allocCode = iterPrefix @ [                val wrld = "wrld"
283                  val allocCode = [
284                        CL.mkComment["allocate initial block of strands"],                        CL.mkComment["allocate initial block of strands"],
285                        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)),
286                        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)),
287                        CL.mkDecl(CL.int32,"numDims",SOME(CL.I_Exp(CL.mkInt(IntInf.fromInt nDims, CL.int32))))                        CL.mkDecl(worldTy, wrld,
288                      ]                          SOME(CL.I_Exp(CL.E_Apply(N.allocInitially, [
289                val numStrandsLoopBody =                              CL.mkVar "ProgramName",
290                      CL.mkExpStm(CL.mkAssignOp(CL.mkVar numStrandsVar, CL.*=,CL.mkSubscript(CL.mkVar "size",CL.mkVar "i")))                              CL.mkUnOp(CL.%&, CL.E_Var(N.strandDesc name)),
291                val numStrandsLoop =  CL.mkFor([(CL.intTy, "i", CL.mkInt(0,CL.intTy))],                              CL.E_Bool isArray,
292                      CL.mkBinOp(CL.mkVar "i", CL.#<, CL.mkVar "numDims"),                              CL.E_Int(IntInf.fromInt nDims, CL.int32),
293                      [CL.mkPostOp(CL.mkVar "i", CL.^++)], numStrandsLoopBody)                              CL.E_Var "base",
294                in                              CL.E_Var "size"
295                  numDims := nDims;                            ]))))
                 initially := allocCode @ [numStrandsLoop]  
               end  
   
   
       (***** OUTPUT *****)  
         fun genStrandInit (Strand{name,tyName,state,output,code,...}, nDims) = let  
               val params = [  
                       CL.PARAM([], CL.T_Ptr(CL.uint32), "sizes" ),  
                       CL.PARAM([], CL.intTy, "width"),  
                       CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "strands")  
296                      ]                      ]
297                val body = let              (* create the loop nest for the initially iterations *)
298                      fun loopParams 3 = ["x", "y", "k"]                val indexVar = "ix"
299                        | loopParams 2 = ["x", "y"]                val strandTy = CL.T_Ptr(CL.T_Named(N.strandTy name))
300                        | loopParams 1 = ["x"]                fun mkLoopNest [] = CL.mkBlock(createPrefix @ [
301                        | loopParams _ = raise Fail "genStrandInit: missing size dim"                        CL.mkDecl(strandTy, "sp",
302                      fun mkLoopNest ([], _, nDims) = if nDims = 1                          SOME(CL.I_Exp(
303                            then CL.mkBlock [                            CL.E_Cast(strandTy,
304                                CL.mkCall(RN.strandInit name, [                            CL.E_Apply(N.inState, [CL.E_Var "wrld", CL.E_Var indexVar]))))),
305                                  CL.mkUnOp(CL.%&,CL.mkSubscript(CL.mkVar "strands",CL.mkStr "x")),                        CL.mkCall(N.strandInit name, CL.E_Var "sp" :: args),
306                                                  CL.mkVar "x"])                        CL.mkAssign(CL.E_Var indexVar, CL.mkBinOp(CL.E_Var indexVar, CL.#+, CL.E_Int(1, CL.uint32)))
307                              ]                      ])
308                            else let                  | mkLoopNest ((CL.V(ty, param), lo, hi)::iters) = let
309                              val index = CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")                      val body = mkLoopNest iters
                             in  
                               CL.mkBlock([CL.mkCall(RN.strandInit name, [CL.mkUnOp(CL.%&,CL.mkSubscript(CL.mkVar "strands",index)),  
                               CL.mkVar "x", CL.mkVar"y"])])  
                             end  
                       | mkLoopNest (param::rest,count,nDims) = let  
                           val body = mkLoopNest (rest, count + 1,nDims)  
310                            in                            in
311                              CL.mkFor(                              CL.mkFor(
312                                  [(CL.intTy, param, CL.mkInt(0,CL.intTy))],                          [(ty, param, lo)],
313                                  CL.mkBinOp(CL.mkVar param, CL.#<, CL.mkSubscript(CL.mkVar "sizes",CL.mkInt(count,CL.intTy))),                          CL.mkBinOp(CL.E_Var param, CL.#<=, hi),
314                                  [CL.mkPostOp(CL.mkVar param, CL.^++)],                          [CL.mkPostOp(CL.E_Var param, CL.^++)],
315                                  body)                                  body)
316                            end                            end
317                  val iterCode = [
318                          CL.mkComment["initially"],
319                          CL.mkDecl(CL.uint32, indexVar, SOME(CL.I_Exp(CL.E_Int(0, CL.uint32)))),
320                          mkLoopNest iters
321                        ]
322                  val body = CL.mkBlock(
323                        iterPrefix @
324                        allocCode @
325                        iterCode @
326                        [CL.mkReturn(SOME(CL.E_Var "wrld"))])
327                  val initFn = CL.D_Func([], worldTy, N.initially, [], body)
328                      in                      in
329                        [mkLoopNest ((loopParams nDims),0,nDims)]                  initially := initFn
                     end  
                 in  
                   CL.D_Func(["static"], CL.voidTy, RN.strandInitSetup, params,CL.mkBlock(body))  
330                  end                  end
331    
332          fun genStrandPrint (Strand{name, tyName, state, output, code,...},nDims) = let        (***** OUTPUT *****)
333            fun genStrandPrint (Strand{name, tyName, state, output, code,...}) = let
334              (* the print function *)              (* the print function *)
335                val prFnName = concat[name, "_print"]                val prFnName = concat[name, "_print"]
336                val prFn = let                val prFn = let
337                      val params = [                      val params = [
338                            CL.PARAM([], CL.T_Ptr(CL.T_Named "FILE"), "outS"),                            CL.PARAM([], CL.T_Ptr(CL.T_Named "FILE"), "outS"),
                           CL.PARAM([], CL.T_Ptr(CL.uint32), "sizes" ),  
                           CL.PARAM([], CL.intTy, "width"),  
339                            CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "self")                            CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "self")
340                          ]                          ]
   
341                     val SOME(ty, x) = !output                     val SOME(ty, x) = !output
342                     val outState = if nDims = 1 then                      val outState = CL.mkIndirect(CL.mkVar "self", x)
                           CL.mkSelect(CL.mkSubscript(CL.mkVar "self",CL.mkVar "x"), x)  
                         else if nDims = 2 then  
                                 CL.mkSelect(CL.mkSubscript(CL.mkVar "self",  
                                    CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")), x)  
   
                         else CL.mkSelect(CL.mkVar "self",x)  
   
343                      val prArgs = (case ty                      val prArgs = (case ty
344                             of Ty.IVecTy 1 => [CL.mkStr(!RN.gIntFormat ^ "\n"), outState]                             of Ty.IVecTy 1 => [CL.E_Str(!N.gIntFormat ^ "\n"), outState]
345                              | Ty.IVecTy d => let                              | Ty.IVecTy d => let
346                                  val fmt = CL.mkStr(                                  val fmt = CL.E_Str(
347                                        String.concatWith " " (List.tabulate(d, fn _ => !RN.gIntFormat))                                        String.concatWith " " (List.tabulate(d, fn _ => !N.gIntFormat))
348                                        ^ "\n")                                        ^ "\n")
349                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))                                  val args = List.tabulate (d, fn i => ToC.ivecIndex(outState, d, i))
350                                  in                                  in
351                                    fmt :: args                                    fmt :: args
352                                  end                                  end
353                              | Ty.TensorTy[] => [CL.mkStr "%f\n", outState]                              | Ty.TensorTy[] => [CL.E_Str "%f\n", outState]
354                              | Ty.TensorTy[d] => let                              | Ty.TensorTy[d] => let
355                                  val fmt = CL.mkStr(                                  val fmt = CL.E_Str(
356                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))
357                                        ^ "\n")                                        ^ "\n")
358                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))                                  val args = List.tabulate (d, fn i => ToC.vecIndex(outState, d, i))
359                                  in                                  in
360                                    fmt :: args                                    fmt :: args
361                                  end                                  end
362                              | _ => raise Fail("genStrand: unsupported output type " ^ Ty.toString ty)                              | _ => raise Fail("genStrand: unsupported output type " ^ Ty.toString ty)
363                            (* end case *))                            (* end case *))
   
                           val body = let  
   
                             fun loopParams (3) =  
                                  "x"::"y"::"k"::[]  
                               | loopParams (2) =  
                                  "x"::"y"::[]  
                               | loopParams (1) =  
                                  "x"::[]  
                               | loopParams (_) =  
                                 raise Fail("genStrandPrint: unsupported output type " ^ Ty.toString ty)  
   
                            fun mkLoopNest ([],_) =  
                                                 CL.mkCall("fprintf", CL.mkVar "outS" :: prArgs)  
                                 | mkLoopNest (param::rest,count) = let  
                                         val body = mkLoopNest (rest, count + 1)  
                                    in  
                                                 CL.mkFor(  
                                                         [(CL.intTy, param, CL.mkInt(0,CL.intTy))],  
                                                 CL.mkBinOp(CL.mkVar param, CL.#<, CL.mkSubscript(CL.mkVar "sizes",CL.mkInt(count,CL.intTy))),  
                                                 [CL.mkPostOp(CL.mkVar param, CL.^++)],  
                                                 body)  
                                    end  
364                          in                          in
365                                  [mkLoopNest ((loopParams nDims),0)]                        CL.D_Func(["static"], CL.voidTy, prFnName, params,
366                          end                          CL.mkCall("fprintf", CL.mkVar "outS" :: prArgs))
   
                     in  
                       CL.D_Func(["static"], CL.voidTy, prFnName, params,CL.mkBlock(body))  
367                      end                      end
368                in                in
369                                   prFn                                   prFn
370                end                end
371          fun genStrandTyDef (Strand{tyName, state,...}) =  
372            fun genStrandTyDef (targetTy, Strand{tyName, state,...}) =
373              (* the type declaration for the strand's state struct *)              (* the type declaration for the strand's state struct *)
374                CL.D_StructDef(                CL.D_StructDef(
375                        List.rev (List.map (fn ToCL.V(ty, x) => (ty, x)) (!state)),                  List.rev (List.map (fn x => (targetTy x, #var x)) (!state)),
376                        tyName)                        tyName)
377    
   
378          (* generates the load kernel function *)          (* generates the load kernel function *)
379  (* FIXME: this code might be part of the runtime system *)  
         fun genKernelLoader() =  
                 CL.D_Verbatim ( ["/* Loads the Kernel from a file */",  
                                                 "char * loadKernel (const char * filename) {",  
                                                 "struct stat statbuf;",  
                                                 "FILE *fh;",  
                                                 "char *source;",  
                                                 "fh = fopen(filename, \"r\");",  
                                                 "if (fh == 0)",  
                                                 "   return 0;",  
                                                 "stat(filename, &statbuf);",  
                                                 "source = (char *) malloc(statbuf.st_size + 1);",  
                                                 "fread(source, statbuf.st_size, 1, fh);",  
                                                 "fread(source, statbuf.st_size, 1, fh);",  
                                                 "return source;",  
                                                 "}"])  
380  (* generates the opencl buffers for the image data *)  (* generates the opencl buffers for the image data *)
381          fun getGlobalDataBuffers(globals,count,contextVar,errVar) = let          fun getGlobalDataBuffers(globals,contextVar,errVar) = let
382                  val globalBufferDecl =  CL.mkDecl(clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)                  val globalBufferDecl =  CL.mkDecl(clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)
383                  val globalBuffer = CL.mkAssign(CL.mkVar(concat[RN.globalsVarName,"_cl"]), CL.mkApply("clCreateBuffer",                val globalBuffer = CL.mkAssign(CL.mkVar(concat[RN.globalsVarName,"_cl"]),
384                                                                  [CL.mkVar contextVar,                      CL.mkApply("clCreateBuffer", [
385                            CL.mkVar contextVar,
386                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",
387                                                                  CL.mkApply("sizeof",[CL.mkVar RN.globalsTy]),                                                                  CL.mkApply("sizeof",[CL.mkVar RN.globalsTy]),
388                                                                  CL.mkVar RN.globalsVarName,                                                                  CL.mkVar RN.globalsVarName,
389                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))                          CL.mkUnOp(CL.%&,CL.mkVar errVar)
390                          ]))
391          fun genDataBuffers([],_,_,_) = []                fun genDataBuffers([],_,_) = []
392            | genDataBuffers((var,nDims)::globals,count,contextVar,errVar) = let                  | genDataBuffers((var,nDims)::globals,contextVar,errVar) = let
393          (* FIXME: use CL constructors to  build expressions (not strings) *)          (* FIXME: use CL constructors to  build expressions (not strings) *)
394                    val size = if nDims = 1 then                      val size = if nDims = 1
395                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,                              then CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,
396                                           CL.mkIndirect(CL.mkVar var, "size[0]"))                                           CL.mkIndirect(CL.mkVar var, "size[0]"))
397                                          else if nDims = 2 then                            else if nDims = 2
398                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,                              then CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,
399                                            CL.mkIndirect(CL.mkVar var, concat["size[0]", " * ", var, "->size[1]"]))                                            CL.mkIndirect(CL.mkVar var, concat["size[0]", " * ", var, "->size[1]"]))
400                                          else                              else CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,
                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,  
401                                            CL.mkIndirect(CL.mkVar var,concat["size[0]", " * ", var, "->size[1] * ", var, "->size[2]"]))                                            CL.mkIndirect(CL.mkVar var,concat["size[0]", " * ", var, "->size[1] * ", var, "->size[2]"]))
   
402                   in                   in
403                     CL.mkDecl(clMemoryTy, RN.addBufferSuffix var ,NONE)::                     CL.mkDecl(clMemoryTy, RN.addBufferSuffix var ,NONE)::
404                     CL.mkDecl(clMemoryTy, RN.addBufferSuffixData var ,NONE)::                     CL.mkDecl(clMemoryTy, RN.addBufferSuffixData var ,NONE)::
# Line 435  Line 413 
413                                                                   CL.mkVar "CL_MEM_COPY_HOST_PTR",                                                                   CL.mkVar "CL_MEM_COPY_HOST_PTR",
414                                                                  size,                                                                  size,
415                                                                  CL.mkIndirect(CL.mkVar var,"data"),                                                                  CL.mkIndirect(CL.mkVar var,"data"),
416                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)])):: genDataBuffers(globals,count + 2,contextVar,errVar)                          CL.mkUnOp(CL.%&,CL.mkVar errVar)])):: genDataBuffers(globals,contextVar,errVar)
417                  end                  end
418          in          in
419                  [globalBufferDecl] @ [globalBuffer] @ genDataBuffers(globals,count + 2,contextVar,errVar)                  globalBufferDecl :: globalBuffer :: genDataBuffers(globals,contextVar,errVar)
420          end          end
421    
   
422  (* generates the kernel arguments for the image data *)  (* generates the kernel arguments for the image data *)
423          fun genGlobalArguments(globals,count,kernelVar,errVar) = let          fun genGlobalArguments(globals,count,kernelVar,errVar) = let
424          val globalArgument = CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,CL.mkApply("clSetKernelArg",          val globalArgument = CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,
425                  CL.mkApply("clSetKernelArg",
426                                                                  [CL.mkVar kernelVar,                                                                  [CL.mkVar kernelVar,
427                                                                   CL.mkInt(count,CL.intTy),                   CL.mkPostOp(CL.E_Var count, CL.^++),
428                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
429                                                                   CL.mkUnOp(CL.%&,CL.mkVar(concat[RN.globalsVarName,"_cl"]))])))                                                                   CL.mkUnOp(CL.%&,CL.mkVar(concat[RN.globalsVarName,"_cl"]))])))
   
430          fun genDataArguments([],_,_,_) = []          fun genDataArguments([],_,_,_) = []
431            | genDataArguments((var,nDims)::globals,count,kernelVar,errVar) =            | genDataArguments((var,nDims)::globals,count,kernelVar,errVar) =
432                  CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,
433                  CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=, CL.mkApply("clSetKernelArg",                  CL.mkApply("clSetKernelArg",
434                                                                  [CL.mkVar kernelVar,                                                                  [CL.mkVar kernelVar,
435                                                                   CL.mkInt(count,CL.intTy),                     CL.mkPostOp(CL.E_Var count, CL.^++),
436                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
437                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffix var))])))::                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffix var))])))::
438                  CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,
439                          CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,CL.mkApply("clSetKernelArg",                  CL.mkApply("clSetKernelArg",
440                                                                  [CL.mkVar kernelVar,                                                                  [CL.mkVar kernelVar,
441                                                                   CL.mkInt((count + 1),CL.intTy),                     CL.mkPostOp(CL.E_Var count, CL.^++),
442                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
443                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffixData var))]))):: genDataArguments (globals, count + 2,kernelVar,errVar)                     CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffixData var))]))) ::
444                  genDataArguments (globals,count,kernelVar,errVar)
445          in          in
446              globalArgument :: genDataArguments(globals, count, kernelVar, errVar)
                 [globalArgument] @ genDataArguments(globals,count + 1,kernelVar,errVar)  
   
447          end          end
448    
449          (* generates the main function of host code *)        (* generates the globals buffers and arguments function *)
450          fun genHostMain() = let          fun genGlobalBuffersArgs (imgGlobals) = let
               val setupCall = [CL.mkCall(RN.setupFName,[CL.mkVar RN.globalsVarName])]  
               val globalsDecl = CL.mkDecl(  
                     CL.T_Ptr(CL.T_Named RN.globalsTy),  
                     RN.globalsVarName,  
                     SOME(CL.I_Exp(CL.mkApply("malloc", [CL.mkApply("sizeof",[CL.mkVar RN.globalsTy])]))))  
               val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.mkVar RN.globalsVarName])  
               val returnStm = [CL.mkReturn(SOME(CL.mkInt(0,CL.intTy)))]  
               val params = [  
                      CL.PARAM([],CL.intTy, "argc"),  
                      CL.PARAM([],CL.charArrayPtr,"argv")  
                    ]  
               val body = CL.mkBlock([globalsDecl] @ [initGlobalsCall]  @ setupCall @ returnStm)  
               in  
                 CL.D_Func([],CL.intTy,"main",params,body)  
               end  
   
       (* generates the host-side setup function *)  
         fun genHostSetupFunc (strand as Strand{name,tyName,...}, filename, nDims, initially, imgGlobals) = let  
451              (* Delcare opencl setup objects *)              (* Delcare opencl setup objects *)
               val programVar= "program"  
               val kernelVar = "kernel"  
               val cmdVar = "queue"  
               val inStateVar = "selfin"  
               val outStateVar = "selfout"  
               val stateSizeVar= "state_mem_size"  
               val clInstateVar = "clSelfIn"  
               val clOutStateVar = "clSelfOut"  
               val clGlobals = "clGlobals"  
               val sourcesVar = "sources"  
               val contextVar = "context"  
452                val errVar = "err"                val errVar = "err"
453                val imgDataSizeVar = "image_dataSize"                val imgDataSizeVar = "image_dataSize"
               val globalVar = "global_work_size"  
               val localVar = "local_work_size"  
               val clFNVar = "filename"  
               val numStrandsVar = "numStrandsVar"  
               val headerFNVar = "header"  
               val deviceVar = "device"  
               val platformsVar = "platforms"  
               val numPlatformsVar = "num_platforms"  
               val numDevicesVar = "num_devices"  
               val assertStm = CL.mkCall("assert",[CL.mkBinOp(CL.mkVar errVar, CL.#==, CL.mkVar "CL_SUCCESS")])  
454                val params = [                val params = [
455                        CL.PARAM([],CL.T_Named("cl_device_id"), deviceVar)                        CL.PARAM([],CL.T_Named("cl_context"), "context"),
456                      ]                        CL.PARAM([],CL.T_Named("cl_kernel"), "kernel"),
457                val declarations = [                        CL.PARAM([],CL.T_Named("int"), "argStart")
                     CL.mkDecl(clProgramTy, programVar, NONE),  
                     CL.mkDecl(clKernelTy, kernelVar, NONE),  
                     CL.mkDecl(clCmdQueueTy, cmdVar, NONE),  
                     CL.mkDecl(clContextTy, contextVar, NONE),  
                     CL.mkDecl(CL.intTy, errVar, NONE),  
                     CL.mkDecl(CL.intTy, numStrandsVar, SOME(CL.I_Exp(CL.mkInt(1,CL.intTy)))),  
                     CL.mkDecl(CL.intTy, stateSizeVar, NONE),  
                     CL.mkDecl(CL.intTy, "width", NONE),  
                     CL.mkDecl(CL.intTy, imgDataSizeVar, NONE),  
                     (*CL.mkDecl(clDeviceIdTy, deviceVar, NONE), *)  
                     CL.mkDecl(CL.T_Ptr(CL.T_Named tyName), inStateVar,NONE),  
                     CL.mkDecl(clMemoryTy,clInstateVar,NONE),  
                     CL.mkDecl(clMemoryTy,clOutStateVar,NONE),  
                     CL.mkDecl(CL.T_Ptr(CL.T_Named tyName), outStateVar,NONE),  
                     CL.mkDecl(CL.charPtr, clFNVar,SOME(CL.I_Exp(CL.mkStr filename))),  
 (* FIXME:  use Paths.diderotInclude *)  
                     CL.mkDecl(CL.charPtr, headerFNVar,SOME(CL.I_Exp(CL.mkStr "../src/include/Diderot/cl-types.h"))),  
                     CL.mkDecl(CL.T_Array(CL.charPtr,SOME(2)),sourcesVar,NONE),  
                     CL.mkDecl(CL.T_Array(CL.T_Named "size_t",SOME(nDims)),globalVar,NONE),  
                     CL.mkDecl(CL.T_Array(CL.T_Named "size_t",SOME(nDims)),localVar,NONE),  
                     CL.mkDecl(CL.intTy,numDevicesVar,SOME(CL.I_Exp(CL.mkInt(~1,CL.intTy)))),  
                     CL.mkDecl(CL.T_Array(CL.T_Named "cl_platform_id", SOME(1)), platformsVar, NONE),  
                     CL.mkDecl(CL.intTy,"num_platforms",SOME(CL.I_Exp(CL.mkInt(~1,CL.intTy))))  
458                  ]                  ]
459              (* Setup Global Variables *)                val clGlobalBuffers = getGlobalDataBuffers(!imgGlobals, "context", "err")
460                val globalsDecl = CL.mkDecl(                val clGlobalArguments = genGlobalArguments(!imgGlobals, "argStart", "kernel", "err")
                     CL.T_Ptr(CL.T_Named RN.globalsTy),  
                     RN.globalsVarName,  
                     SOME(CL.I_Exp(CL.mkApply("malloc", [CL.mkApply("sizeof",[CL.mkVar RN.globalsTy])]))))  
               val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.mkVar RN.globalsVarName])  
   
                 (* Retrieve the platforms  
                 val platformStm = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetPlatformIDs",  
                                                   [CL.mkInt(10,CL.intTy),  
                                                    CL.mkVar platformsVar,  
                                                    CL.mkUnOp(CL.%&,CL.mkVar numPlatformsVar)])),  
                                                    assertStm]  
   
                 val devicesStm = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetDeviceIDs",  
                                                   [CL.mkSubscript(CL.mkVar platformsVar,CL.mkInt(0,CL.intTy)),  
                                                    CL.mkVar "CL_DEVICE_TYPE_GPU",  
                                                    CL.mkInt(1,CL.intTy),  
                                                    CL.mkUnOp(CL.%&,CL.mkVar deviceVar),  
                                                    CL.mkUnOp(CL.%&,CL.mkVar numDevicesVar)])),  
                                                    assertStm] *)  
   
                 (* Create Context *)  
                 val contextStm = [CL.mkAssign(CL.mkVar contextVar, CL.mkApply("clCreateContext",  
                                                   [CL.mkInt(0,CL.intTy),  
                                                   CL.mkInt(1,CL.intTy),  
                                                   CL.mkUnOp(CL.%&,CL.mkVar deviceVar),  
                                                   CL.mkVar "NULL",  
                                                   CL.mkVar "NULL",  
                                                   CL.mkUnOp(CL.%&,CL.mkVar errVar)])),  
                                                   assertStm]  
   
                 (* Create Command Queue *)  
                 val commandStm = [CL.mkAssign(CL.mkVar cmdVar, CL.mkApply("clCreateCommandQueue",  
                                                   [CL.mkVar contextVar,  
                                                   CL.mkVar deviceVar,  
                                                   CL.mkInt(0,CL.intTy),  
                                                   CL.mkUnOp(CL.%&,CL.mkVar errVar)])),  
                                                   assertStm]  
   
   
                 (*Create Program/Build/Kernel with Source statement *)  
                 val createProgStm = CL.mkAssign(CL.mkVar programVar, CL.mkApply("clCreateProgramWithSource",  
                                                                                                                 [CL.mkVar contextVar,  
                                                                                                                  CL.mkInt(2,CL.intTy),  
                                                                                                                  CL.mkCast(CL.T_Ptr(CL.T_Named("const char *")),CL.mkUnOp(CL.%&,CL.mkVar sourcesVar)),  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))  
   
                 (* FIXME: Remove after testing purposes, Build Log for OpenCL*)  
                 val buildLog = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clBuildProgram",  
                                                                                                                 [CL.mkVar programVar,  
                                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkVar "NULL"])),  
                                           CL.mkDecl(CL.charPtr, "build", NONE),  
                                           CL.mkDecl(CL.T_Named("size_t"),"ret_val_size",NONE),  
                                            CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetProgramBuildInfo",  
                                                                                                                 [CL.mkVar programVar,  
                                                                                                                 CL.mkVar deviceVar,  
                                                                                                                  CL.mkVar "CL_PROGRAM_BUILD_LOG",  
                                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                                  CL.mkVar "NULL",  
                                                                                                                  CL.mkUnOp(CL.%&,CL.mkVar "ret_val_size")])),  
                                           CL.mkAssign(CL.mkVar "build", CL.mkApply("malloc", [CL.mkVar "ret_val_size"])),  
                                                 CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetProgramBuildInfo",  
                                                                                                                 [CL.mkVar programVar,  
                                                                                                                 CL.mkVar deviceVar,  
                                                                                                                  CL.mkVar "CL_PROGRAM_BUILD_LOG",  
                                                                                                                  CL.mkVar "ret_val_size",  
                                                                                                                  CL.mkVar "build",  
                                                                                                                  CL.mkVar "NULL"])),  
                                                 CL.mkAssign(CL.mkSubscript(CL.mkVar "build",CL.mkVar "ret_val_size"),CL.mkVar ("'\\" ^ "0'")),  
                                                 CL.mkCall("printf",[CL.mkStr ( "Build Log:" ^ "\n" ^ "%s" ^ "\n"), CL.mkVar "build"])]  
   
   
   
   
                 val createKernel = CL.mkAssign(CL.mkVar kernelVar, CL.mkApply("clCreateKernel",  
                                                                                                                 [CL.mkVar programVar,  
                                                                                                                  CL.mkStr RN.kernelFuncName,  
                                                                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))  
   
   
                 val create_build_stms = [createProgStm,assertStm] @ buildLog @ [assertStm,createKernel,assertStm]  
   
   
   
                 (* Create Memory Buffers for Strand States and Globals *)  
                 val strandSize = CL.mkAssign(CL.mkVar stateSizeVar,CL.mkBinOp(CL.mkApply("sizeof",  
                                                                         [CL.mkVar tyName]), CL.#*,CL.mkVar numStrandsVar))  
   
                 val clStrandObjects = [CL.mkAssign(CL.mkVar clInstateVar, CL.mkApply("clCreateBuffer",  
                                                                 [CL.mkVar contextVar,  
                                                                 CL.mkVar "CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR",  
                                                                 CL.mkVar stateSizeVar,  
                                                                 CL.mkVar "NULL",  
                                                                 CL.mkUnOp(CL.%&,CL.mkVar errVar)])),  
                                                          CL.mkAssign(CL.mkVar clOutStateVar, CL.mkApply("clCreateBuffer",  
                                                                 [CL.mkVar contextVar,  
                                                                 CL.mkVar "CL_MEM_READ_WRITE",  
                                                                 CL.mkVar stateSizeVar,  
                                                                 CL.mkVar "NULL",  
                                                                 CL.mkUnOp(CL.%&,CL.mkVar errVar)]))]  
   
   
                 (* Setup up selfOut variable *)  
                 val strandsArrays = [CL.mkAssign(CL.mkVar outStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.mkVar numStrandsVar,  
                                                                         CL.#*, CL.mkApply("sizeof",[CL.mkVar tyName]))])),  
                                                                 CL.mkAssign(CL.mkVar inStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.mkVar numStrandsVar,  
                                                                         CL.#*, CL.mkApply("sizeof",[CL.mkVar tyName]))]))]  
   
   
                 (* Initialize Width Parameter *)  
                 val widthDel = if nDims = 2 then  
                           CL.mkAssign(CL.mkVar "width",CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1, CL.intTy)))  
                    else  
                           CL.mkAssign(CL.mkVar "width",CL.mkInt(0,CL.intTy))  
   
   
                 val strands_init = CL.mkCall(RN.strandInitSetup, [  
                         CL.mkVar "size", CL.mkVar "width", CL.mkVar inStateVar  
                       ])  
   
             val clGlobalBuffers = getGlobalDataBuffers(!imgGlobals,3,contextVar,errVar)  
   
   
                 (* Load the Kernel and Header Files *)  
                 val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(1,CL.intTy)),  
                                                                           CL.mkApply(RN.clLoaderFN, [CL.mkVar clFNVar])),  
            CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(0,CL.intTy)),  
                                                                           CL.mkApply(RN.clLoaderFN, [CL.mkVar headerFNVar]))]  
   
                 (* val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(1,CL.intTy)),  
                                                                           CL.mkApply(RN.clLoaderFN, [CL.mkVar clFNVar]))] *)  
   
   
                 (* Created Enqueue Statements *)  
 (* FIXME: simplify this code by function abstraction *)  
         val enqueueStm = if nDims = 1  
                         then [CL.mkAssign(CL.mkVar errVar,  
                                                           CL.mkApply("clEnqueueNDRangeKernel",  
                                                                                                 [CL.mkVar cmdVar,  
                                                                                                  CL.mkVar kernelVar,  
                                                                                                  CL.mkInt(1,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar globalVar,  
                                                                                                  CL.mkVar localVar,  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]  
                         else if nDims = 2  then  
                          [CL.mkAssign(CL.mkVar errVar,  
                                                         CL.mkApply("clEnqueueNDRangeKernel",  
                                                                                                 [CL.mkVar cmdVar,  
                                                                                                  CL.mkVar kernelVar,  
                                                                                                  CL.mkInt(2,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar globalVar,  
                                                                                                  CL.mkVar localVar,  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]  
                         else  
                           [CL.mkAssign(CL.mkVar errVar,  
                                                         CL.mkApply("clEnqueueNDRangeKernel",  
                                                                                                 [CL.mkVar cmdVar,  
                                                                                                  CL.mkVar kernelVar,  
                                                                                                  CL.mkInt(3,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar globalVar,  
                                                                                                  CL.mkVar localVar,  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]  
   
   
   
                 (* Setup Global and Local variables *)  
   
                 val globalAndlocalStms = if nDims = 1 then  
                         [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),  
                          CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),  
                                                                   CL.mkVar "16")]  
   
   
                 else if nDims = 2 then  
                         [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(1,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),  
                                                                   CL.mkVar "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(1,CL.intTy)),  
                                                                   CL.mkVar "16")]  
   
                 else  
                         [CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(0,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(1,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(2,CL.intTy)),  
                                                                    CL.mkSubscript(CL.mkVar "size", CL.mkInt(2,CL.intTy))),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(0,CL.intTy)),  
                                                                   CL.mkVar "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(1,CL.intTy)),  
                                                                   CL.mkVar "16"),  
                         CL.mkAssign(CL.mkSubscript(CL.mkVar localVar, CL.mkInt(2,CL.intTy)),  
                                                                   CL.mkVar "16")]  
   
   
   
                 (* Setup Kernel arguments *)  
                 val kernelArguments = [CL.mkAssign(CL.mkVar errVar,CL.mkApply("clSetKernelArg",  
                                                                 [CL.mkVar kernelVar,  
                                                                  CL.mkInt(0,CL.intTy),  
                                                                  CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),  
                                                                  CL.mkUnOp(CL.%&,CL.mkVar clInstateVar)])),  
                                                             CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar, CL.|=,CL.mkApply("clSetKernelArg",  
                                                                 [CL.mkVar kernelVar,  
                                                                  CL.mkInt(1,CL.intTy),  
                                                                  CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),  
                                                                  CL.mkUnOp(CL.%&,CL.mkVar clOutStateVar)]))),  
                                                                   CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar, CL.|=,CL.mkApply("clSetKernelArg",  
                                                                 [CL.mkVar kernelVar,  
                                                                  CL.mkInt(2,CL.intTy),  
                                                                  CL.mkApply("sizeof",[CL.mkVar "int"]),  
                                                                  CL.mkUnOp(CL.%&,CL.mkVar "width")])))]  
   
            val clGlobalArguments = genGlobalArguments(!imgGlobals,3,kernelVar,errVar) @ [assertStm]  
   
                 (* Retrieve output *)  
                 val outputStm = CL.mkAssign(CL.mkVar errVar,  
                                                         CL.mkApply("clEnqueueReadBuffer",  
                                                                                                 [CL.mkVar cmdVar,  
                                                                                                  CL.mkVar clOutStateVar,  
                                                                                                  CL.mkVar "CL_TRUE",  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar stateSizeVar,  
                                                                                                  CL.mkVar outStateVar,  
                                                                                                  CL.mkInt(0,CL.intTy),  
                                                                                                  CL.mkVar "NULL",  
                                                                                                  CL.mkVar "NULL"]))  
   
                 (* Free all the objects *)  
                 val freeStms = [CL.mkCall("clReleaseKernel",[CL.mkVar kernelVar]),  
                                                 CL.mkCall("clReleaseProgram",[CL.mkVar programVar ]),  
                                                 CL.mkCall("clReleaseCommandQueue",[CL.mkVar cmdVar]),  
                                                 CL.mkCall("clReleaseContext",[CL.mkVar contextVar]),  
                                                 CL.mkCall("clReleaseMemObject",[CL.mkVar clInstateVar]),  
                                                 CL.mkCall("clReleaseMemObject",[CL.mkVar clOutStateVar])]  
   
   
                 (*Setup Strand Print Function *)  
                 val outputData = [CL.mkDecl(CL.T_Ptr(CL.T_Named("FILE")), "outS", SOME(CL.I_Exp(CL.mkApply("fopen",  
                                                 [CL.mkStr "mip.txt",  
                                                 CL.mkStr "w"])))),  
                                                 CL.mkCall(concat[name, "_print"],  
                                                                         [CL.mkVar "outS",  
                                                                          CL.mkVar "size",  
                                                                          CL.mkVar "width",  
                                                                          CL.mkVar outStateVar])]  
   
   
   
461                  (* Body put all the statments together *)                  (* Body put all the statments together *)
462                  val body =  declarations @ [globalsDecl,initGlobalsCall] (*@ platformStm @ devicesStm *) @ contextStm @ commandStm @ !initially @ [strandSize] @                val body =  clGlobalBuffers @ clGlobalArguments
                                    strandsArrays @ globalAndlocalStms @ [widthDel,strands_init]  @ clStrandObjects @ clGlobalBuffers @ sourceStms  @ create_build_stms  (*@  
                                    kernelArguments @ clGlobalArguments @ enqueueStm @  [outputStm] @ freeStms @ outputData *)  
   
463                  in                  in
464                    CL.D_Func([],CL.voidTy,RN.globalsSetupName,params,CL.mkBlock(body))
                 CL.D_Func([],CL.voidTy,RN.setupFName,params,CL.mkBlock(body))  
   
465                  end                  end
466    
467  (* generate the data and global parameters *)  (* generate the data and global parameters *)
468          fun genKeneralGlobalParams ((name,tyname)::rest) =          fun genKeneralGlobalParams ((name,tyname)::rest) =
469                  CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), concat[RN.globalsVarName]) ::                  CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), concat[RN.globalsVarName]) ::
470                  CL.PARAM([], CL.T_Ptr(CL.T_Named (RN.imageTy tyname)),RN.addBufferSuffix name) ::                  CL.PARAM([], CL.T_Ptr(CL.T_Named (RN.imageTy tyname)),RN.addBufferSuffix name) ::
471                  CL.PARAM([], CL.T_Ptr(CL.voidTy),RN.addBufferSuffixData name) ::                  CL.PARAM([], CL.T_Ptr(CL.voidTy),RN.addBufferSuffixData name) ::
472                  genKeneralGlobalParams(rest)                genKeneralGlobalParams rest
473              | genKeneralGlobalParams [] = []
           | genKeneralGlobalParams ([]) = []  
474    
475          (*generate code for intilizing kernel global data *)          (*generate code for intilizing kernel global data *)
476          fun initKernelGlobals (globals,imgGlobals) = let          fun initKernelGlobals (globals,imgGlobals) = let
477                  fun initGlobalStruct (CL.D_Var(_, _ , name, _)::rest) =  (* FIXME: should use List.map here *)
478                                  CL.mkAssign(CL.mkVar name, CL.mkIndirect(CL.mkVar RN.globalsVarName, name)) ::                fun initGlobalStruct ({hostTy, gpuTy, var}::rest) =
479                                  initGlobalStruct(rest)                      CL.mkAssign(CL.mkVar var, CL.mkIndirect(CL.mkVar RN.globalsVarName, var)) ::
480                    | initGlobalStruct ( _::rest) = initGlobalStruct(rest)                      initGlobalStruct rest
481                    | initGlobalStruct([]) = []                  | initGlobalStruct [] = []
   
482                  fun initGlobalImages((name,tyname)::rest) =                  fun initGlobalImages((name,tyname)::rest) =
483                                  CL.mkAssign(CL.mkVar name, CL.mkVar (RN.addBufferSuffix name)) ::                                  CL.mkAssign(CL.mkVar name, CL.mkVar (RN.addBufferSuffix name)) ::
484                                  CL.mkAssign(CL.mkIndirect(CL.mkVar name,"data"),CL.mkVar (RN.addBufferSuffixData name)) ::                                  CL.mkAssign(CL.mkIndirect(CL.mkVar name,"data"),CL.mkVar (RN.addBufferSuffixData name)) ::
485                                  initGlobalImages(rest)                      initGlobalImages rest
486                    | initGlobalImages([]) = []                    | initGlobalImages [] = []
487                  in                  in
488                    initGlobalStruct(globals) @ initGlobalImages(imgGlobals)                  initGlobalStruct globals @ initGlobalImages(imgGlobals)
489                  end                  end
490    
491          (* generate the main kernel function for the .cl file *)          (* generate the main kernel function for the .cl file *)
492          fun genKernelFun(Strand{name, tyName, state, output, code,...},nDims,globals,imgGlobals) = let          fun genKernelFun (strand, nDims, globals, imgGlobals) = let
493                  val Strand{name, tyName, state, output, code,...} = strand
494                   val fName = RN.kernelFuncName;                   val fName = RN.kernelFuncName;
495                   val inState = "strand_in"                   val inState = "strand_in"
496                   val outState = "strand_out"                   val outState = "strand_out"
# Line 858  Line 500 
500                        CL.PARAM(["__global"], CL.intTy, "width")                        CL.PARAM(["__global"], CL.intTy, "width")
501                      ] @ genKeneralGlobalParams(!imgGlobals)                      ] @ genKeneralGlobalParams(!imgGlobals)
502                    val thread_ids = if nDims = 1                    val thread_ids = if nDims = 1
503                          then [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),                      then [
504                                    CL.mkAssign(CL.mkVar "x",CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(0,CL.intTy)]))]                          CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),
505                          else                          CL.mkAssign(CL.mkVar "x",CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(0,CL.intTy)]))
506                                  [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),                        ]
507                        else [
508                            CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),
509                                   CL.mkDecl(CL.intTy, "y", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),                                   CL.mkDecl(CL.intTy, "y", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy)))),
510                                    CL.mkAssign(CL.mkVar "x",  CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(0,CL.intTy)])),                                    CL.mkAssign(CL.mkVar "x",  CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(0,CL.intTy)])),
511                                    CL.mkAssign(CL.mkVar "y",CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(1,CL.intTy)]))]                          CL.mkAssign(CL.mkVar "y",CL.mkApply(RN.getGlobalThreadId,[CL.mkInt(1,CL.intTy)]))
512                          ]
513                    val strandDecl = [CL.mkDecl(CL.T_Named tyName, inState, NONE),                val strandDecl = [
514                        CL.mkDecl(CL.T_Named tyName, inState, NONE),
515                                                          CL.mkDecl(CL.T_Named tyName, outState,NONE)]                                                          CL.mkDecl(CL.T_Named tyName, outState,NONE)]
516                    val strandObjects  = if nDims = 1                    val strandObjects  = if nDims = 1
517                          then [CL.mkAssign(CL.mkSubscript(CL.mkVar "selfIn",CL.mkStr "x"),                      then [
518                                                                           CL.mkVar inState),                          CL.mkAssign( CL.mkVar inState, CL.mkSubscript(CL.mkVar "selfIn", CL.mkStr "x")),
519                                    CL.mkAssign(CL.mkSubscript(CL.mkVar "selfOut",CL.mkStr "x"),                          CL.mkAssign(CL.mkVar outState,CL.mkSubscript(CL.mkVar "selfOut", CL.mkStr "x"))
520                                                                           CL.mkVar outState)]                        ]
521                          else let                          else let
522                                  val index = CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")                                  val index = CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")
523                                  in                        in [
524                                          [CL.mkAssign(CL.mkSubscript(CL.mkVar "selfIn",index),                          CL.mkAssign(CL.mkVar inState, CL.mkSubscript(CL.mkVar "selfIn",index)),
525                                                                          CL.mkVar inState),                          CL.mkAssign(CL.mkVar outState,CL.mkSubscript(CL.mkVar "selfOut",index))
526                                           CL.mkAssign(CL.mkSubscript(CL.mkVar "selfOut",index),                        ] end
                                                                         CL.mkVar outState)]  
                                 end  
   
   
527                    val status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy))))                    val status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy))))
528                    val local_vars =  thread_ids @ initKernelGlobals(!globals,!imgGlobals)  @ strandDecl @ strandObjects @ [status]                    val local_vars =  thread_ids @ initKernelGlobals(!globals,!imgGlobals)  @ strandDecl @ strandObjects @ [status]
529                    val while_exp = CL.mkBinOp(CL.mkBinOp(CL.mkVar "status",CL.#!=, CL.mkVar RN.kStabilize),CL.#||,CL.mkBinOp(CL.mkVar "status", CL.#!=, CL.mkVar RN.kDie))                val while_exp = CL.mkBinOp(
530                    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)])),                      CL.mkBinOp(CL.mkVar "status",CL.#!=, CL.mkVar RN.kStabilize),
531                                                          CL.mkCall(RN.strandStabilize name,[ CL.mkUnOp(CL.%&,CL.mkVar inState),  CL.mkUnOp(CL.%&,CL.mkVar outState)])]                      CL.#||,
532                        CL.mkBinOp(CL.mkVar "status", CL.#!=, CL.mkVar RN.kDie))
533                    val whileBlock = [CL.mkWhile(while_exp,CL.mkBlock while_body)]                val whileBody = CL.mkBlock [
534                          CL.mkAssign(CL.mkVar "status",
535                            CL.mkApply(RN.strandUpdate name,
536                              [CL.mkUnOp(CL.%&,CL.mkVar inState), CL.mkUnOp(CL.%&,CL.mkVar outState)])),
537                          CL.mkCall(RN.strandStabilize name,
538                            [CL.mkUnOp(CL.%&,CL.mkVar inState), CL.mkUnOp(CL.%&,CL.mkVar outState)])
539                        ]
540                  val whileBlock = [CL.mkWhile(while_exp, whileBody)]
541                    val body = CL.mkBlock(local_vars  @ whileBlock)                    val body = CL.mkBlock(local_vars  @ whileBlock)
542                  in                  in
543                     CL.D_Func(["__kernel"], CL.voidTy, fName, params, body)                     CL.D_Func(["__kernel"], CL.voidTy, fName, params, body)
544                  end                  end
545          (* generate a global structure from the globals *)          (* generate a global structure from the globals *)
546          fun genGlobalStruct(globals) = let          fun genGlobalStruct (targetTy, globals) = let
547                   fun getGlobals(CL.D_Var(_,ty,globalVar,_)::rest) = (ty,globalVar)::getGlobals(rest)                val globs = List.map (fn (x : mirror_var) => (targetTy x, #var x)) globals
                    | getGlobals([]) = []  
                    | getGlobals(_::rest) = getGlobals(rest)  
548                   in                   in
549                          CL.D_StructDef(getGlobals(globals),RN.globalsTy)                  CL.D_StructDef(globs, RN.globalsTy)
550                    end                    end
551            fun genGlobals (declFn, targetTy, globals) = let
552        (* generate the table of strand descriptors *)                fun doVar (x : mirror_var) = declFn (CL.D_Var([], targetTy x, #var x, NONE))
         fun genStrandTable (ppStrm, strands) = let  
               val nStrands = length strands  
               fun genInit (Strand{name, ...}) = CL.I_Exp(CL.mkUnOp(CL.%&, CL.mkVar(RN.strandDesc name)))  
               fun genInits (_, []) = []  
                 | genInits (i, s::ss) = (i, genInit s) :: genInits(i+1, ss)  
               fun ppDecl dcl = PrintAsC.output(ppStrm, dcl)  
553                in                in
554                  ppDecl (CL.D_Var([], CL.int32, RN.numStrands,                  List.app doVar globals
                   SOME(CL.I_Exp(CL.mkInt(IntInf.fromInt nStrands, CL.int32)))));  
                 ppDecl (CL.D_Var([],  
                   CL.T_Array(CL.T_Ptr(CL.T_Named RN.strandDescTy), SOME nStrands),  
                   RN.strands,  
                   SOME(CL.I_Array(genInits (0, strands)))))  
555                end                end
556    
557            fun genSrc (baseName, prog) = let
558          fun genSrc (baseName, Prog{double,globals, topDecls, strands, initially,imgGlobals,numDims,...}) = let                val Prog{double, globals, topDecls, strands, initially, imgGlobals, numDims, ...} = prog
559                val clFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "cl"}                val clFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "cl"}
560                val cFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "c"}                val cFileName = OS.Path.joinBaseExt{base=baseName, ext=SOME "c"}
561                val clOutS = TextIO.openOut clFileName                val clOutS = TextIO.openOut clFileName
562                val cOutS = TextIO.openOut cFileName                val cOutS = TextIO.openOut cFileName
563  (* FIXME: need to use PrintAsC and PrintAsCL *)  (* FIXME: need to use PrintAsC and PrintAsCL *)
564                val clppStrm = PrintAsC.new clOutS                val clppStrm = PrintAsCL.new clOutS
565                val cppStrm = PrintAsC.new cOutS                val cppStrm = PrintAsC.new cOutS
566                fun cppDecl dcl = PrintAsC.output(cppStrm, dcl)                fun cppDecl dcl = PrintAsC.output(cppStrm, dcl)
567                fun clppDecl dcl = PrintAsCL.output(clppStrm, dcl)                fun clppDecl dcl = PrintAsCL.output(clppStrm, dcl)
# Line 940  Line 574 
574                        then "#define DIDEROT_DOUBLE_PRECISION"                        then "#define DIDEROT_DOUBLE_PRECISION"
575                        else "#define DIDEROT_SINGLE_PRECISION",                        else "#define DIDEROT_SINGLE_PRECISION",
576                      "#define DIDEROT_TARGET_CL",                      "#define DIDEROT_TARGET_CL",
577                      "#include \"Diderot/cl-types.h\""                      "#include \"Diderot/cl-diderot.h\""
578                    ]));                    ]));
579                  List.app clppDecl (List.rev (!globals));                  genGlobals (clppDecl, #gpuTy, !globals);
580                  clppDecl (genGlobalStruct (!globals));                  clppDecl (genGlobalStruct (#gpuTy, !globals));
581                  clppDecl (genStrandTyDef strand);                  clppDecl (genStrandTyDef(#gpuTy, strand));
582                  List.app clppDecl (!code);                  List.app clppDecl (!code);
583                  clppDecl (genKernelFun (strand,!numDims,globals,imgGlobals));                  clppDecl (genKernelFun (strand,!numDims,globals,imgGlobals));
584                (* Generate the Host file .c *)  
585                  (* Generate the Host C file *)
586                  cppDecl (CL.D_Verbatim([                  cppDecl (CL.D_Verbatim([
587                      if double                      if double
588                        then "#define DIDEROT_DOUBLE_PRECISION"                        then "#define DIDEROT_DOUBLE_PRECISION"
# Line 955  Line 590 
590                      "#define DIDEROT_TARGET_CL",                      "#define DIDEROT_TARGET_CL",
591                      "#include \"Diderot/diderot.h\""                      "#include \"Diderot/diderot.h\""
592                    ]));                    ]));
593                  List.app cppDecl (List.rev (!globals));                  genGlobals (cppDecl, #hostTy, !globals);
594                  cppDecl (genGlobalStruct (!globals));                  cppDecl (genGlobalStruct (#hostTy, !globals));
595                  cppDecl (genStrandTyDef strand);                  cppDecl (genStrandTyDef (#hostTy, strand));
596                  cppDecl  (!init_code);                  cppDecl  (!init_code);
597                  cppDecl (genStrandInit(strand,!numDims));                  cppDecl (genStrandPrint strand);
                 cppDecl (genStrandPrint(strand,!numDims));  
                 (* cppDecl (genKernelLoader());*)  
598                  List.app cppDecl (List.rev (!topDecls));                  List.app cppDecl (List.rev (!topDecls));
599                  cppDecl (genHostSetupFunc (strand, clFileName, !numDims, initially, imgGlobals));                  cppDecl (genGlobalBuffersArgs (imgGlobals));
600                    cppDecl (!initially);
601                  PrintAsC.close cppStrm;                  PrintAsC.close cppStrm;
602                  PrintAsCL.close clppStrm;                  PrintAsCL.close clppStrm;
603                  TextIO.closeOut cOutS;                  TextIO.closeOut cOutS;

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

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