Add support for raw vectors as buffers and kernel output

R/clBuffer.R

@@ -1,5 +1,5 @@
 # Creating a buffer in a context
-clBuffer <- function(context, length, mode = c("numeric", "single", "double", "integer"))
+clBuffer <- function(context, length, mode = c("numeric", "single", "double", "integer", "raw"))
 {
     if (!inherits(context, "clContext"))
         stop("Invalid context")
@@ -20,6 +20,9 @@ as.double.clBuffer <- function(x, ...) {
 as.integer.clBuffer <- function(x, ...) {
     as.integer(.Call(cl_read_buffer, x, NULL))
 }
+as.raw.clBuffer <- function(x, ...) {
+    as.raw(.Call(cl_read_buffer, x, NULL))
+}
 is.clBuffer <- function(any) inherits(any, "clBuffer")
 
 # Printing information about the buffer
@@ -67,7 +70,7 @@ length.clBuffer<- function(x) {
     # Determine expected class for value.
     targetClass <- switch(attributes(x)$mode,
                           single=, double="numeric",
-                          integer="integer",
+                          integer="integer", raw="raw",
                           stop("Invalid buffer class ", attributes(x)$mode))
 
     # Convert if necessary - target is either numeric or integer

R/ocl.R

@@ -90,7 +90,7 @@ oclContext <- function(device = "default", precision = c("best", "single", "doub
 }
 
 # Compile a "simple kernel"
-oclSimpleKernel <- function(context, name, code, output.mode = c("numeric", "single", "double", "integer")) {
+oclSimpleKernel <- function(context, name, code, output.mode = c("numeric", "single", "double", "integer", "raw")) {
     if (!inherits(context,"clContext"))
         stop("invalid context")
     output.mode <- match.arg(output.mode)

man/clBuffer.Rd

@@ -3,6 +3,7 @@
 \alias{as.clBuffer}
 \alias{as.double.clBuffer}
 \alias{as.integer.clBuffer}
+\alias{as.raw.clBuffer}
 \alias{is.clBuffer}
 \alias{print.clBuffer}
 \alias{length.clBuffer}
@@ -18,15 +19,15 @@
 
   Just like vectors in R, OpenCL buffers have a mode, which is (as of now) one
   of "double" or "numeric" (corresponds to \code{double} in OpenCL C), "single"
-  (\code{float}) or "integer" (\code{int}).
+  (\code{float}), "integer" (\code{int}), or "raw" (\code{char}).
 
   The constructor \code{clBuffer} takes a context as created by
   \code{oclContext}, a length and a mode argument.
 
   The conversion function \code{as.clBuffer} creates an OpenCL buffer of the
   same length and mode as the argument and copies the data. Conversely,
-  \code{as.double} (= \code{as.numeric}) and \code{as.integer} read a buffer and
-  coerce the result as vector the appropriate mode.
+  \code{as.double} (= \code{as.numeric}), \code{as.integer} and \code{as.raw}
+  read a buffer and coerce the result as vector the appropriate mode.
 
   With \code{is.clBuffer} one can check if an object is an OpenCL buffer.
 
@@ -48,7 +49,7 @@
   \code{value} must have the same length as the buffer (no recycling).
 }
 \usage{
-clBuffer(context, length, mode = c("numeric", "single", "double", "integer"))
+clBuffer(context, length, mode = c("numeric", "single", "double", "integer", "raw"))
 as.clBuffer(vector, context, mode = class(vector))
 is.clBuffer(any)
 \method{as.double}{clBuffer}(x, \dots)

man/oclRun.Rd

@@ -26,9 +26,8 @@ oclRun(kernel, size, ..., dim = size)
 \details{
   \code{oclRun} pushes kernel arguments, executes the kernel and
   retrieves the result. The kernel is expected to have either
-  \code{__global double *} or \code{__global float *}
-  type (write-only) as the first argument which will be used for the
-  result and \code{const unsigned int} second argument denoting the result
+  \code{__global T *} type (write-only) as the first argument which will be used
+  for the result and \code{unsigned int} second argument denoting the result
   length. All other arguments are assumed to be read-only and will be
   filled according to the \code{\dots} values.
   These can either be OpenCL buffers as generated by \code{\link{clBuffer}} for

man/oclSimpleKernel.Rd

@@ -9,7 +9,7 @@
 }
 \usage{
 oclSimpleKernel(context, name, code,
-    output.mode = c("numeric", "single", "double", "integer"))
+    output.mode = c("numeric", "single", "double", "integer", "raw"))
 }
 \arguments{
   \item{context}{Context (as created by \code{\link{oclContext}})
@@ -21,8 +21,8 @@ oclSimpleKernel(context, name, code,
     concatenated (as-is, no newlines are added!) by the engine.}
   \item{output.mode}{Mode of the output argument of the kernel, as in
     \code{\link{clBuffer}}. This can be one of "single", "double", "integer",
-    or "numeric". The default value "numeric" maps to the default precision of
-    the context.
+    "raw" or "numeric". The default value "numeric" maps to the default
+    precision of the context.
 
     The kernel code may use a type \code{numeric} that is typedef'd to the
     given precision, i.e. either \code{float} or \code{double}. The OpenCL
@@ -35,8 +35,8 @@ oclSimpleKernel(context, name, code,
 
   The kernel built by this function is simple in that it can have
   exactly one vector output and arbitrarily many inputs. The first
-  argument of the kernel must be \code{__global double*} or
-  \code{__global float*} for the output and the second argument must be
+  argument of the kernel must be \code{__global T*} with \code{T} corresponding
+  to \code{output.mode} for the output and the second argument must be
   \code{const unsigned int} for the length of the output vector.
   Additional numeric scalar arguments are assumed to have the same mode as the
   output, i.e. if the output shall have "double" precision, then numeric scalar

src/buffer.c

@@ -20,6 +20,8 @@ ClType get_type(SEXP mode_exp)
         Rf_error("invalid mode");
 
     const char* mode = CHAR(STRING_ELT(mode_exp, 0));
+    if (!strcmp(mode, "raw"))
+        return CLT_RAW;
     if (!strcmp(mode, "integer"))
         return CLT_INT;
     if (!strcmp(mode, "single"))
@@ -33,6 +35,7 @@ ClType get_type(SEXP mode_exp)
 SEXP get_type_description(ClType type)
 {
     switch (type) {
+    case CLT_RAW: return Rf_mkString("raw");
     case CLT_INT: return Rf_mkString("integer");
     case CLT_FLOAT: return Rf_mkString("single");
     case CLT_DOUBLE: return Rf_mkString("double");
@@ -44,6 +47,7 @@ SEXP get_type_description(ClType type)
 static size_t get_element_size(ClType type)
 {
     switch (type) {
+    case CLT_RAW: return sizeof(cl_char);
     case CLT_INT: return sizeof(cl_int);
     case CLT_FLOAT: return sizeof(cl_float);
     case CLT_DOUBLE: return sizeof(cl_double);
@@ -55,6 +59,7 @@ static size_t get_element_size(ClType type)
 static SEXPTYPE get_sexptype(ClType type)
 {
     switch (type) {
+    case CLT_RAW: return RAWSXP;
     case CLT_INT: return INTSXP;
     case CLT_FLOAT: return REALSXP;
     case CLT_DOUBLE: return REALSXP;
@@ -152,6 +157,17 @@ attribute_visible SEXP cl_supported_index(SEXP indices) {
 #endif
 }
 
+static void* getDataPtr(ClType type, SEXP values)
+{
+    switch (type) {
+    case CLT_RAW: return RAW(values);
+    case CLT_INT: return INTEGER(values);
+    case CLT_FLOAT: break; // handled elsewhere
+    case CLT_DOUBLE: return REAL(values);
+    }
+    return NULL;
+}
+
 /* Read data from an OpenCL buffer */
 attribute_visible SEXP cl_read_buffer(SEXP buffer_exp, SEXP indices)
 {
@@ -208,8 +224,7 @@ attribute_visible SEXP cl_read_buffer(SEXP buffer_exp, SEXP indices)
     }
 
     last_ocl_error = clEnqueueReadBuffer(queue, buffer, CL_TRUE, i0, read_size,
-					 (type == CLT_FLOAT) ? (void*)intermediate :
-					 ((type == CLT_INT) ? (void*)INTEGER(res) : (void*)REAL(res)),
+					 (type == CLT_FLOAT) ? (void*)intermediate : getDataPtr(type, res),
 					 wait ? 1 : 0, wait ? &wait : NULL, NULL);
     if (last_ocl_error != CL_SUCCESS) {
         if (type == CLT_FLOAT)
@@ -229,6 +244,17 @@ attribute_visible SEXP cl_read_buffer(SEXP buffer_exp, SEXP indices)
     return res;
 }
 
+static const void* getDataPtrReadOnly(ClType type, SEXP values)
+{
+    switch (type) {
+    case CLT_RAW: return RAW_RO(values);
+    case CLT_INT: return INTEGER_RO(values);
+    case CLT_FLOAT: break; // handled elsewhere
+    case CLT_DOUBLE: return REAL_RO(values);
+    }
+    return NULL;
+}
+
 /* Write data to an OpenCL buffer */
 attribute_visible SEXP cl_write_buffer(SEXP buffer_exp, SEXP indices, SEXP values)
 {
@@ -302,8 +328,7 @@ attribute_visible SEXP cl_write_buffer(SEXP buffer_exp, SEXP indices, SEXP value
 
     /* Note that we do not have to block here (other than for mem mgmt reasons) */
     last_ocl_error = clEnqueueWriteBuffer(queue, buffer, CL_TRUE, i0, write_size,
-					  (type == CLT_FLOAT) ? (const void*)intermediate :
-					  ((type == CLT_INT) ? (const void*)INTEGER_RO(values) : (const void*)REAL_RO(values)),
+					  (type == CLT_FLOAT) ? (const void*)intermediate : getDataPtrReadOnly(type, values),
 					  0, NULL, NULL);
     if (last_ocl_error != CL_SUCCESS) {
         if (type == CLT_FLOAT)

src/ocl.h

@@ -20,6 +20,7 @@ extern SEXP oclMessageSymbol;
 
 /* Supported buffer data types */
 typedef enum {
+    CLT_RAW,
     CLT_INT,
     CLT_FLOAT,
     CLT_DOUBLE

tests/buffer.R

@@ -48,11 +48,23 @@ print(attributes(ints)$mode)
 print(ints)
 length(ints)
 
+#    The same for a raw buffer
+bytes <- clBuffer(ctx, 32, "raw")
+bytes[] <- as.raw(16:47)
+
+#    Inspect the resulting buffer
+class(bytes)
+print(attributes(bytes)$mode)
+print(bytes)
+length(bytes)
+
 # 3. Let's see if we can guess the type correctly
 numeric.buf <- as.clBuffer(sqrt(3:18), ctx)
 print(numeric.buf)
 integer.buf <- as.clBuffer(-1:14, ctx)
 print(integer.buf)
+byte.buf <- as.clBuffer(as.raw(1:16), ctx)
+print(byte.buf)
 
 # 4. Other functions
 c(is.clBuffer(ints), is.clBuffer(1:16), is.clBuffer(ctx))

tests/buffer.Rout.save

@@ -80,6 +80,22 @@ OpenCL buffer, 32 elements of type integer
 > length(ints)
 [1] 32
 > 
+> #    The same for a raw buffer
+> bytes <- clBuffer(ctx, 32, "raw")
+> bytes[] <- as.raw(16:47)
+> 
+> #    Inspect the resulting buffer
+> class(bytes)
+[1] "clBuffer"
+> print(attributes(bytes)$mode)
+[1] "raw"
+> print(bytes)
+OpenCL buffer, 32 elements of type raw 
+ [1] 10 11 12 13 14 15 16 17 18 19 1a 1b 1c 1d 1e 1f 20 21 22 23 24 25 26 27 28
+[26] 29 2a 2b 2c 2d 2e 2f
+> length(bytes)
+[1] 32
+> 
 > # 3. Let's see if we can guess the type correctly
 > numeric.buf <- as.clBuffer(sqrt(3:18), ctx)
 > print(numeric.buf)
@@ -90,6 +106,10 @@ OpenCL buffer, 16 elements of type double
 > print(integer.buf)
 OpenCL buffer, 16 elements of type integer 
  [1] -1  0  1  2  3  4  5  6  7  8  9 10 11 12 13 14
+> byte.buf <- as.clBuffer(as.raw(1:16), ctx)
+> print(byte.buf)
+OpenCL buffer, 16 elements of type raw 
+ [1] 01 02 03 04 05 06 07 08 09 0a 0b 0c 0d 0e 0f 10
 > 
 > # 4. Other functions
 > c(is.clBuffer(ints), is.clBuffer(1:16), is.clBuffer(ctx))

tests/kernel.R

@@ -26,3 +26,9 @@ multiply <- oclSimpleKernel(ctx, "multiply", code)
 input <- as.clBuffer((1:16)^2, ctx)
 output <- oclRun(multiply, 16, input, 2.5)
 output
+
+# 4. Run kernel with a raw buffer argument
+swap <- oclSimpleKernel(ctx, "swap", code, "raw")
+input <- as.clBuffer(as.raw(c(0x12, 0x34, 0x56, 0x78)), ctx)
+output <- oclRun(swap, 4, input)
+output

tests/kernel.Rout.save

@@ -50,6 +50,14 @@ OpenCL buffer, 16 elements of type double
  [1]   2.5  10.0  22.5  40.0  62.5  90.0 122.5 160.0 202.5 250.0 302.5 360.0
 [13] 422.5 490.0 562.5 640.0
 > 
+> # 4. Run kernel with a raw buffer argument
+> swap <- oclSimpleKernel(ctx, "swap", code, "raw")
+> input <- as.clBuffer(as.raw(c(0x12, 0x34, 0x56, 0x78)), ctx)
+> output <- oclRun(swap, 4, input)
+> output
+OpenCL buffer, 4 elements of type raw 
+[1] 21 43 65 87
+> 
 > proc.time()
    user  system elapsed 
   0.133   0.043   0.180 

tests/kernel.cl

@@ -24,3 +24,11 @@ __kernel void multiply(__global numeric* output,
 	size_t i = get_global_id(0);
 	output[i] = factor * input[i];
 }
+
+__kernel void swap(__global char* output,
+                   const unsigned int count,
+                   __global const char* input)
+{
+	size_t i = get_global_id(0);
+	output[i] = ((input[i] & 0x0f) << 4) | (input[i] >> 4);
+}