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 1256, Tue May 24 17:58:28 2011 UTC revision 1261, Sat May 28 23:16:30 2011 UTC
# Line 172  Line 172 
172                          stms @ getGlobals(globals)                          stms @ getGlobals(globals)
173                  end                  end
174    
175          (* register the code that is used to register command-line options for input variables *)
176            fun inputs (Prog{topDecls, ...}, stm) = let
177                  val inputsFn = CL.D_Func(
178                        [], CL.voidTy, RN.registerOpts,
179                        [CL.PARAM([], CL.T_Ptr(CL.T_Named RN.optionsTy), "opts")],
180                        stm)
181                  in
182                    topDecls := inputsFn :: !topDecls
183                  end
184    
185          fun init (Prog{globals,topDecls,...}, CL.S_Block(init)) = let          fun init (Prog{globals,topDecls,...}, CL.S_Block(init)) = let
186                val params = [                val params = [
187                            CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName)                            CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName)
188                          ]                          ]
189                     val body = CL.S_Block(globalIndirects(!globals,init))                     val body = CL.S_Block(globalIndirects(!globals,init))
190                     val initFn = CL.D_Func([], CL.voidTy, RN.initGlobals, params, body)                     val initFn = CL.D_Func([], CL.voidTy, RN.initGlobals, params, body)
   
191                in                in
192                          topDecls := initFn :: !topDecls                          topDecls := initFn :: !topDecls
193                end                end
   
194            | init (Prog{globals,topDecls,...}, init) = let            | init (Prog{globals,topDecls,...}, init) = let
195                val params = [                val params = [
196                            CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName)                            CL.PARAM([], CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName)
197                          ]                          ]
198                     val initFn = CL.D_Func([], CL.voidTy, RN.initGlobals, params, init)                     val initFn = CL.D_Func([], CL.voidTy, RN.initGlobals, params, init)
   
199                in                in
200                  topDecls := initFn :: !topDecls                  topDecls := initFn :: !topDecls
201                end                end
# Line 295  Line 302 
302    
303    
304          (* generates the load kernel function *)          (* generates the load kernel function *)
305    (* FIXME: this code might be part of the runtime system *)
306          fun genKernelLoader() =          fun genKernelLoader() =
307                  CL.D_Verbatim ( ["/* Loads the Kernel from a file */",                  CL.D_Verbatim ( ["/* Loads the Kernel from a file */",
308                                                  "char * loadKernel (const char * filename) {",                                                  "char * loadKernel (const char * filename) {",
# Line 322  Line 330 
330    
331          fun genDataBuffers([],_,_,_) = []          fun genDataBuffers([],_,_,_) = []
332            | genDataBuffers((var,nDims)::globals,count,contextVar,errVar) = let            | genDataBuffers((var,nDims)::globals,count,contextVar,errVar) = let
333    (* FIXME: use CL constructors to  build expressions (not strings) *)
334                     val size = if nDims = 1 then                     val size = if nDims = 1 then
335                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.E_Var "float"]), CL.#*,                                          CL.mkBinOp(CL.mkApply("sizeof",[CL.E_Var "float"]), CL.#*,
336                                           CL.mkIndirect(CL.E_Var var, "size[0]"))                                           CL.mkIndirect(CL.E_Var var, "size[0]"))
# Line 398  Line 407 
407                  end                  end
408          (* generates the host-side setup function *)          (* generates the host-side setup function *)
409          fun genHostSetupFunc(strand as Strand{name,tyName,...}, filename, nDims, initially, imgGlobals, oneDim, twoDim, thirdDim) = let          fun genHostSetupFunc(strand as Strand{name,tyName,...}, filename, nDims, initially, imgGlobals, oneDim, twoDim, thirdDim) = let
410                  (*Delcare opencl setup objects *)                (* Declare opencl setup objects *)
411                  val programVar= "program"                  val programVar= "program"
412                  val kernelVar = "kernel"                  val kernelVar = "kernel"
413                  val cmdVar = "queue"                  val cmdVar = "queue"
# Line 425  Line 434 
434                  val params = [                  val params = [
435                           CL.PARAM([],CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName)                           CL.PARAM([],CL.T_Ptr(CL.T_Named RN.globalsTy), RN.globalsVarName)
436                           ]                           ]
437                  val delcarations = [CL.mkDecl(CL.clProgramTy, programVar, NONE),                  val declarations = [CL.mkDecl(CL.clProgramTy, programVar, NONE),
438                            CL.mkDecl(CL.clKernelTy, kernelVar, NONE),                            CL.mkDecl(CL.clKernelTy, kernelVar, NONE),
439                            CL.mkDecl(CL.clCmdQueueTy, cmdVar, NONE),                            CL.mkDecl(CL.clCmdQueueTy, cmdVar, NONE),
440                            CL.mkDecl(CL.clContextTy, contextVar, NONE),                            CL.mkDecl(CL.clContextTy, contextVar, NONE),
# Line 512  Line 521 
521                                                                            CL.mkApply(RN.clLoaderFN, [CL.E_Var headerFNVar]))]                                                                            CL.mkApply(RN.clLoaderFN, [CL.E_Var headerFNVar]))]
522    
523                  (* Created Enqueue Statements *)                  (* Created Enqueue Statements *)
524    (* FIXME: simplify this code by function abstraction *)
525                  val enqueueStm = if nDims = 1                  val enqueueStm = if nDims = 1
526                          then [CL.mkAssign(CL.E_Var errVar,                          then [CL.mkAssign(CL.E_Var errVar,
527                                                            CL.mkApply("clEnqueueNDRangeKernel",                                                            CL.mkApply("clEnqueueNDRangeKernel",
# Line 552  Line 562 
562                  (* Setup up selfOut variable *)                  (* Setup up selfOut variable *)
563                  val selfOutStm = CL.mkAssign(CL.E_Var outStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.E_Var numStrandsVar,                  val selfOutStm = CL.mkAssign(CL.E_Var outStateVar, CL.mkApply("malloc", [CL.mkBinOp(CL.E_Var numStrandsVar,
564                                                                          CL.#*, CL.mkApply("sizeof",[CL.E_Var tyName]))]))                                                                          CL.#*, CL.mkApply("sizeof",[CL.E_Var tyName]))]))
   
565                  (* Setup Global and Local variables *)                  (* Setup Global and Local variables *)
566                  val globalAndlocalStms = if nDims = 1
567                  val globalAndlocalStms = if nDims = 1 then                      then [CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(0,CL.intTy)),
                         [CL.mkAssign(CL.mkSubscript(CL.E_Var globalVar, CL.E_Int(0,CL.intTy)),  
568                                                                     CL.mkSubscript(CL.E_Var "size", CL.E_Int(0,CL.intTy))),                                                                     CL.mkSubscript(CL.E_Var "size", CL.E_Int(0,CL.intTy))),
569                           CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(0,CL.intTy)),                           CL.mkAssign(CL.mkSubscript(CL.E_Var localVar, CL.E_Int(0,CL.intTy)),
570                                                                    CL.E_Var "16")]                                                                    CL.E_Var "16")]
# Line 608  Line 616 
616             val clGlobalArguments = genGlobalArguments(!imgGlobals,3,kernelVar,errVar)             val clGlobalArguments = genGlobalArguments(!imgGlobals,3,kernelVar,errVar)
617    
618                  (* Retrieve output *)                  (* Retrieve output *)
619                  val outputStm = CL.mkAssign(CL.E_Var errVar,                val outputStm = CL.mkAssign(CL.E_Var errVar, CL.mkApply("clEnqueueReadBuffer",
                                                         CL.mkApply("clEnqueueReadBuffer",  
620                                                                                                  [CL.E_Var cmdVar,                                                                                                  [CL.E_Var cmdVar,
621                                                                                                   CL.E_Var clOutStateVar,                                                                                                   CL.E_Var clOutStateVar,
622                                                                                                   CL.E_Var "CL_TRUE",                                                                                                   CL.E_Var "CL_TRUE",
# Line 619  Line 626 
626                                                                                                   CL.E_Int(0,CL.intTy),                                                                                                   CL.E_Int(0,CL.intTy),
627                                                                                                   CL.E_Var "NULL",                                                                                                   CL.E_Var "NULL",
628                                                                                                   CL.E_Var "NULL"]))                                                                                                   CL.E_Var "NULL"]))
   
629                  (* Free all the objects *)                  (* Free all the objects *)
630                  val freeStms = [CL.mkCall("clReleaseKernel",[CL.E_Var kernelVar]),                  val freeStms = [CL.mkCall("clReleaseKernel",[CL.E_Var kernelVar]),
631                                                  CL.mkCall("clReleaseProgram",[CL.E_Var programVar ]),                                                  CL.mkCall("clReleaseProgram",[CL.E_Var programVar ]),
# Line 627  Line 633 
633                                                  CL.mkCall("clReleaseContext",[CL.E_Var contextVar]),                                                  CL.mkCall("clReleaseContext",[CL.E_Var contextVar]),
634                                                  CL.mkCall("clReleaseMemObject",[CL.E_Var clInstateVar]),                                                  CL.mkCall("clReleaseMemObject",[CL.E_Var clInstateVar]),
635                                                  CL.mkCall("clReleaseMemObject",[CL.E_Var clOutStateVar])]                                                  CL.mkCall("clReleaseMemObject",[CL.E_Var clOutStateVar])]
   
636                  (* Body put all the statments together *)                  (* Body put all the statments together *)
637                  val body =  delcarations @ platformStm @ devicesStm @ contextStm @ commandStm @ !initially @ [strandSize] @                val body =  declarations @ platformStm @ devicesStm @ contextStm @ commandStm @ !initially @ [strandSize] @
638                                     clStrandObjects @ clGlobalBuffers @ sourceStms  @ [selfOutStm] @ globalAndlocalStms @                                     clStrandObjects @ clGlobalBuffers @ sourceStms  @ [selfOutStm] @ globalAndlocalStms @
639                                     kernelArguments @ clGlobalArguments @ enqueueStm @  [outputStm] @ freeStms                                     kernelArguments @ clGlobalArguments @ enqueueStm @  [outputStm] @ freeStms
640    
641                  in                  in
   
642                  CL.D_Func([],CL.voidTy,RN.setupFName,params,CL.mkBlock(body))                  CL.D_Func([],CL.voidTy,RN.setupFName,params,CL.mkBlock(body))
   
643                  end                  end
644    
645    
# Line 653  Line 656 
656                    val thread_ids = if nDims = 1                    val thread_ids = if nDims = 1
657                          then [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy)))),                          then [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy)))),
658                                    CL.mkAssign(CL.E_Var "x",CL.mkApply(RN.getGlobalThreadId,[CL.E_Int(0,CL.intTy)]))]                                    CL.mkAssign(CL.E_Var "x",CL.mkApply(RN.getGlobalThreadId,[CL.E_Int(0,CL.intTy)]))]
659                          else                       else [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy)))),
                                 [CL.mkDecl(CL.intTy, "x", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy)))),  
660                                   CL.mkDecl(CL.intTy, "y", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy)))),                                   CL.mkDecl(CL.intTy, "y", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy)))),
661                                    CL.mkAssign(CL.E_Var "x",  CL.mkApply(RN.getGlobalThreadId,[CL.E_Int(0,CL.intTy)])),                                    CL.mkAssign(CL.E_Var "x",  CL.mkApply(RN.getGlobalThreadId,[CL.E_Int(0,CL.intTy)])),
662                                    CL.mkAssign(CL.E_Var "y",CL.mkApply(RN.getGlobalThreadId,[CL.E_Int(1,CL.intTy)]))]                                    CL.mkAssign(CL.E_Var "y",CL.mkApply(RN.getGlobalThreadId,[CL.E_Int(1,CL.intTy)]))]
   
663                    val strandDecl = [CL.mkDecl(CL.T_Named tyName, inState, NONE),                    val strandDecl = [CL.mkDecl(CL.T_Named tyName, inState, NONE),
664                                                          CL.mkDecl(CL.T_Named tyName, outState,NONE)]                                                          CL.mkDecl(CL.T_Named tyName, outState,NONE)]
665                    val strandObjects  = if nDims = 1                    val strandObjects  = if nDims = 1
666                          then [CL.mkAssign(CL.mkSubscript(CL.E_Var "selfIn",CL.E_Str "x"),                       then [CL.mkAssign(CL.mkSubscript(CL.E_Var "selfIn",CL.E_Str "x"), CL.E_Var inState),
667                                                                           CL.E_Var inState),                             CL.mkAssign(CL.mkSubscript(CL.E_Var "selfOut",CL.E_Str "x"), CL.E_Var outState)]
                                   CL.mkAssign(CL.mkSubscript(CL.E_Var "selfOut",CL.E_Str "x"),  
                                                                          CL.E_Var outState)]  
668                          else let                          else let
669                                  val index = CL.mkBinOp(CL.mkBinOp(CL.E_Var "y",CL.#*,CL.E_Var "width"),CL.#+,CL.E_Var "x")                                  val index = CL.mkBinOp(CL.mkBinOp(CL.E_Var "y",CL.#*,CL.E_Var "width"),CL.#+,CL.E_Var "x")
670                                  in                                  in
671                                          [CL.mkAssign(CL.mkSubscript(CL.E_Var "selfIn",index),                           [CL.mkAssign(CL.mkSubscript(CL.E_Var "selfIn",index), CL.E_Var inState),
672                                                                          CL.E_Var inState),                            CL.mkAssign(CL.mkSubscript(CL.E_Var "selfOut",index), CL.E_Var outState)]
                                          CL.mkAssign(CL.mkSubscript(CL.E_Var "selfOut",index),  
                                                                         CL.E_Var outState)]  
673                                  end                                  end
674                    val status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy))))                    val status = CL.mkDecl(CL.intTy, "status", SOME(CL.I_Exp(CL.E_Int(0, CL.intTy))))
675                    val strand_init_function = CL.mkCall(RN.strandInit name, [CL.E_Var inState])                    val strand_init_function = CL.mkCall(RN.strandInit name, [CL.E_Var inState])
# Line 683  Line 680 
680                                                          CL.mkIfThen(CL.mkBinOp(CL.E_Var "status",CL.#==, CL.E_Var RN.kStabilize),CL.mkBreak)]                                                          CL.mkIfThen(CL.mkBinOp(CL.E_Var "status",CL.#==, CL.E_Var RN.kStabilize),CL.mkBreak)]
681    
682                    val whileBlock = [CL.mkWhile(while_exp,CL.mkBlock while_body)]                    val whileBlock = [CL.mkWhile(while_exp,CL.mkBlock while_body)]
   
683                    val body = CL.mkBlock(local_vars  @ whileBlock)                    val body = CL.mkBlock(local_vars  @ whileBlock)
684                  in                  in
685                     CL.D_Func(["__kernel"], CL.voidTy, fName, params, body)                     CL.D_Func(["__kernel"], CL.voidTy, fName, params, body)
# Line 722  Line 718 
718                fun cppDecl dcl = PrintAsC.output(cppStrm, dcl)                fun cppDecl dcl = PrintAsC.output(cppStrm, dcl)
719                fun clppDecl dcl = PrintAsC.output(clppStrm, dcl)                fun clppDecl dcl = PrintAsC.output(clppStrm, dcl)
720                val strands = AtomTable.listItems strands                val strands = AtomTable.listItems strands
721                val single_strand as Strand{name, tyName, code, ...}= hd(strands)                val singleStrand as Strand{name, tyName, code, ...} = hd(strands)
722                in                in
723              (* Generate the Host file .c *)              (* Generate the Host file .c *)
724              cppDecl (CL.D_Verbatim([ "#include <OpenCL/OpenCL.h>",                  cppDecl (CL.D_Verbatim([
725                                                                   "#include Diderot/diderot.h"]));                      "#include <OpenCL/OpenCL.h>",
726                        "#include Diderot/diderot.h"
727                      ]));
728                  List.app cppDecl (List.rev (!globals));                  List.app cppDecl (List.rev (!globals));
729              cppDecl (genGlobalStruct (!globals));              cppDecl (genGlobalStruct (!globals));
730              cppDecl (genStrandTyDef single_strand);                  cppDecl (genStrandTyDef singleStrand);
731              cppDecl (genKernelLoader());              cppDecl (genKernelLoader());
732              List.app cppDecl (List.rev (!topDecls));              List.app cppDecl (List.rev (!topDecls));
733              cppDecl (genHostSetupFunc (single_strand,clFileName,!numDims,initially,imgGlobals,oneDim,twoDim,thirdDim));                  cppDecl (genHostSetupFunc (
734                        singleStrand, clFileName, !numDims, initially,
735                        imgGlobals, oneDim, twoDim, thirdDim));
736              cppDecl (genHostMain());              cppDecl (genHostMain());
737    
738              (* Generate the OpenCl file *)              (* Generate the OpenCl file *)
739              clppDecl (genGlobalStruct (!globals));              clppDecl (genGlobalStruct (!globals));
740              clppDecl (genStrandTyDef single_strand);                  clppDecl (genStrandTyDef singleStrand);
741              List.app clppDecl (!code);              List.app clppDecl (!code);
742              clppDecl (genKernelFun (single_strand,!numDims));                  clppDecl (genKernelFun (singleStrand,!numDims));
743    
744                  (*List.app (fn strand => List.app ppDecl (genStrand strand)) strands;                  (*List.app (fn strand => List.app ppDecl (genStrand strand)) strands;
745                   genStrandTable (ppStrm, strands);                   genStrandTable (ppStrm, strands);
# Line 779  Line 779 
779                  (*RunCC.compile (basename, cflags);                  (*RunCC.compile (basename, cflags);
780                  RunCC.link (basename, ldOpts)*)                  RunCC.link (basename, ldOpts)*)
781    
   
782        end        end
783    
784    (* strands *)    (* strands *)

Legend:
Removed from v.1256  
changed lines
  Added in v.1261

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