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

SCM Repository

[diderot] View of /branches/vis12-cl/src/compiler/cl-target/cl-util.sml
ViewVC logotype

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

Parent Directory Parent Directory | Revision Log Revision Log


Revision 2694 - (download) (annotate)
Wed Sep 10 22:55:58 2014 UTC (4 years, 10 months ago) by jhr
File size: 5889 byte(s)
  Woeking on OpenCL support
(* cl-util.sml
 *
 * COPYRIGHT (c) 2014 The Diderot Project (http://diderot-language.cs.uchicago.edu)
 * All rights reserved.
 *
 * Utility functions for generating C and OpenCL code.
 *)

structure CLUtil : sig

  (* helper functions for specifying parameters in various address spaces *)
    val globalParam : (CLang.ty * CLang.var) -> CLang.param
    val constantParam : (CLang.ty * CLang.var) -> CLang.param
    val localParam : (CLang.ty * CLang.var) -> CLang.param
    val privateParam : (CLang.ty * CLang.var) -> CLang.param
    val clParam : (CLang.ty * CLang.var) -> CLang.param

  (* OpenCL global pointer type *)
    val globalPtr : CLang.ty -> CLang.ty

  (* make OpenCL kernel functions *)
    val mkKernel : string * CLang.param list * CLang.stm -> CLang.decl
    val mkTaskKernel : string * CLang.param list * CLang.stm -> CLang.decl

  (* make OpenCL kernel invocations *)
    val mkEnqueueKernel : {
	    retSts : CLang.var option,
	    cmdQ : CLang.exp,
	    kernel : CLang.exp,
	    workDim : int,
	    globalWorkSz : CLang.exp,
	    localWorkSz : CLang.exp,
	    args : (CLang.var * CLang.exp) list,	(* (arg, size) *)
	    onError : CLang.var -> CLang.stm
	  } -> CLang.stm list
    val mkEnqueueTask : {
	    retSts : CLang.var option,
	    cmdQ : CLang.exp,
	    kernel : CLang.exp,
	    args : (CLang.var * CLang.exp) list,	(* (arg, size) *)
	    onError : CLang.var -> CLang.stm
	  } -> CLang.stm list

  end = struct

    local
      structure CL = CLang

      fun mapi f l = let
	    fun mapf (i, [], l) = List.rev l
	      | mapf (i, x::xs, l) = mapf (i+1, xs, f(i, x)::l)
	    in
	      mapf (0, l, [])
	    end

      fun mkInt i = CL.mkInt(IntInf.fromInt i)
    in

  (* helper functions for specifying parameters in various address spaces *)
    local
      fun param spc (ty, x) = CL.PARAM([spc], ty, x)
    in
    val globalParam = param "__global"
    val constantParam = param "__constant"
    val localParam = param "__local"
    val privateParam = param "__private"
    fun clParam (ty, x) = CL.PARAM([], ty, x)
    end (* local *)

  (* OpenCL global pointer type *)
    fun globalPtr ty = CL.T_Qual("__global", CL.T_Ptr ty)

  (* OpenCL kernel declaration *)
    fun mkKernel (f, params, body) = CL.D_Func(
	  ["__kernel", "__attribute__((reqd_work_group_size(BLK_SZ, 1, 1)))"],
	  CL.voidTy,
	  f,
	  params,
	  body)

  (* OpenCL task kernel declaration; (i.e., a kernel that has a single instance) *)
    fun mkTaskKernel (f, params, body) = CL.D_Func(
	  ["__kernel"],
	  CL.voidTy,
	  f,
	  params,
	  body)

  (* helper function for generating a kernel invocation; the code has the
   * following form:
	if (((errSts = clSetKernelArg(...)) != CL_SUCCESS)
        ||  ...
        ||  ((errSts = clSetKernelArg(...)) != CL_SUCCESS)
        ||  ((errSts = clEnqueueNDRangeKernel(...)) != CL_SUCCESS) {
            /* error handler */
        }
   *)
    fun mkEnqueueKernel {
	    retSts : CL.var option,
	    cmdQ : CL.exp,
	    kernel : CL.exp,
	    workDim : int,
	    globalWorkSz : CL.exp,
	    localWorkSz : CL.exp,
	    args : (CL.var * CL.exp) list,
	    onError : CL.var -> CL.stm
	  } = let
	  val (retSts, retStsDecl) = (case retSts
		 of SOME sts => (sts, [])
		  | NONE => let
		      val sts = "errSts"
		      in
			(sts, [CL.mkDecl(CLang.T_Named "cl_int", sts, NONE)])
		      end
		(* end case *))
	  fun mkApply (f, args) = CL.mkAssignOp(CL.mkVar retSts, CL.$=, CL.mkApply(f, args))
	  fun initArg (i, (arg, argSize)) = let
		val setArg = mkApply("clSetKernelArg", [
			kernel,
			CL.mkInt(IntInf.fromInt i),
			argSize,
			CL.mkAddrOf(CL.mkVar arg)
		      ])
		in
		  setArg
		end
	  val argExps = mapi initArg args
	  val applyExp = mkApply("clEnqueueNDRangeKernel", [
		  cmdQ,			(* cl_command_queue command_queue *)
		  kernel,		(* cl_kernel kernel *)
		  mkInt workDim,	(* cl_uint work_dim *)
		  mkInt 0,		(* const size_t *global_work_offset; // must be null *)
		  globalWorkSz,		(* const size_t *global_work_size *)
		  localWorkSz,		(* const size_t *local_work_size *)
		  mkInt 0,		(* cl_uint num_events_in_wait_list *)
		  mkInt 0,		(* const cl_event *event_wait_list *)
		  mkInt 0		(* cl_event *event *)
		])
	  val condExp = List.foldr (fn (e1, e2) => CL.mkBinOp(e1, CL.#||, e2)) applyExp argExps
	  in
	    retStsDecl @ [
		CL.mkIfThen(
		  condExp,
		  onError retSts)
	      ]
	  end

  (* helper function for generating a kernel task invocation; the code has the
   * following form:
	if (((errSts = clSetKernelArg(...)) != CL_SUCCESS)
        ||  ...
        ||  ((errSts = clSetKernelArg(...)) != CL_SUCCESS)
        ||  ((errSts = clEnqueueTask(...)) != CL_SUCCESS) {
            /* error handler */
        }
   *)
    fun mkEnqueueTask {
	    retSts : CL.var option,
	    cmdQ : CL.exp,
	    kernel : CL.exp,
	    args : (CL.var * CL.exp) list,
	    onError : CL.var -> CL.stm
	  } = let
	  val (retSts, retStsDecl) = (case retSts
		 of SOME sts => (sts, [])
		  | NONE => let
		      val sts = "errSts"
		      in
			(sts, [CL.mkDecl(CLang.T_Named "cl_int", sts, NONE)])
		      end
		(* end case *))
	  fun mkApply (f, args) = CL.mkAssignOp(CL.mkVar retSts, CL.$=, CL.mkApply(f, args))
	  fun initArg (i, (arg, argSize)) = let
		val setArg = mkApply("clSetKernelArg", [
			kernel,
			CL.mkInt(IntInf.fromInt i),
			argSize,
			CL.mkAddrOf(CL.mkVar arg)
		      ])
		in
		  setArg
		end
	  val argExps = mapi initArg args
	  val applyExp = mkApply("clEnqueueTask", [
		  cmdQ,			(* cl_command_queue command_queue *)
		  kernel,		(* cl_kernel kernel *)
		  mkInt 0,		(* cl_uint num_events_in_wait_list *)
		  mkInt 0,		(* const cl_event *event_wait_list *)
		  mkInt 0		(* cl_event *event *)
		])
	  val condExp = List.foldr (fn (e1, e2) => CL.mkBinOp(e1, CL.#||, e2)) applyExp argExps
	  in
	    retStsDecl @ [
		CL.mkIfThen(
		  condExp,
		  onError retSts)
	      ]
	  end

    end (* local *)

  end

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