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 1277, Mon Jun 6 15:20:18 2011 UTC revision 1278, Mon Jun 6 16:27:28 2011 UTC
# Line 28  Line 28 
28        }        }
29    
30      datatype program = Prog of {      datatype program = Prog of {
31            name : string,                  (* stem of source file *)
32          double : bool,                  (* true for double-precision support *)          double : bool,                  (* true for double-precision support *)
33          parallel : bool,                (* true for multithreaded (or multi-GPU) target *)          parallel : bool,                (* true for multithreaded (or multi-GPU) target *)
34          debug : bool,                   (* true for debug support in executable *)          debug : bool,                   (* true for debug support in executable *)
# Line 138  Line 139 
139    (* programs *)    (* programs *)
140      structure Program =      structure Program =
141        struct        struct
142          fun new {double, parallel, debug} = (          fun new {name, double, parallel, debug} = (
143                RN.initTargetSpec double;                RN.initTargetSpec double;
144                Prog{                Prog{
145                      name = name,
146                    double = double, parallel = parallel, debug = debug,                    double = double, parallel = parallel, debug = debug,
147                    globals = ref [],                    globals = ref [],
148                    topDecls = ref [],                    topDecls = ref [],
# Line 152  Line 154 
154                  })                  })
155        (* register the global initialization part of a program *)        (* register the global initialization part of a program *)
156            fun globalIndirects (globals,stms) = let            fun globalIndirects (globals,stms) = let
157                  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) = CL.mkAssign(CL.mkIndirect(CL.mkVar RN.globalsVarName,globalVar),CL.mkVar globalVar)::getGlobals(rest)
158                    | getGlobals([]) = []                    | getGlobals([]) = []
159                    | getGlobals(_::rest) = getGlobals(rest)                    | getGlobals(_::rest) = getGlobals(rest)
160                  in                  in
# Line 208  Line 210 
210                val baseInit = mapi (fn (i, (_, e, _)) => (i, CL.I_Exp e)) iters                val baseInit = mapi (fn (i, (_, e, _)) => (i, CL.I_Exp e)) iters
211                val sizeInit = mapi                val sizeInit = mapi
212                      (fn (i, (ToCL.V(ty, _), lo, hi)) =>                      (fn (i, (ToCL.V(ty, _), lo, hi)) =>
213                          (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))))
214                      ) iters                      ) iters
215                    val numStrandsVar = "numStrandsVar"                    val numStrandsVar = "numStrandsVar"
216                val allocCode = iterPrefix @ [                val allocCode = iterPrefix @ [
217                        CL.mkComment["allocate initial block of strands"],                        CL.mkComment["allocate initial block of strands"],
218                        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)),
219                        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)),
220                        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))))
221                      ]                      ]
222                val numStrandsLoopBody = CL.mkExpStm(CL.mkAssignOp(CL.E_Var numStrandsVar, CL.*=,CL.mkSubscript(CL.E_Var "size",CL.E_Var "i")))                val numStrandsLoopBody = CL.mkExpStm(CL.mkAssignOp(CL.mkVar numStrandsVar, CL.*=,CL.mkSubscript(CL.mkVar "size",CL.mkVar "i")))
223                val numStrandsLoop =  CL.mkFor([(CL.intTy, "i", CL.E_Int(0,CL.intTy))],                val numStrandsLoop =  CL.mkFor([(CL.intTy, "i", CL.mkInt(0,CL.intTy))],
224                      CL.mkBinOp(CL.E_Var "i", CL.#<, CL.E_Var "numDims"),                      CL.mkBinOp(CL.mkVar "i", CL.#<, CL.mkVar "numDims"),
225                      [CL.mkPostOp(CL.E_Var "i", CL.^++)], numStrandsLoopBody)                      [CL.mkPostOp(CL.mkVar "i", CL.^++)], numStrandsLoopBody)
226                in                in
227                  numDims := nDims;                  numDims := nDims;
228                  initially := allocCode @ [numStrandsLoop]                  initially := allocCode @ [numStrandsLoop]
# Line 234  Line 236 
236                            CL.PARAM([], CL.intTy, "width"),                            CL.PARAM([], CL.intTy, "width"),
237                            CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "strands")                            CL.PARAM([], CL.T_Ptr(CL.T_Named tyName), "strands")
238                          ]                          ]
   
239                  val body = let                  val body = let
240                              fun loopParams (3) =                      fun loopParams 3 = ["x", "y", "k"]
241                                   "x"::"y"::"k"::[]                        | loopParams 2 = ["x", "y"]
242                                | loopParams (2) =                        | loopParams 1 = ["x"]
243                                   "x"::"y"::[]                        | loopParams _ = raise Fail "genStrandInit: missing size dim"
244                                | loopParams (1) =                      fun mkLoopNest ([], _, nDims) = if nDims = 1
245                                   "x"::[]                            then CL.mkBlock [
246                                | loopParams (_) =                                CL.mkCall(RN.strandInit name, [
247                                  raise Fail("genStrandInit: missing size dim")                                  CL.mkUnOp(CL.%&,CL.mkSubscript(CL.mkVar "strands",CL.mkStr "x")),
248                             fun mkLoopNest ([],_,nDims) =  if nDims = 1 then                                                  CL.mkVar "x"])
249                                ]
                                         CL.mkBlock ([CL.mkCall(RN.strandInit name, [CL.E_UnOp(CL.%&,CL.mkSubscript(CL.E_Var "strands",CL.E_Str "x")),  
                                                 CL.E_Var "x"])])  
250                                          else let                                          else let
251                                                  val index = CL.mkBinOp(CL.mkBinOp(CL.E_Var "x",CL.#*,CL.E_Var "width"),CL.#+,CL.E_Var "y")                                                  val index = CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")
252                                          in                                          in
253                                                  CL.mkBlock([CL.mkCall(RN.strandInit name, [CL.E_UnOp(CL.%&,CL.mkSubscript(CL.E_Var "strands",index)),                                                  CL.mkBlock([CL.mkCall(RN.strandInit name, [CL.mkUnOp(CL.%&,CL.mkSubscript(CL.mkVar "strands",index)),
254                                                  CL.E_Var "x", CL.E_Var"y"])])                                                  CL.mkVar "x", CL.mkVar"y"])])
255                                          end                                          end
256    
257                                  | mkLoopNest (param::rest,count,nDims) = let                                  | mkLoopNest (param::rest,count,nDims) = let
258                                          val body = mkLoopNest (rest, count + 1,nDims)                                          val body = mkLoopNest (rest, count + 1,nDims)
259                                     in                                     in
260                                          CL.mkFor(                                          CL.mkFor(
261                                                          [(CL.intTy, param, CL.E_Int(0,CL.intTy))],                                                          [(CL.intTy, param, CL.mkInt(0,CL.intTy))],
262                                                  CL.mkBinOp(CL.E_Var param, CL.#<, CL.mkSubscript(CL.E_Var "sizes",CL.E_Int(count,CL.intTy))),                                                  CL.mkBinOp(CL.mkVar param, CL.#<, CL.mkSubscript(CL.mkVar "sizes",CL.mkInt(count,CL.intTy))),
263                                                  [CL.mkPostOp(CL.E_Var param, CL.^++)],                                                  [CL.mkPostOp(CL.mkVar param, CL.^++)],
264                                                  body)                                                  body)
265                                     end                                     end
266                          in                          in
# Line 284  Line 283 
283    
284                     val SOME(ty, x) = !output                     val SOME(ty, x) = !output
285                     val outState = if nDims = 1 then                     val outState = if nDims = 1 then
286                            CL.mkSelect(CL.mkSubscript(CL.mkVar "self",CL.E_Var "x"), x)                            CL.mkSelect(CL.mkSubscript(CL.mkVar "self",CL.mkVar "x"), x)
287                          else if nDims = 2 then                          else if nDims = 2 then
288                                  CL.mkSelect(CL.mkSubscript(CL.mkVar "self",                                  CL.mkSelect(CL.mkSubscript(CL.mkVar "self",
289                                     CL.mkBinOp(CL.mkBinOp(CL.E_Var "x",CL.#*,CL.E_Var "width"),CL.#+,CL.E_Var "y")), x)                                     CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")), x)
290    
291                          else CL.mkSelect(CL.mkVar "self",x)                          else CL.mkSelect(CL.mkVar "self",x)
292    
293                      val prArgs = (case ty                      val prArgs = (case ty
294                             of Ty.IVecTy 1 => [CL.E_Str(!RN.gIntFormat ^ "\n"), outState]                             of Ty.IVecTy 1 => [CL.mkStr(!RN.gIntFormat ^ "\n"), outState]
295                              | Ty.IVecTy d => let                              | Ty.IVecTy d => let
296                                  val fmt = CL.E_Str(                                  val fmt = CL.mkStr(
297                                        String.concatWith " " (List.tabulate(d, fn _ => !RN.gIntFormat))                                        String.concatWith " " (List.tabulate(d, fn _ => !RN.gIntFormat))
298                                        ^ "\n")                                        ^ "\n")
299                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))
300                                  in                                  in
301                                    fmt :: args                                    fmt :: args
302                                  end                                  end
303                              | Ty.TensorTy[] => [CL.E_Str "%f\n", outState]                              | Ty.TensorTy[] => [CL.mkStr "%f\n", outState]
304                              | Ty.TensorTy[d] => let                              | Ty.TensorTy[d] => let
305                                  val fmt = CL.E_Str(                                  val fmt = CL.mkStr(
306                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))                                        String.concatWith " " (List.tabulate(d, fn _ => "%f"))
307                                        ^ "\n")                                        ^ "\n")
308                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))                                  val args = List.tabulate (d, fn i => ToCL.vecIndex(outState, i))
# Line 330  Line 329 
329                                          val body = mkLoopNest (rest, count + 1)                                          val body = mkLoopNest (rest, count + 1)
330                                     in                                     in
331                                                  CL.mkFor(                                                  CL.mkFor(
332                                                          [(CL.intTy, param, CL.E_Int(0,CL.intTy))],                                                          [(CL.intTy, param, CL.mkInt(0,CL.intTy))],
333                                                  CL.mkBinOp(CL.E_Var param, CL.#<, CL.mkSubscript(CL.E_Var "sizes",CL.E_Int(count,CL.intTy))),                                                  CL.mkBinOp(CL.mkVar param, CL.#<, CL.mkSubscript(CL.mkVar "sizes",CL.mkInt(count,CL.intTy))),
334                                                  [CL.mkPostOp(CL.E_Var param, CL.^++)],                                                  [CL.mkPostOp(CL.mkVar param, CL.^++)],
335                                                  body)                                                  body)
336                                     end                                     end
337                          in                          in
# Line 372  Line 371 
371  (* generates the opencl buffers for the image data *)  (* generates the opencl buffers for the image data *)
372          fun getGlobalDataBuffers(globals,count,contextVar,errVar) = let          fun getGlobalDataBuffers(globals,count,contextVar,errVar) = let
373                  val globalBufferDecl =  CL.mkDecl(CL.clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)                  val globalBufferDecl =  CL.mkDecl(CL.clMemoryTy,concat[RN.globalsVarName,"_cl"],NONE)
374                  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",
375                                                                  [CL.E_Var contextVar,                                                                  [CL.mkVar contextVar,
376                                                                  CL.E_Var "CL_MEM_COPY_HOST_PTR",                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",
377                                                                  CL.mkApply("sizeof",[CL.E_Var RN.globalsTy]),                                                                  CL.mkApply("sizeof",[CL.mkVar RN.globalsTy]),
378                                                                  CL.E_Var RN.globalsVarName,                                                                  CL.mkVar RN.globalsVarName,
379                                                                  CL.E_UnOp(CL.%&,CL.E_Var errVar)]))                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))
380    
381          fun genDataBuffers([],_,_,_) = []          fun genDataBuffers([],_,_,_) = []
382            | genDataBuffers((var,nDims)::globals,count,contextVar,errVar) = let            | genDataBuffers((var,nDims)::globals,count,contextVar,errVar) = let
383          (* FIXME: use CL constructors to  build expressions (not strings) *)          (* FIXME: use CL constructors to  build expressions (not strings) *)
384                    val size = if nDims = 1 then                    val size = if nDims = 1 then
385                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.E_Var "float"]), CL.#*,                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,
386                                           CL.mkIndirect(CL.E_Var var, "size[0]"))                                           CL.mkIndirect(CL.mkVar var, "size[0]"))
387                                          else if nDims = 2 then                                          else if nDims = 2 then
388                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.E_Var "float"]), CL.#*,                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,
389                                            CL.mkIndirect(CL.E_Var var, concat["size[0]", " * ", var, "->size[1]"]))                                            CL.mkIndirect(CL.mkVar var, concat["size[0]", " * ", var, "->size[1]"]))
390                                          else                                          else
391                                           CL.mkBinOp(CL.mkApply("sizeof",[CL.E_Var "float"]), CL.#*,                                           CL.mkBinOp(CL.mkApply("sizeof",[CL.mkVar "float"]), CL.#*,
392                                            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]"]))
393    
394                   in                   in
395                     CL.mkDecl(CL.clMemoryTy,RN.addBufferSuffix var ,NONE)::                     CL.mkDecl(CL.clMemoryTy,RN.addBufferSuffix var ,NONE)::
396                     CL.mkDecl(CL.clMemoryTy,RN.addBufferSuffixData var ,NONE)::                     CL.mkDecl(CL.clMemoryTy,RN.addBufferSuffixData var ,NONE)::
397                     CL.mkAssign(CL.E_Var(RN.addBufferSuffix var), CL.mkApply("clCreateBuffer",                     CL.mkAssign(CL.mkVar(RN.addBufferSuffix var), CL.mkApply("clCreateBuffer",
398                                                                  [CL.E_Var contextVar,                                                                  [CL.mkVar contextVar,
399                                                                  CL.E_Var "CL_MEM_COPY_HOST_PTR",                                                                  CL.mkVar "CL_MEM_COPY_HOST_PTR",
400                                                                  CL.mkApply("sizeof",[CL.E_Var (RN.imageTy nDims)]),                                                                  CL.mkApply("sizeof",[CL.mkVar (RN.imageTy nDims)]),
401                                                                  CL.E_Var var,                                                                  CL.mkVar var,
402                                                                  CL.E_UnOp(CL.%&,CL.E_Var errVar)])) ::                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)])) ::
403                          CL.mkAssign(CL.E_Var(RN.addBufferSuffixData var), CL.mkApply("clCreateBuffer",                          CL.mkAssign(CL.mkVar(RN.addBufferSuffixData var), CL.mkApply("clCreateBuffer",
404                                                                  [CL.E_Var contextVar,                                                                  [CL.mkVar contextVar,
405                                                                   CL.E_Var "CL_MEM_COPY_HOST_PTR",                                                                   CL.mkVar "CL_MEM_COPY_HOST_PTR",
406                                                                  size,                                                                  size,
407                                                                  CL.mkIndirect(CL.E_Var var,"data"),                                                                  CL.mkIndirect(CL.mkVar var,"data"),
408                                                                  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)
409                  end                  end
410          in          in
411                  [globalBufferDecl] @ [globalBuffer] @ genDataBuffers(globals,count + 2,contextVar,errVar)                  [globalBufferDecl] @ [globalBuffer] @ genDataBuffers(globals,count + 2,contextVar,errVar)
# Line 415  Line 414 
414    
415  (* generates the kernel arguments for the image data *)  (* generates the kernel arguments for the image data *)
416          fun genGlobalArguments(globals,count,kernelVar,errVar) = let          fun genGlobalArguments(globals,count,kernelVar,errVar) = let
417          val globalArgument = CL.mkExpStm(CL.mkAssignOp(CL.E_Var errVar,CL.|=,CL.mkApply("clSetKernelArg",          val globalArgument = CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,CL.mkApply("clSetKernelArg",
418                                                                  [CL.E_Var kernelVar,                                                                  [CL.mkVar kernelVar,
419                                                                   CL.E_Int(count,CL.intTy),                                                                   CL.mkInt(count,CL.intTy),
420                                                                   CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
421                                                                   CL.E_UnOp(CL.%&,CL.E_Var(concat[RN.globalsVarName,"_cl"]))])))                                                                   CL.mkUnOp(CL.%&,CL.mkVar(concat[RN.globalsVarName,"_cl"]))])))
422    
423          fun genDataArguments([],_,_,_) = []          fun genDataArguments([],_,_,_) = []
424            | genDataArguments((var,nDims)::globals,count,kernelVar,errVar) =            | genDataArguments((var,nDims)::globals,count,kernelVar,errVar) =
425    
426                  CL.mkExpStm(CL.mkAssignOp(CL.E_Var errVar,CL.|=, CL.mkApply("clSetKernelArg",                  CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=, CL.mkApply("clSetKernelArg",
427                                                                  [CL.E_Var kernelVar,                                                                  [CL.mkVar kernelVar,
428                                                                   CL.E_Int(count,CL.intTy),                                                                   CL.mkInt(count,CL.intTy),
429                                                                   CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
430                                                                   CL.E_UnOp(CL.%&,CL.E_Var(RN.addBufferSuffix var))])))::                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffix var))])))::
431    
432                          CL.mkExpStm(CL.mkAssignOp(CL.E_Var errVar,CL.|=,CL.mkApply("clSetKernelArg",                          CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar,CL.|=,CL.mkApply("clSetKernelArg",
433                                                                  [CL.E_Var kernelVar,                                                                  [CL.mkVar kernelVar,
434                                                                   CL.E_Int((count + 1),CL.intTy),                                                                   CL.mkInt((count + 1),CL.intTy),
435                                                                   CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
436                                                                   CL.E_UnOp(CL.%&,CL.E_Var(RN.addBufferSuffixData var))]))):: genDataArguments (globals, count + 2,kernelVar,errVar)                                                                   CL.mkUnOp(CL.%&,CL.mkVar(RN.addBufferSuffixData var))]))):: genDataArguments (globals, count + 2,kernelVar,errVar)
437    
438          in          in
439    
# Line 444  Line 443 
443    
444          (* generates the main function of host code *)          (* generates the main function of host code *)
445          fun genHostMain() = let          fun genHostMain() = let
446                val setupCall = [CL.mkCall(RN.setupFName,[CL.E_Var RN.globalsVarName])]                val setupCall = [CL.mkCall(RN.setupFName,[CL.mkVar RN.globalsVarName])]
447                val globalsDecl = CL.mkDecl(                val globalsDecl = CL.mkDecl(
448                      CL.T_Ptr(CL.T_Named RN.globalsTy),                      CL.T_Ptr(CL.T_Named RN.globalsTy),
449                      RN.globalsVarName,                      RN.globalsVarName,
450                      SOME(CL.I_Exp(CL.mkApply("malloc", [CL.mkApply("sizeof",[CL.E_Var RN.globalsTy])]))))                      SOME(CL.I_Exp(CL.mkApply("malloc", [CL.mkApply("sizeof",[CL.mkVar RN.globalsTy])]))))
451                val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.E_Var RN.globalsVarName])                val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.mkVar RN.globalsVarName])
452                val returnStm = [CL.mkReturn(SOME(CL.E_Int(0,CL.intTy)))]                val returnStm = [CL.mkReturn(SOME(CL.mkInt(0,CL.intTy)))]
453                val params = [                val params = [
454                       CL.PARAM([],CL.intTy, "argc"),                       CL.PARAM([],CL.intTy, "argc"),
455                       CL.PARAM([],CL.charArrayPtr,"argv")                       CL.PARAM([],CL.charArrayPtr,"argv")
# Line 485  Line 484 
484                val platformsVar = "platforms"                val platformsVar = "platforms"
485                val numPlatformsVar = "num_platforms"                val numPlatformsVar = "num_platforms"
486                val numDevicesVar = "num_devices"                val numDevicesVar = "num_devices"
487                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")])
488                val params = [                val params = [
489                        CL.PARAM([],CL.T_Named("cl_device_id"), deviceVar)                        CL.PARAM([],CL.T_Named("cl_device_id"), deviceVar)
490                      ]                      ]
# Line 495  Line 494 
494                      CL.mkDecl(CL.clCmdQueueTy, cmdVar, NONE),                      CL.mkDecl(CL.clCmdQueueTy, cmdVar, NONE),
495                      CL.mkDecl(CL.clContextTy, contextVar, NONE),                      CL.mkDecl(CL.clContextTy, contextVar, NONE),
496                      CL.mkDecl(CL.intTy, errVar, NONE),                      CL.mkDecl(CL.intTy, errVar, NONE),
497                      CL.mkDecl(CL.intTy, numStrandsVar, SOME(CL.I_Exp(CL.E_Int(1,CL.intTy)))),                      CL.mkDecl(CL.intTy, numStrandsVar, SOME(CL.I_Exp(CL.mkInt(1,CL.intTy)))),
498                      CL.mkDecl(CL.intTy, stateSizeVar, NONE),                      CL.mkDecl(CL.intTy, stateSizeVar, NONE),
499                      CL.mkDecl(CL.intTy, "width", NONE),                      CL.mkDecl(CL.intTy, "width", NONE),
500                      CL.mkDecl(CL.intTy, imgDataSizeVar, NONE),                      CL.mkDecl(CL.intTy, imgDataSizeVar, NONE),
# Line 504  Line 503 
503                      CL.mkDecl(CL.clMemoryTy,clInstateVar,NONE),                      CL.mkDecl(CL.clMemoryTy,clInstateVar,NONE),
504                      CL.mkDecl(CL.clMemoryTy,clOutStateVar,NONE),                      CL.mkDecl(CL.clMemoryTy,clOutStateVar,NONE),
505                      CL.mkDecl(CL.T_Ptr(CL.T_Named tyName), outStateVar,NONE),                      CL.mkDecl(CL.T_Ptr(CL.T_Named tyName), outStateVar,NONE),
506                      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))),
507  (* FIXME:  use Paths.diderotInclude *)  (* FIXME:  use Paths.diderotInclude *)
508                      CL.mkDecl(CL.charPtr, headerFNVar,SOME(CL.I_Exp(CL.E_Str "../src/include/Diderot/cl-types.h"))),                      CL.mkDecl(CL.charPtr, headerFNVar,SOME(CL.I_Exp(CL.mkStr "../src/include/Diderot/cl-types.h"))),
509                      CL.mkDecl(CL.T_Array(CL.charPtr,SOME(2)),sourcesVar,NONE),                      CL.mkDecl(CL.T_Array(CL.charPtr,SOME(2)),sourcesVar,NONE),
510                      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),
511                      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),
512                      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)))),
513                      CL.mkDecl(CL.T_Array(CL.T_Named "cl_platform_id", SOME(1)), platformsVar, NONE),                      CL.mkDecl(CL.T_Array(CL.T_Named "cl_platform_id", SOME(1)), platformsVar, NONE),
514                      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))))
515                  ]                  ]
516              (* Setup Global Variables *)              (* Setup Global Variables *)
517                val globalsDecl = CL.mkDecl(                val globalsDecl = CL.mkDecl(
518                      CL.T_Ptr(CL.T_Named RN.globalsTy),                      CL.T_Ptr(CL.T_Named RN.globalsTy),
519                      RN.globalsVarName,                      RN.globalsVarName,
520                      SOME(CL.I_Exp(CL.mkApply("malloc", [CL.mkApply("sizeof",[CL.E_Var RN.globalsTy])]))))                      SOME(CL.I_Exp(CL.mkApply("malloc", [CL.mkApply("sizeof",[CL.mkVar RN.globalsTy])]))))
521                val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.E_Var RN.globalsVarName])                val initGlobalsCall = CL.mkCall(RN.initGlobals,[CL.mkVar RN.globalsVarName])
522    
523                  (* Retrieve the platforms                  (* Retrieve the platforms
524                  val platformStm = [CL.mkAssign(CL.E_Var errVar, CL.mkApply("clGetPlatformIDs",                  val platformStm = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetPlatformIDs",
525                                                    [CL.E_Int(10,CL.intTy),                                                    [CL.mkInt(10,CL.intTy),
526                                                     CL.E_Var platformsVar,                                                     CL.mkVar platformsVar,
527                                                     CL.E_UnOp(CL.%&,CL.E_Var numPlatformsVar)])),                                                     CL.mkUnOp(CL.%&,CL.mkVar numPlatformsVar)])),
528                                                     assertStm]                                                     assertStm]
529    
530                  val devicesStm = [CL.mkAssign(CL.E_Var errVar, CL.mkApply("clGetDeviceIDs",                  val devicesStm = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetDeviceIDs",
531                                                    [CL.mkSubscript(CL.E_Var platformsVar,CL.E_Int(0,CL.intTy)),                                                    [CL.mkSubscript(CL.mkVar platformsVar,CL.mkInt(0,CL.intTy)),
532                                                     CL.E_Var "CL_DEVICE_TYPE_GPU",                                                     CL.mkVar "CL_DEVICE_TYPE_GPU",
533                                                     CL.E_Int(1,CL.intTy),                                                     CL.mkInt(1,CL.intTy),
534                                                     CL.E_UnOp(CL.%&,CL.E_Var deviceVar),                                                     CL.mkUnOp(CL.%&,CL.mkVar deviceVar),
535                                                     CL.E_UnOp(CL.%&,CL.E_Var numDevicesVar)])),                                                     CL.mkUnOp(CL.%&,CL.mkVar numDevicesVar)])),
536                                                     assertStm] *)                                                     assertStm] *)
537    
538                  (* Create Context *)                  (* Create Context *)
539                  val contextStm = [CL.mkAssign(CL.E_Var contextVar, CL.mkApply("clCreateContext",                  val contextStm = [CL.mkAssign(CL.mkVar contextVar, CL.mkApply("clCreateContext",
540                                                    [CL.E_Int(0,CL.intTy),                                                    [CL.mkInt(0,CL.intTy),
541                                                    CL.E_Int(1,CL.intTy),                                                    CL.mkInt(1,CL.intTy),
542                                                    CL.E_UnOp(CL.%&,CL.E_Var deviceVar),                                                    CL.mkUnOp(CL.%&,CL.mkVar deviceVar),
543                                                    CL.E_Var "NULL",                                                    CL.mkVar "NULL",
544                                                    CL.E_Var "NULL",                                                    CL.mkVar "NULL",
545                                                    CL.E_UnOp(CL.%&,CL.E_Var errVar)])),                                                    CL.mkUnOp(CL.%&,CL.mkVar errVar)])),
546                                                    assertStm]                                                    assertStm]
547    
548                  (* Create Command Queue *)                  (* Create Command Queue *)
549                  val commandStm = [CL.mkAssign(CL.E_Var cmdVar, CL.mkApply("clCreateCommandQueue",                  val commandStm = [CL.mkAssign(CL.mkVar cmdVar, CL.mkApply("clCreateCommandQueue",
550                                                    [CL.E_Var contextVar,                                                    [CL.mkVar contextVar,
551                                                    CL.E_Var deviceVar,                                                    CL.mkVar deviceVar,
552                                                    CL.E_Int(0,CL.intTy),                                                    CL.mkInt(0,CL.intTy),
553                                                    CL.E_UnOp(CL.%&,CL.E_Var errVar)])),                                                    CL.mkUnOp(CL.%&,CL.mkVar errVar)])),
554                                                    assertStm]                                                    assertStm]
555    
556    
557                  (*Create Program/Build/Kernel with Source statement *)                  (*Create Program/Build/Kernel with Source statement *)
558                  val createProgStm = CL.mkAssign(CL.E_Var programVar, CL.mkApply("clCreateProgramWithSource",                  val createProgStm = CL.mkAssign(CL.mkVar programVar, CL.mkApply("clCreateProgramWithSource",
559                                                                                                                  [CL.E_Var contextVar,                                                                                                                  [CL.mkVar contextVar,
560                                                                                                                   CL.E_Int(2,CL.intTy),                                                                                                                   CL.mkInt(2,CL.intTy),
561                                                                                                                   CL.E_Cast(CL.T_Ptr(CL.T_Named("const char *")),CL.E_UnOp(CL.%&,CL.E_Var sourcesVar)),                                                                                                                   CL.mkCast(CL.T_Ptr(CL.T_Named("const char *")),CL.mkUnOp(CL.%&,CL.mkVar sourcesVar)),
562                                                                                                                   CL.E_Var "NULL",                                                                                                                   CL.mkVar "NULL",
563                                                                                                                   CL.E_UnOp(CL.%&,CL.E_Var errVar)]))                                                                                                                   CL.mkUnOp(CL.%&,CL.mkVar errVar)]))
564    
565                  (* FIXME: Remove after testing purposes, Build Log for OpenCL*)                  (* FIXME: Remove after testing purposes, Build Log for OpenCL*)
566                  val buildLog = [CL.mkAssign(CL.E_Var errVar, CL.mkApply("clBuildProgram",                  val buildLog = [CL.mkAssign(CL.mkVar errVar, CL.mkApply("clBuildProgram",
567                                                                                                                  [CL.E_Var programVar,                                                                                                                  [CL.mkVar programVar,
568                                                                                                                   CL.E_Int(0,CL.intTy),                                                                                                                   CL.mkInt(0,CL.intTy),
569                                                                                                                   CL.E_Var "NULL",                                                                                                                   CL.mkVar "NULL",
570                                                                                                                   CL.E_Var "NULL",                                                                                                                   CL.mkVar "NULL",
571                                                                                                                   CL.E_Var "NULL",                                                                                                                   CL.mkVar "NULL",
572                                                                                                                   CL.E_Var "NULL"])),                                                                                                                   CL.mkVar "NULL"])),
573                                            CL.mkDecl(CL.charPtr, "build", NONE),                                            CL.mkDecl(CL.charPtr, "build", NONE),
574                                            CL.mkDecl(CL.T_Named("size_t"),"ret_val_size",NONE),                                            CL.mkDecl(CL.T_Named("size_t"),"ret_val_size",NONE),
575                                             CL.mkAssign(CL.E_Var errVar, CL.mkApply("clGetProgramBuildInfo",                                             CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetProgramBuildInfo",
576                                                                                                                  [CL.E_Var programVar,                                                                                                                  [CL.mkVar programVar,
577                                                                                                                  CL.E_Var deviceVar,                                                                                                                  CL.mkVar deviceVar,
578                                                                                                                   CL.E_Var "CL_PROGRAM_BUILD_LOG",                                                                                                                   CL.mkVar "CL_PROGRAM_BUILD_LOG",
579                                                                                                                   CL.E_Int(0,CL.intTy),                                                                                                                   CL.mkInt(0,CL.intTy),
580                                                                                                                   CL.E_Var "NULL",                                                                                                                   CL.mkVar "NULL",
581                                                                                                                   CL.E_UnOp(CL.%&,CL.E_Var "ret_val_size")])),                                                                                                                   CL.mkUnOp(CL.%&,CL.mkVar "ret_val_size")])),
582                                            CL.mkAssign(CL.E_Var "build", CL.mkApply("malloc", [CL.E_Var "ret_val_size"])),                                            CL.mkAssign(CL.mkVar "build", CL.mkApply("malloc", [CL.mkVar "ret_val_size"])),
583                                                  CL.mkAssign(CL.E_Var errVar, CL.mkApply("clGetProgramBuildInfo",                                                  CL.mkAssign(CL.mkVar errVar, CL.mkApply("clGetProgramBuildInfo",
584                                                                                                                  [CL.E_Var programVar,                                                                                                                  [CL.mkVar programVar,
585                                                                                                                  CL.E_Var deviceVar,                                                                                                                  CL.mkVar deviceVar,
586                                                                                                                   CL.E_Var "CL_PROGRAM_BUILD_LOG",                                                                                                                   CL.mkVar "CL_PROGRAM_BUILD_LOG",
587                                                                                                                   CL.E_Var "ret_val_size",                                                                                                                   CL.mkVar "ret_val_size",
588                                                                                                                   CL.E_Var "build",                                                                                                                   CL.mkVar "build",
589                                                                                                                   CL.E_Var "NULL"])),                                                                                                                   CL.mkVar "NULL"])),
590                                                  CL.mkAssign(CL.mkSubscript(CL.E_Var "build",CL.E_Var "ret_val_size"),CL.E_Var ("'\\" ^ "0'")),                                                  CL.mkAssign(CL.mkSubscript(CL.mkVar "build",CL.mkVar "ret_val_size"),CL.mkVar ("'\\" ^ "0'")),
591                                                  CL.mkCall("printf",[CL.E_Str ( "Build Log:" ^ "\n" ^ "%s" ^ "\n"), CL.E_Var "build"])]                                                  CL.mkCall("printf",[CL.mkStr ( "Build Log:" ^ "\n" ^ "%s" ^ "\n"), CL.mkVar "build"])]
592    
593    
594    
595    
596                  val createKernel = CL.mkAssign(CL.E_Var kernelVar, CL.mkApply("clCreateKernel",                  val createKernel = CL.mkAssign(CL.mkVar kernelVar, CL.mkApply("clCreateKernel",
597                                                                                                                  [CL.E_Var programVar,                                                                                                                  [CL.mkVar programVar,
598                                                                                                                   CL.E_Str RN.kernelFuncName,                                                                                                                   CL.mkStr RN.kernelFuncName,
599                                                                                                                   CL.E_UnOp(CL.%&,CL.E_Var errVar)]))                                                                                                                   CL.mkUnOp(CL.%&,CL.mkVar errVar)]))
600    
601    
602                  val create_build_stms = [createProgStm,assertStm] @ buildLog @ [assertStm,createKernel,assertStm]                  val create_build_stms = [createProgStm,assertStm] @ buildLog @ [assertStm,createKernel,assertStm]
# Line 605  Line 604 
604    
605    
606                  (* Create Memory Buffers for Strand States and Globals *)                  (* Create Memory Buffers for Strand States and Globals *)
607                  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",
608                                                                          [CL.E_Var tyName]), CL.#*,CL.E_Var numStrandsVar))                                                                          [CL.mkVar tyName]), CL.#*,CL.mkVar numStrandsVar))
609    
610                  val clStrandObjects = [CL.mkAssign(CL.E_Var clInstateVar, CL.mkApply("clCreateBuffer",                  val clStrandObjects = [CL.mkAssign(CL.mkVar clInstateVar, CL.mkApply("clCreateBuffer",
611                                                                  [CL.E_Var contextVar,                                                                  [CL.mkVar contextVar,
612                                                                  CL.E_Var "CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR",                                                                  CL.mkVar "CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR",
613                                                                  CL.E_Var stateSizeVar,                                                                  CL.mkVar stateSizeVar,
614                                                                  CL.E_Var "NULL",                                                                  CL.mkVar "NULL",
615                                                                  CL.E_UnOp(CL.%&,CL.E_Var errVar)])),                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)])),
616                                                           CL.mkAssign(CL.E_Var clOutStateVar, CL.mkApply("clCreateBuffer",                                                           CL.mkAssign(CL.mkVar clOutStateVar, CL.mkApply("clCreateBuffer",
617                                                                  [CL.E_Var contextVar,                                                                  [CL.mkVar contextVar,
618                                                                  CL.E_Var "CL_MEM_READ_WRITE",                                                                  CL.mkVar "CL_MEM_READ_WRITE",
619                                                                  CL.E_Var stateSizeVar,                                                                  CL.mkVar stateSizeVar,
620                                                                  CL.E_Var "NULL",                                                                  CL.mkVar "NULL",
621                                                                  CL.E_UnOp(CL.%&,CL.E_Var errVar)]))]                                                                  CL.mkUnOp(CL.%&,CL.mkVar errVar)]))]
622    
623    
624                  (* Setup up selfOut variable *)                  (* Setup up selfOut variable *)
625                  val strandsArrays = [CL.mkAssign(CL.E_Var outStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.E_Var numStrandsVar,                  val strandsArrays = [CL.mkAssign(CL.mkVar outStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.mkVar numStrandsVar,
626                                                                          CL.#*, CL.mkApply("sizeof",[CL.E_Var tyName]))])),                                                                          CL.#*, CL.mkApply("sizeof",[CL.mkVar tyName]))])),
627                                                                  CL.mkAssign(CL.E_Var inStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.E_Var numStrandsVar,                                                                  CL.mkAssign(CL.mkVar inStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.mkVar numStrandsVar,
628                                                                          CL.#*, CL.mkApply("sizeof",[CL.E_Var tyName]))]))]                                                                          CL.#*, CL.mkApply("sizeof",[CL.mkVar tyName]))]))]
629    
630    
631                  (* Initialize Width Parameter *)                  (* Initialize Width Parameter *)
632                  val widthDel = if nDims = 2 then                  val widthDel = if nDims = 2 then
633                            CL.mkAssign(CL.E_Var "width",CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(1, CL.intTy)))                            CL.mkAssign(CL.mkVar "width",CL.mkSubscript(CL.mkVar globalVar, CL.mkInt(1, CL.intTy)))
634                     else                     else
635                            CL.mkAssign(CL.E_Var "width",CL.E_Int(0,CL.intTy))                            CL.mkAssign(CL.mkVar "width",CL.mkInt(0,CL.intTy))
636    
637    
638                  val strands_init =      CL.mkCall(RN.strandInitSetup,                  val strands_init = CL.mkCall(RN.strandInitSetup, [
639                                                                           [CL.E_Var "size",                          CL.mkVar "size", CL.mkVar "width", CL.mkVar inStateVar
640                                                                           CL.E_Var "width",                        ])
                                                                          CL.E_Var inStateVar])  
641    
642              val clGlobalBuffers = getGlobalDataBuffers(!imgGlobals,3,contextVar,errVar)              val clGlobalBuffers = getGlobalDataBuffers(!imgGlobals,3,contextVar,errVar)
643    
644    
645                  (* Load the Kernel and Header Files *)                  (* Load the Kernel and Header Files *)
646                  val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.E_Var sourcesVar,CL.E_Int(1,CL.intTy)),                  val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(1,CL.intTy)),
647                                                                            CL.mkApply(RN.clLoaderFN, [CL.E_Var clFNVar])),                                                                            CL.mkApply(RN.clLoaderFN, [CL.mkVar clFNVar])),
648             CL.mkAssign(CL.mkSubscript(CL.E_Var sourcesVar,CL.E_Int(0,CL.intTy)),             CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(0,CL.intTy)),
649                                                                            CL.mkApply(RN.clLoaderFN, [CL.E_Var headerFNVar]))]                                                                            CL.mkApply(RN.clLoaderFN, [CL.mkVar headerFNVar]))]
650    
651                  (* val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.E_Var sourcesVar,CL.E_Int(1,CL.intTy)),                  (* val sourceStms = [CL.mkAssign(CL.mkSubscript(CL.mkVar sourcesVar,CL.mkInt(1,CL.intTy)),
652                                                                            CL.mkApply(RN.clLoaderFN, [CL.E_Var clFNVar]))] *)                                                                            CL.mkApply(RN.clLoaderFN, [CL.mkVar clFNVar]))] *)
653    
654    
655                  (* Created Enqueue Statements *)                  (* Created Enqueue Statements *)
656  (* FIXME: simplify this code by function abstraction *)  (* FIXME: simplify this code by function abstraction *)
657          val enqueueStm = if nDims = 1          val enqueueStm = if nDims = 1
658                          then [CL.mkAssign(CL.E_Var errVar,                          then [CL.mkAssign(CL.mkVar errVar,
659                                                            CL.mkApply("clEnqueueNDRangeKernel",                                                            CL.mkApply("clEnqueueNDRangeKernel",
660                                                                                                  [CL.E_Var cmdVar,                                                                                                  [CL.mkVar cmdVar,
661                                                                                                   CL.E_Var kernelVar,                                                                                                   CL.mkVar kernelVar,
662                                                                                                   CL.E_Int(1,CL.intTy),                                                                                                   CL.mkInt(1,CL.intTy),
663                                                                                                   CL.E_Var "NULL",                                                                                                   CL.mkVar "NULL",
664                                                                                                   CL.E_Var globalVar,                                                                                                   CL.mkVar globalVar,
665                                                                                                   CL.E_Var localVar,                                                                                                   CL.mkVar localVar,
666                                                                                                   CL.E_Int(0,CL.intTy),                                                                                                   CL.mkInt(0,CL.intTy),
667                                                                                                   CL.E_Var "NULL",                                                                                                   CL.mkVar "NULL",
668                                                                                                   CL.E_Var "NULL"])),CL.mkCall("clFinish",[CL.E_Var cmdVar])]                                                                                                   CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]
669                          else if nDims = 2  then                          else if nDims = 2  then
670                           [CL.mkAssign(CL.E_Var errVar,                           [CL.mkAssign(CL.mkVar errVar,
671                                                          CL.mkApply("clEnqueueNDRangeKernel",                                                          CL.mkApply("clEnqueueNDRangeKernel",
672                                                                                                  [CL.E_Var cmdVar,                                                                                                  [CL.mkVar cmdVar,
673                                                                                                   CL.E_Var kernelVar,                                                                                                   CL.mkVar kernelVar,
674                                                                                                   CL.E_Int(2,CL.intTy),                                                                                                   CL.mkInt(2,CL.intTy),
675                                                                                                   CL.E_Var "NULL",                                                                                                   CL.mkVar "NULL",
676                                                                                                   CL.E_Var globalVar,                                                                                                   CL.mkVar globalVar,
677                                                                                                   CL.E_Var localVar,                                                                                                   CL.mkVar localVar,
678                                                                                                   CL.E_Int(0,CL.intTy),                                                                                                   CL.mkInt(0,CL.intTy),
679                                                                                                   CL.E_Var "NULL",                                                                                                   CL.mkVar "NULL",
680                                                                                                   CL.E_Var "NULL"])),CL.mkCall("clFinish",[CL.E_Var cmdVar])]                                                                                                   CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]
681                          else                          else
682                            [CL.mkAssign(CL.E_Var errVar,                            [CL.mkAssign(CL.mkVar errVar,
683                                                          CL.mkApply("clEnqueueNDRangeKernel",                                                          CL.mkApply("clEnqueueNDRangeKernel",
684                                                                                                  [CL.E_Var cmdVar,                                                                                                  [CL.mkVar cmdVar,
685                                                                                                   CL.E_Var kernelVar,                                                                                                   CL.mkVar kernelVar,
686                                                                                                   CL.E_Int(3,CL.intTy),                                                                                                   CL.mkInt(3,CL.intTy),
687                                                                                                   CL.E_Var "NULL",                                                                                                   CL.mkVar "NULL",
688                                                                                                   CL.E_Var globalVar,                                                                                                   CL.mkVar globalVar,
689                                                                                                   CL.E_Var localVar,                                                                                                   CL.mkVar localVar,
690                                                                                                   CL.E_Int(0,CL.intTy),                                                                                                   CL.mkInt(0,CL.intTy),
691                                                                                                   CL.E_Var "NULL",                                                                                                   CL.mkVar "NULL",
692                                                                                                   CL.E_Var "NULL"])),CL.mkCall("clFinish",[CL.E_Var cmdVar])]                                                                                                   CL.mkVar "NULL"])),CL.mkCall("clFinish",[CL.mkVar cmdVar])]
693    
694    
695    
696                  (* Setup Global and Local variables *)                  (* Setup Global and Local variables *)
697    
698                  val globalAndlocalStms = if nDims = 1 then                  val globalAndlocalStms = if nDims = 1 then
699                          [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)),
700                                                                     CL.mkSubscript(CL.E_Var "size", CL.E_Int(0,CL.intTy))),                                                                     CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),
701                           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)),
702                                                                    CL.E_Var "16")]                                                                    CL.mkVar "16")]
703    
704    
705                  else if nDims = 2 then                  else if nDims = 2 then
706                          [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)),
707                                                                     CL.mkSubscript(CL.E_Var "size", CL.E_Int(0,CL.intTy))),                                                                     CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),
708                          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)),
709                                                                     CL.mkSubscript(CL.E_Var "size", CL.E_Int(1,CL.intTy))),                                                                     CL.mkSubscript(CL.mkVar "size", CL.mkInt(1,CL.intTy))),
710                          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)),
711                                                                    CL.E_Var "16"),                                                                    CL.mkVar "16"),
712                          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)),
713                                                                    CL.E_Var "16")]                                                                    CL.mkVar "16")]
714    
715                  else                  else
716                          [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)),
717                                                                     CL.mkSubscript(CL.E_Var "size", CL.E_Int(0,CL.intTy))),                                                                     CL.mkSubscript(CL.mkVar "size", CL.mkInt(0,CL.intTy))),
718                          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)),
719                                                                     CL.mkSubscript(CL.E_Var "size", CL.E_Int(1,CL.intTy))),                                                                     CL.mkSubscript(CL.mkVar "size", CL.mkInt(1,CL.intTy))),
720                          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)),
721                                                                     CL.mkSubscript(CL.E_Var "size", CL.E_Int(2,CL.intTy))),                                                                     CL.mkSubscript(CL.mkVar "size", CL.mkInt(2,CL.intTy))),
722                          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)),
723                                                                    CL.E_Var "16"),                                                                    CL.mkVar "16"),
724                          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)),
725                                                                    CL.E_Var "16"),                                                                    CL.mkVar "16"),
726                          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)),
727                                                                    CL.E_Var "16")]                                                                    CL.mkVar "16")]
728    
729    
730    
731                  (* Setup Kernel arguments *)                  (* Setup Kernel arguments *)
732                  val kernelArguments = [CL.mkAssign(CL.E_Var errVar,CL.mkApply("clSetKernelArg",                  val kernelArguments = [CL.mkAssign(CL.mkVar errVar,CL.mkApply("clSetKernelArg",
733                                                                  [CL.E_Var kernelVar,                                                                  [CL.mkVar kernelVar,
734                                                                   CL.E_Int(0,CL.intTy),                                                                   CL.mkInt(0,CL.intTy),
735                                                                   CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
736                                                                   CL.E_UnOp(CL.%&,CL.E_Var clInstateVar)])),                                                                   CL.mkUnOp(CL.%&,CL.mkVar clInstateVar)])),
737                                                              CL.mkExpStm(CL.mkAssignOp(CL.E_Var errVar, CL.|=,CL.mkApply("clSetKernelArg",                                                              CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar, CL.|=,CL.mkApply("clSetKernelArg",
738                                                                  [CL.E_Var kernelVar,                                                                  [CL.mkVar kernelVar,
739                                                                   CL.E_Int(1,CL.intTy),                                                                   CL.mkInt(1,CL.intTy),
740                                                                   CL.mkApply("sizeof",[CL.E_Var "cl_mem"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "cl_mem"]),
741                                                                   CL.E_UnOp(CL.%&,CL.E_Var clOutStateVar)]))),                                                                   CL.mkUnOp(CL.%&,CL.mkVar clOutStateVar)]))),
742                                                                    CL.mkExpStm(CL.mkAssignOp(CL.E_Var errVar, CL.|=,CL.mkApply("clSetKernelArg",                                                                    CL.mkExpStm(CL.mkAssignOp(CL.mkVar errVar, CL.|=,CL.mkApply("clSetKernelArg",
743                                                                  [CL.E_Var kernelVar,                                                                  [CL.mkVar kernelVar,
744                                                                   CL.E_Int(2,CL.intTy),                                                                   CL.mkInt(2,CL.intTy),
745                                                                   CL.mkApply("sizeof",[CL.E_Var "int"]),                                                                   CL.mkApply("sizeof",[CL.mkVar "int"]),
746                                                                   CL.E_UnOp(CL.%&,CL.E_Var "width")])))]                                                                   CL.mkUnOp(CL.%&,CL.mkVar "width")])))]
747    
748             val clGlobalArguments = genGlobalArguments(!imgGlobals,3,kernelVar,errVar) @ [assertStm]             val clGlobalArguments = genGlobalArguments(!imgGlobals,3,kernelVar,errVar) @ [assertStm]
749    
750                  (* Retrieve output *)                  (* Retrieve output *)
751                  val outputStm = CL.mkAssign(CL.E_Var errVar,                  val outputStm = CL.mkAssign(CL.mkVar errVar,
752                                                          CL.mkApply("clEnqueueReadBuffer",                                                          CL.mkApply("clEnqueueReadBuffer",
753                                                                                                  [CL.E_Var cmdVar,                                                                                                  [CL.mkVar cmdVar,
754                                                                                                   CL.E_Var clOutStateVar,                                                                                                   CL.mkVar clOutStateVar,
755                                                                                                   CL.E_Var "CL_TRUE",                                                                                                   CL.mkVar "CL_TRUE",
756                                                                                                   CL.E_Int(0,CL.intTy),                                                                                                   CL.mkInt(0,CL.intTy),
757                                                                                                   CL.E_Var stateSizeVar,                                                                                                   CL.mkVar stateSizeVar,
758                                                                                                   CL.E_Var outStateVar,                                                                                                   CL.mkVar outStateVar,
759                                                                                                   CL.E_Int(0,CL.intTy),                                                                                                   CL.mkInt(0,CL.intTy),
760                                                                                                   CL.E_Var "NULL",                                                                                                   CL.mkVar "NULL",
761                                                                                                   CL.E_Var "NULL"]))                                                                                                   CL.mkVar "NULL"]))
762    
763                  (* Free all the objects *)                  (* Free all the objects *)
764                  val freeStms = [CL.mkCall("clReleaseKernel",[CL.E_Var kernelVar]),                  val freeStms = [CL.mkCall("clReleaseKernel",[CL.mkVar kernelVar]),
765                                                  CL.mkCall("clReleaseProgram",[CL.E_Var programVar ]),                                                  CL.mkCall("clReleaseProgram",[CL.mkVar programVar ]),
766                                                  CL.mkCall("clReleaseCommandQueue",[CL.E_Var cmdVar]),                                                  CL.mkCall("clReleaseCommandQueue",[CL.mkVar cmdVar]),
767                                                  CL.mkCall("clReleaseContext",[CL.E_Var contextVar]),                                                  CL.mkCall("clReleaseContext",[CL.mkVar contextVar]),
768                                                  CL.mkCall("clReleaseMemObject",[CL.E_Var clInstateVar]),                                                  CL.mkCall("clReleaseMemObject",[CL.mkVar clInstateVar]),
769                                                  CL.mkCall("clReleaseMemObject",[CL.E_Var clOutStateVar])]                                                  CL.mkCall("clReleaseMemObject",[CL.mkVar clOutStateVar])]
770    
771    
772                  (*Setup Strand Print Function *)                  (*Setup Strand Print Function *)
773                  val outputData = [CL.mkDecl(CL.T_Ptr(CL.T_Named("FILE")), "outS", SOME(CL.I_Exp(CL.mkApply("fopen",                  val outputData = [CL.mkDecl(CL.T_Ptr(CL.T_Named("FILE")), "outS", SOME(CL.I_Exp(CL.mkApply("fopen",
774                                                  [CL.E_Str "mip.txt",                                                  [CL.mkStr "mip.txt",
775                                                  CL.E_Str "w"])))),                                                  CL.mkStr "w"])))),
776                                                  CL.mkCall(concat[name, "_print"],                                                  CL.mkCall(concat[name, "_print"],
777                                                                          [CL.E_Var "outS",                                                                          [CL.mkVar "outS",
778                                                                           CL.E_Var "size",                                                                           CL.mkVar "size",
779                                                                           CL.E_Var "width",                                                                           CL.mkVar "width",
780                                                                           CL.E_Var outStateVar])]                                                                           CL.mkVar outStateVar])]
781    
782    
783    
# Line 805  Line 803 
803          (*generate code for intilizing kernel global data *)          (*generate code for intilizing kernel global data *)
804          fun initKernelGlobals (globals,imgGlobals) = let          fun initKernelGlobals (globals,imgGlobals) = let
805                  fun initGlobalStruct (CL.D_Var(_, _ , name, _)::rest) =                  fun initGlobalStruct (CL.D_Var(_, _ , name, _)::rest) =
806                                  CL.mkAssign(CL.E_Var name, CL.mkIndirect(CL.E_Var RN.globalsVarName, name)) ::                                  CL.mkAssign(CL.mkVar name, CL.mkIndirect(CL.mkVar RN.globalsVarName, name)) ::
807                                  initGlobalStruct(rest)                                  initGlobalStruct(rest)
808                    | initGlobalStruct ( _::rest) = initGlobalStruct(rest)                    | initGlobalStruct ( _::rest) = initGlobalStruct(rest)
809                    | initGlobalStruct([]) = []                    | initGlobalStruct([]) = []
810    
811                  fun initGlobalImages((name,tyname)::rest) =                  fun initGlobalImages((name,tyname)::rest) =
812                                  CL.mkAssign(CL.E_Var name, CL.E_Var (RN.addBufferSuffix name)) ::                                  CL.mkAssign(CL.mkVar name, CL.mkVar (RN.addBufferSuffix name)) ::
813                                  CL.mkAssign(CL.mkIndirect(CL.E_Var name,"data"),CL.E_Var (RN.addBufferSuffixData name)) ::                                  CL.mkAssign(CL.mkIndirect(CL.mkVar name,"data"),CL.mkVar (RN.addBufferSuffixData name)) ::
814                                  initGlobalImages(rest)                                  initGlobalImages(rest)
815                    | initGlobalImages([]) = []                    | initGlobalImages([]) = []
816                  in                  in
# Line 830  Line 828 
828                        CL.PARAM(["__global"], CL.intTy, "width")                        CL.PARAM(["__global"], CL.intTy, "width")
829                      ] @ genKeneralGlobalParams(!imgGlobals)                      ] @ genKeneralGlobalParams(!imgGlobals)
830                    val thread_ids = if nDims = 1                    val thread_ids = if nDims = 1
831                          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)))),
832                                    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)]))]
833                          else                          else
834                                  [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)))),
835                                   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)))),
836                                    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)])),
837                                    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)]))]
838    
839                    val strandDecl = [CL.mkDecl(CL.T_Named tyName, inState, NONE),                    val strandDecl = [CL.mkDecl(CL.T_Named tyName, inState, NONE),
840                                                          CL.mkDecl(CL.T_Named tyName, outState,NONE)]                                                          CL.mkDecl(CL.T_Named tyName, outState,NONE)]
841                    val strandObjects  = if nDims = 1                    val strandObjects  = if nDims = 1
842                          then [CL.mkAssign(CL.mkSubscript(CL.E_Var "selfIn",CL.E_Str "x"),                          then [CL.mkAssign(CL.mkSubscript(CL.mkVar "selfIn",CL.mkStr "x"),
843                                                                           CL.E_Var inState),                                                                           CL.mkVar inState),
844                                    CL.mkAssign(CL.mkSubscript(CL.E_Var "selfOut",CL.E_Str "x"),                                    CL.mkAssign(CL.mkSubscript(CL.mkVar "selfOut",CL.mkStr "x"),
845                                                                           CL.E_Var outState)]                                                                           CL.mkVar outState)]
846                          else let                          else let
847                                  val index = CL.mkBinOp(CL.mkBinOp(CL.E_Var "x",CL.#*,CL.E_Var "width"),CL.#+,CL.E_Var "y")                                  val index = CL.mkBinOp(CL.mkBinOp(CL.mkVar "x",CL.#*,CL.mkVar "width"),CL.#+,CL.mkVar "y")
848                                  in                                  in
849                                          [CL.mkAssign(CL.mkSubscript(CL.E_Var "selfIn",index),                                          [CL.mkAssign(CL.mkSubscript(CL.mkVar "selfIn",index),
850                                                                          CL.E_Var inState),                                                                          CL.mkVar inState),
851                                           CL.mkAssign(CL.mkSubscript(CL.E_Var "selfOut",index),                                           CL.mkAssign(CL.mkSubscript(CL.mkVar "selfOut",index),
852                                                                          CL.E_Var outState)]                                                                          CL.mkVar outState)]
853                                  end                                  end
854    
855    
856                    val status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy))))                    val status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.mkInt(0, CL.intTy))))
857                    val local_vars =  thread_ids @ initKernelGlobals(!globals,!imgGlobals)  @ strandDecl @ strandObjects @ [status]                    val local_vars =  thread_ids @ initKernelGlobals(!globals,!imgGlobals)  @ strandDecl @ strandObjects @ [status]
858                    val while_exp = CL.mkBinOp(CL.mkBinOp(CL.E_Var "status",CL.#!=, CL.E_Var RN.kStabilize),CL.#||,CL.mkBinOp(CL.E_Var "status", CL.#!=, CL.E_Var RN.kDie))                    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))
859                    val while_body = [CL.mkAssign(CL.E_Var "status", CL.mkApply(RN.strandUpdate name,[ CL.E_UnOp(CL.%&,CL.E_Var inState), CL.E_UnOp(CL.%&,CL.E_Var outState)])),                    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)])),
860                                                          CL.mkCall(RN.strandStabilize name,[ CL.E_UnOp(CL.%&,CL.E_Var inState),  CL.E_UnOp(CL.%&,CL.E_Var outState)])]                                                          CL.mkCall(RN.strandStabilize name,[ CL.mkUnOp(CL.%&,CL.mkVar inState),  CL.mkUnOp(CL.%&,CL.mkVar outState)])]
861    
862                    val whileBlock = [CL.mkWhile(while_exp,CL.mkBlock while_body)]                    val whileBlock = [CL.mkWhile(while_exp,CL.mkBlock while_body)]
863    
# Line 879  Line 877 
877        (* generate the table of strand descriptors *)        (* generate the table of strand descriptors *)
878          fun genStrandTable (ppStrm, strands) = let          fun genStrandTable (ppStrm, strands) = let
879                val nStrands = length strands                val nStrands = length strands
880                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)))
881                fun genInits (_, []) = []                fun genInits (_, []) = []
882                  | genInits (i, s::ss) = (i, genInit s) :: genInits(i+1, ss)                  | genInits (i, s::ss) = (i, genInit s) :: genInits(i+1, ss)
883                fun ppDecl dcl = PrintAsC.output(ppStrm, dcl)                fun ppDecl dcl = PrintAsC.output(ppStrm, dcl)
884                in                in
885                  ppDecl (CL.D_Var([], CL.int32, RN.numStrands,                  ppDecl (CL.D_Var([], CL.int32, RN.numStrands,
886                    SOME(CL.I_Exp(CL.E_Int(IntInf.fromInt nStrands, CL.int32)))));                    SOME(CL.I_Exp(CL.mkInt(IntInf.fromInt nStrands, CL.int32)))));
887                  ppDecl (CL.D_Var([],                  ppDecl (CL.D_Var([],
888                    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),
889                    RN.strands,                    RN.strands,

Legend:
Removed from v.1277  
changed lines
  Added in v.1278

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