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/cl-util.sml
ViewVC logotype

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

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

revision 2647, Thu May 29 17:48:44 2014 UTC revision 2648, Thu May 29 18:22:21 2014 UTC
# Line 9  Line 9 
9    
10      local      local
11        structure CL = CLang        structure CL = CLang
12    
13          fun mapi f l = let
14                fun mapf (i, [], l) = List.rev l
15                  | mapf (i, x::xs, l) = mapf (i+1, xs, f(i, x)::l)
16                in
17                  mapf (0, l, [])
18                end
19    
20          fun mkInt i = CL.mkInt(IntInf.fromInt i)
21      in      in
22    
23    (* helper functions for specifying parameters in various address spaces *)    (* helper functions for specifying parameters in various address spaces *)
# Line 25  Line 34 
34    (* OpenCL global pointer type *)    (* OpenCL global pointer type *)
35      fun globalPtr ty = CL.T_Qual("__global", CL.T_Ptr ty)      fun globalPtr ty = CL.T_Qual("__global", CL.T_Ptr ty)
36    
37      (* OpenCL kernel declaration *)
38        fun mkKernel (f, params, body) = CL.D_Func(
39              ["__kernel", "__attribute__((reqd_work_group_size(BLK_SZ, 1, 1)))"],
40              CL.voidTy,
41              f,
42              params,
43              body)
44    
45      (* OpenCL task kernel declaration; (i.e., a kernel that has a single instance) *)
46        fun mkTaskKernel (f, params, body) = CL.D_Func(
47              ["__kernel"],
48              CL.voidTy,
49              f,
50              params,
51              body)
52    
53      (* helper function for generating a kernel invocation; the code has the
54       * following form:
55            if (((errSts = clSetKernelArg(...)) != CL_SUCCESS)
56            ||  ...
57            ||  ((errSts = clSetKernelArg(...)) != CL_SUCCESS)
58            ||  ((errSts = clEnqueueNDRangeKernel(...)) != CL_SUCCESS) {
59                /* error handler */
60            }
61       *)
62        fun mkEnqueueKernel {
63                retSts : CL.var option,
64                cmdQ : CL.exp,
65                kernel : CL.exp,
66                workDim : int,
67                globalWorkSz : CL.exp,
68                localWorkSz : CL.exp,
69                args : ?? list,
70                onError : CL.var -> CL.stm
71              } = let
72              val (retSts, retStsDecl) = (case retSts
73                     of SOME sts => (sts, [])
74                      | NONE => let
75                          val sts = "errSts"
76                          in
77                            (sts, [CL.mkDecl(CLang.T_Named "cl_int", sts, NONE)])
78                          end
79                    (* end case *))
80              fun mkApply (f, args) = CL.mkAssignOp(CL.mkVar retSts, CL.$=, CL.mkApply(f, args))
81              fun initArg (i, ??) = let
82                    val setArg = mkApply("clSetKernelArg", [
83                            kernel,
84                            CL.mkInt(IntInf.fromInt i),
85                            argSize,
86                            argValue
87                          ])
88                    in
89                      setArg
90                    end
91              val argExps = mapi initArg args
92              val applyExp = mkApply("clEnqueueNDRangeKernel", [
93                      cmdQ,                 (* cl_command_queue command_queue *)
94                      kernel,               (* cl_kernel kernel *)
95                      mkInt workDim,        (* cl_uint work_dim *)
96                      mkInt 0,              (* const size_t *global_work_offset; // must be null *)
97                      globalWorkSz,         (* const size_t *global_work_size *)
98                      localWorkSz,          (* const size_t *local_work_size *)
99                      mkInt 0,              (* cl_uint num_events_in_wait_list *)
100                      mkInt 0,              (* const cl_event *event_wait_list *)
101                      mkInt 0               (* cl_event *event *)
102                    ])
103              val condExp = List.foldr (fn (e1, e2) => CL.mkBinOp(e1, CL.#||, e2)) applyExp argExps
104              in
105                retStsDecl @ [
106                    applyStm,
107                    CL.mkIfThen(
108                      condExp,
109                      onError retSts)
110                  ]
111              end
112    
113      end (* local *)      end (* local *)
114    
115    end    end

Legend:
Removed from v.2647  
changed lines
  Added in v.2648

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