From 21f00482b15d9e6004e47c9ed45bdedd1eb112b6 Mon Sep 17 00:00:00 2001 From: Aaron Puchert Date: Mon, 6 Jul 2026 23:34:51 +0200 Subject: [PATCH] Add support for raw vectors as buffers and kernel output We can treat raw buffers as `char` arrays in OpenCL. Fixes #25. --- R/clBuffer.R | 7 +++++-- R/ocl.R | 2 +- man/clBuffer.Rd | 9 +++++---- man/oclRun.Rd | 5 ++--- man/oclSimpleKernel.Rd | 10 +++++----- src/buffer.c | 33 +++++++++++++++++++++++++++++---- src/ocl.h | 1 + tests/buffer.R | 12 ++++++++++++ tests/buffer.Rout.save | 20 ++++++++++++++++++++ tests/kernel.R | 6 ++++++ tests/kernel.Rout.save | 8 ++++++++ tests/kernel.cl | 8 ++++++++ 12 files changed, 102 insertions(+), 19 deletions(-) diff --git a/R/clBuffer.R b/R/clBuffer.R index 9b7682b..3507420 100644 --- a/R/clBuffer.R +++ b/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 diff --git a/R/ocl.R b/R/ocl.R index 54bb9c7..a313c38 100644 --- a/R/ocl.R +++ b/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) diff --git a/man/clBuffer.Rd b/man/clBuffer.Rd index fc186c3..dc061b0 100644 --- a/man/clBuffer.Rd +++ b/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) diff --git a/man/oclRun.Rd b/man/oclRun.Rd index b33b31b..ce40b44 100644 --- a/man/oclRun.Rd +++ b/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 diff --git a/man/oclSimpleKernel.Rd b/man/oclSimpleKernel.Rd index c46c402..d05426c 100644 --- a/man/oclSimpleKernel.Rd +++ b/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 diff --git a/src/buffer.c b/src/buffer.c index 5f0d55f..410f0f1 100644 --- a/src/buffer.c +++ b/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) diff --git a/src/ocl.h b/src/ocl.h index ea41d93..582257c 100644 --- a/src/ocl.h +++ b/src/ocl.h @@ -20,6 +20,7 @@ extern SEXP oclMessageSymbol; /* Supported buffer data types */ typedef enum { + CLT_RAW, CLT_INT, CLT_FLOAT, CLT_DOUBLE diff --git a/tests/buffer.R b/tests/buffer.R index a94cb2e..250db6f 100644 --- a/tests/buffer.R +++ b/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)) diff --git a/tests/buffer.Rout.save b/tests/buffer.Rout.save index b17e925..1fa9b17 100644 --- a/tests/buffer.Rout.save +++ b/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)) diff --git a/tests/kernel.R b/tests/kernel.R index 695fdd6..e109195 100644 --- a/tests/kernel.R +++ b/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 diff --git a/tests/kernel.Rout.save b/tests/kernel.Rout.save index ee13d65..f8c8eec 100644 --- a/tests/kernel.Rout.save +++ b/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 diff --git a/tests/kernel.cl b/tests/kernel.cl index 39c7097..853cf30 100644 --- a/tests/kernel.cl +++ b/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); +}