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

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

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