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

SCM Repository

[diderot] Diff of /branches/vis12-cl/src/compiler/cl-target/gen-output.sml
ViewVC logotype

Diff of /branches/vis12-cl/src/compiler/cl-target/gen-output.sml

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

revision 3136, Wed Mar 25 18:33:59 2015 UTC revision 3137, Thu Mar 26 14:55:52 2015 UTC
# Line 77  Line 77 
77      fun sizes i = CL.mkSubscript(sizesV, mkInt i)      fun sizes i = CL.mkSubscript(sizesV, mkInt i)
78      fun setSizes (i, v) = CL.mkAssign(sizes i, v)      fun setSizes (i, v) = CL.mkAssign(sizes i, v)
79    
80        local
81          fun get_group_id i = CL.mkApply("get_group_id", [CL.mkInt i])
82          fun get_num_groups i = CL.mkApply("get_num_groups", [CL.mkInt i])
83          fun get_local_id i = CL.mkApply("get_local_id", [CL.mkInt i])
84        in
85      (* the expression to compute the base strand index from the group and local ids.
86       * This calculation is based on a 2D arrangement, where the first dimension is
87       * equal to the number of CUs times the CU width, while the second dimension is
88       * equal to the number of workers/CU.
89       *
90       *    int idx = BLK_SZ * (get_group_id(0) + get_num_groups(0)*get_group_id(1)) + get_local_id(0)
91       *)
92        val indexStm = CL.mkDeclInit(
93              CL.uint32, "idx",
94              CL.mkBinOp(
95                CL.mkBinOp(
96                  CL.mkVar "BLK_SZ",
97                  CL.#*,
98                  CL.mkBinOp(
99                    get_group_id 0,
100                    CL.#+,
101                    CL.mkBinOp(get_num_groups 0, CL.#*, get_group_id 1))),
102                CL.#+,
103                get_local_id 0))
104      (* the expression to compute the strand offset from the group and local ids.
105       *
106       *    int offset = BLK_SZ * get_num_groups(0) * get_num_groups(1)
107       *)
108        val offsetStm = CL.mkDeclInit(
109              CL.uint32, "offset",
110              CL.mkBinOp(get_num_groups 0, CL.#*, CL.mkBinOp(get_num_groups 1, CL.#*, CL.mkVar "BLK_SZ")))
111        end (* local *)
112    
113    (* create a kernel for copying the given output state variable to the output    (* create a kernel for copying the given output state variable to the output
114     * buffer.     * buffer.
115     *)     *)
# Line 87  Line 120 
120      fun mkCopyKernel tgt strandTy (ty : TreeIL.Ty.ty, name) = let      fun mkCopyKernel tgt strandTy (ty : TreeIL.Ty.ty, name) = let
121            val (ty', nElems) = CLTyTranslate.toOutputType ty            val (ty', nElems) = CLTyTranslate.toOutputType ty
122            val body = CL.mkBlock[            val body = CL.mkBlock[
123                    CL.mkDeclInit(CL.uint32, "idx",                    indexStm,
                     CL.mkBinOp(  
                       CL.mkBinOp(CL.mkApply("get_group_id", [CL.mkInt 0]), CL.#*, CL.mkVar "BLK_SZ"),  
                       CL.#+,  
                       CL.mkApply("get_local_id", [CL.mkInt 0]))),  
124                    CL.mkDeclInit(CL.uint32, "offset",                    CL.mkDeclInit(CL.uint32, "offset",
125                      CL.mkBinOp(CL.mkApply("get_num_groups", [CL.mkInt 0]), CL.#*, CL.mkVar "BLK_SZ")),                      CL.mkBinOp(CL.mkApply("get_num_groups", [CL.mkInt 0]), CL.#*, CL.mkVar "BLK_SZ")),
126                    CL.S_Decl(["__global"], strandTy, "state", SOME(CL.I_Exp(CL.mkIndirect(CL.mkVar "sched", "state")))),                    CL.S_Decl(["__global"], strandTy, "state", SOME(CL.I_Exp(CL.mkIndirect(CL.mkVar "sched", "state")))),

Legend:
Removed from v.3136  
changed lines
  Added in v.3137

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