diff --git a/inst/doc/custom_ocl_gpuR.R b/inst/doc/custom_ocl_gpuR.R index 454b9b0..0ce5673 100644 --- a/inst/doc/custom_ocl_gpuR.R +++ b/inst/doc/custom_ocl_gpuR.R @@ -4,6 +4,15 @@ opts_chunk$set( concordance=TRUE ) +## ----kernel_example, eval = FALSE---------------------------------------- +# "__kernel void SAXPY(__global float* x, __global float* y, float a) +# { +# const int i = get_global_id(0); +# +# y [i] += a * x [i]; +# } +# " + ## ----cl_setup, eval=FALSE------------------------------------------------ # cl_args <- setup_opencl(objects = c("vclVector", "vclVector", "scalar"), # intents = c("IN", "OUT", "IN"), diff --git a/inst/doc/custom_ocl_gpuR.Rnw b/inst/doc/custom_ocl_gpuR.Rnw index b858ccb..3391b22 100644 --- a/inst/doc/custom_ocl_gpuR.Rnw +++ b/inst/doc/custom_ocl_gpuR.Rnw @@ -57,14 +57,16 @@ Essentially, these files contain routines/functions that are compiled for use on different devices (e.g. CPU, GPU, FPGA, etc.). One example is the SAXPY function -```{kernel_example} -__kernel void SAXPY(__global float* x, __global float* y, float a) + +<>= +"__kernel void SAXPY(__global float* x, __global float* y, float a) { - const int i = get_global_id (0); + const int i = get_global_id(0); y [i] += a * x [i]; } -``` +" +@ The exact definitions and specifications of OpenCL kernels is beyond the scope of this vignette and the interested user is recommended to consult additional diff --git a/inst/doc/custom_ocl_gpuR.pdf b/inst/doc/custom_ocl_gpuR.pdf index 301027d..4432975 100644 Binary files a/inst/doc/custom_ocl_gpuR.pdf and b/inst/doc/custom_ocl_gpuR.pdf differ diff --git a/inst/doc/gpuR.pdf b/inst/doc/gpuR.pdf index 4a39f2a..edb213a 100644 Binary files a/inst/doc/gpuR.pdf and b/inst/doc/gpuR.pdf differ diff --git a/vignettes/custom_ocl_gpuR.R b/vignettes/custom_ocl_gpuR.R deleted file mode 100644 index 15e2778..0000000 --- a/vignettes/custom_ocl_gpuR.R +++ /dev/null @@ -1,35 +0,0 @@ -## ----setup, include=FALSE, cache=FALSE----------------------------------- -library(knitr) -opts_chunk$set( -concordance=TRUE -) - -## ----kernel_example, eval = FALSE---------------------------------------- -# __kernel void SAXPY(__global float* x, __global float* y, float a) -# { -# const int i = get_global_id(0); -# -# y [i] += a * x [i]; -# } - -## ----cl_setup, eval=FALSE------------------------------------------------ -# cl_args <- setup_opencl(objects = c("vclVector", "vclVector", "scalar"), -# intents = c("IN", "OUT", "IN"), -# queues = list("SAXPY", "SAXPY", "SAXPY"), -# kernel_maps = c("x", "y", "a")) - -## ----compile, eval=FALSE------------------------------------------------- -# custom_opencl("saxpy.cl", cl_args, "float") - -## ----example, eval=FALSE------------------------------------------------- -# -# a <- rnorm(16) -# b <- rnorm(16) -# gpuA <- vclVector(a, type = "float") -# gpuB <- vclVector(b, type = "float") -# scalar <- 2 -# -# # apply custom function -# # equivalent to - scalar*a + b -# saxpy(gpuA, gpuB, scalar) - diff --git a/vignettes/custom_ocl_gpuR.Rnw b/vignettes/custom_ocl_gpuR.Rnw index 419e5f2..3391b22 100644 --- a/vignettes/custom_ocl_gpuR.Rnw +++ b/vignettes/custom_ocl_gpuR.Rnw @@ -59,12 +59,13 @@ function <>= -__kernel void SAXPY(__global float* x, __global float* y, float a) +"__kernel void SAXPY(__global float* x, __global float* y, float a) { const int i = get_global_id(0); y [i] += a * x [i]; } +" @ The exact definitions and specifications of OpenCL kernels is beyond the scope diff --git a/vignettes/custom_ocl_gpuR.pdf b/vignettes/custom_ocl_gpuR.pdf deleted file mode 100644 index 9f69579..0000000 Binary files a/vignettes/custom_ocl_gpuR.pdf and /dev/null differ diff --git a/vignettes/custom_ocl_gpuR.tex b/vignettes/custom_ocl_gpuR.tex deleted file mode 100644 index bac3c1c..0000000 --- a/vignettes/custom_ocl_gpuR.tex +++ /dev/null @@ -1,186 +0,0 @@ -% \VignetteIndexEntry{ Custom OpenCL Kernels } -% \VignettePackage{gpuR} -% \VignetteEngine{knitr::knitr} - -% To compile this document -% library('knitr'); rm(list=ls()); knit('custom_ocl_gpuR.Rnw') - -\documentclass[12pt]{article}\usepackage[]{graphicx}\usepackage[]{color} -%% maxwidth is the original width if it is less than linewidth -%% otherwise use linewidth (to make sure the graphics do not exceed the margin) -\makeatletter -\def\maxwidth{ % - \ifdim\Gin@nat@width>\linewidth - \linewidth - \else - \Gin@nat@width - \fi -} -\makeatother - -\definecolor{fgcolor}{rgb}{0.345, 0.345, 0.345} -\newcommand{\hlnum}[1]{\textcolor[rgb]{0.686,0.059,0.569}{#1}}% -\newcommand{\hlstr}[1]{\textcolor[rgb]{0.192,0.494,0.8}{#1}}% -\newcommand{\hlcom}[1]{\textcolor[rgb]{0.678,0.584,0.686}{\textit{#1}}}% -\newcommand{\hlopt}[1]{\textcolor[rgb]{0,0,0}{#1}}% -\newcommand{\hlstd}[1]{\textcolor[rgb]{0.345,0.345,0.345}{#1}}% -\newcommand{\hlkwa}[1]{\textcolor[rgb]{0.161,0.373,0.58}{\textbf{#1}}}% -\newcommand{\hlkwb}[1]{\textcolor[rgb]{0.69,0.353,0.396}{#1}}% -\newcommand{\hlkwc}[1]{\textcolor[rgb]{0.333,0.667,0.333}{#1}}% -\newcommand{\hlkwd}[1]{\textcolor[rgb]{0.737,0.353,0.396}{\textbf{#1}}}% -\let\hlipl\hlkwb - -\usepackage{framed} -\makeatletter -\newenvironment{kframe}{% - \def\at@end@of@kframe{}% - \ifinner\ifhmode% - \def\at@end@of@kframe{\end{minipage}}% - \begin{minipage}{\columnwidth}% - \fi\fi% - \def\FrameCommand##1{\hskip\@totalleftmargin \hskip-\fboxsep - \colorbox{shadecolor}{##1}\hskip-\fboxsep - % There is no \\@totalrightmargin, so: - \hskip-\linewidth \hskip-\@totalleftmargin \hskip\columnwidth}% - \MakeFramed {\advance\hsize-\width - \@totalleftmargin\z@ \linewidth\hsize - \@setminipage}}% - {\par\unskip\endMakeFramed% - \at@end@of@kframe} -\makeatother - -\definecolor{shadecolor}{rgb}{.97, .97, .97} -\definecolor{messagecolor}{rgb}{0, 0, 0} -\definecolor{warningcolor}{rgb}{1, 0, 1} -\definecolor{errorcolor}{rgb}{1, 0, 0} -\newenvironment{knitrout}{}{} % an empty environment to be redefined in TeX - -\usepackage{alltt} -\usepackage[sc]{mathpazo} -\usepackage[T1]{fontenc} -\usepackage{geometry} -\geometry{verbose,tmargin=2.5cm,bmargin=2.5cm,lmargin=2.5cm,rmargin=2.5cm} -\setcounter{secnumdepth}{2} -\setcounter{tocdepth}{2} -\usepackage{url} -\usepackage[unicode=true,pdfusetitle, - bookmarks=true,bookmarksnumbered=true,bookmarksopen=true,bookmarksopenlevel=2, - breaklinks=false,pdfborder={0 0 1},backref=false,colorlinks=false] - {hyperref} -\hypersetup{ - pdfstartview={XYZ null null 1}} - -\newcommand{\pkg}[1]{{\fontseries{b}\selectfont #1}} -\renewcommand{\pkg}[1]{{\textsf{#1}}} - -\newcommand{\Rpackage}[1]{\textsl{#1}} -\newcommand\CRANpkg[1]{% - {\href{http://cran.fhcrc.org/web/packages/#1/index.html}% - {\Rpackage{#1}}}} -\newcommand\Githubpkg[1]{\GithubSplit#1\relax} -\def\GithubSplit#1/#2\relax{{\href{https://github.com/#1/#2}% - {\Rpackage{#2}}}} - -\newcommand{\Rcode}[1]{\texttt{#1}} -\newcommand{\Rfunction}[1]{\Rcode{#1}} -\newcommand{\Robject}[1]{\Rcode{#1}} -\newcommand{\Rclass}[1]{\textit{#1}} -\IfFileExists{upquote.sty}{\usepackage{upquote}}{} -\begin{document} - - - -\title{Custom OpenCL Kernels} -\author{Dr. Charles Determan Jr. PhD\footnote{cdetermanjr@gmail.com}} -\newpage - -\maketitle -\section{Introduction} -At the heart of portable GPU computing is the use of OpenCL kernel files. -Essentially, these files contain routines/functions that are compiled for use -on different devices (e.g. CPU, GPU, FPGA, etc.). One example is the SAXPY -function - - -\begin{knitrout} -\definecolor{shadecolor}{rgb}{0.969, 0.969, 0.969}\color{fgcolor}\begin{kframe} -\begin{alltt} -__kernel void \hlkwd{SAXPY}(__global float* x, __global float* y, float a) -\{ - const int i = \hlkwd{get_global_id}(0); - - y [i] += a * x [i]; -\} -\end{alltt} -\end{kframe} -\end{knitrout} - -The exact definitions and specifications of OpenCL kernels is beyond the scope -of this vignette and the interested user is recommended to consult additional -resource. - -Now, to compile a function to use the OpenCL kernel requires the definition of -OpenCL contexts, buffers, defining global/local sizes, enqueuing operations, and -managing data memory copies. The intent here is to make the use of a custom OpenCL -kernel as seamless as possible within the \Rpackage{gpuR} package. - -\newpage -\maketitle -\section{Demo} - -In order to create the dynamic functions to leverage the OpenCL kernels, -the user should be able to say what the purpose of each argument is. -Taking the SAXPY example above, the user must know which arguments are device -objects (i.e. OpenCL buffers) and which aren't. In this case, the first two -arguments I will denote as \Rfunction{vclVector} objects and the third as a scalar -argument. Likewise, they are denoted by the objects intent, which kernel they -will be passed to (only relevant when more than one kernel function), and the -corresponding argument name in the OpenCL kernel function. The setup call would -look like the following. - -\begin{knitrout} -\definecolor{shadecolor}{rgb}{0.969, 0.969, 0.969}\color{fgcolor}\begin{kframe} -\begin{alltt} -\hlstd{cl_args} \hlkwb{<-} \hlkwd{setup_opencl}\hlstd{(}\hlkwc{objects} \hlstd{=} \hlkwd{c}\hlstd{(}\hlstr{"vclVector"}\hlstd{,} \hlstr{"vclVector"}\hlstd{,} \hlstr{"scalar"}\hlstd{),} - \hlkwc{intents} \hlstd{=} \hlkwd{c}\hlstd{(}\hlstr{"IN"}\hlstd{,} \hlstr{"OUT"}\hlstd{,} \hlstr{"IN"}\hlstd{),} - \hlkwc{queues} \hlstd{=} \hlkwd{list}\hlstd{(}\hlstr{"SAXPY"}\hlstd{,} \hlstr{"SAXPY"}\hlstd{,} \hlstr{"SAXPY"}\hlstd{),} - \hlkwc{kernel_maps} \hlstd{=} \hlkwd{c}\hlstd{(}\hlstr{"x"}\hlstd{,} \hlstr{"y"}\hlstd{,} \hlstr{"a"}\hlstd{))} -\end{alltt} -\end{kframe} -\end{knitrout} - -Once the argument definitions are setup, the kernel file and argument definitions -can be passed to the \Rfunction{custom\_opencl} function which also takes one additional -argument. OpenCL is a typed language and thefore the user must denote the precision. - -\begin{knitrout} -\definecolor{shadecolor}{rgb}{0.969, 0.969, 0.969}\color{fgcolor}\begin{kframe} -\begin{alltt} -\hlkwd{custom_opencl}\hlstd{(}\hlstr{"saxpy.cl"}\hlstd{, cl_args,} \hlstr{"float"}\hlstd{)} -\end{alltt} -\end{kframe} -\end{knitrout} - -You will notice that there is no assignment set during this call. The internal -function is compiled and exported to the current global environment with the name -of the kernel file (prior to the extension). - -The function can then be called as a normal function on the respective objects. - -\begin{knitrout} -\definecolor{shadecolor}{rgb}{0.969, 0.969, 0.969}\color{fgcolor}\begin{kframe} -\begin{alltt} -\hlstd{a} \hlkwb{<-} \hlkwd{rnorm}\hlstd{(}\hlnum{16}\hlstd{)} -\hlstd{b} \hlkwb{<-} \hlkwd{rnorm}\hlstd{(}\hlnum{16}\hlstd{)} -\hlstd{gpuA} \hlkwb{<-} \hlkwd{vclVector}\hlstd{(a,} \hlkwc{type} \hlstd{=} \hlstr{"float"}\hlstd{)} -\hlstd{gpuB} \hlkwb{<-} \hlkwd{vclVector}\hlstd{(b,} \hlkwc{type} \hlstd{=} \hlstr{"float"}\hlstd{)} -\hlstd{scalar} \hlkwb{<-} \hlnum{2} - -\hlcom{# apply custom function} -\hlcom{# equivalent to - scalar*a + b} -\hlkwd{saxpy}\hlstd{(gpuA, gpuB, scalar)} -\end{alltt} -\end{kframe} -\end{knitrout} - -\end{document}