Merge pull request rapidsai#10 from sklam/ccode

Establish basic C library template

.gitignore

@@ -0,0 +1,5 @@
+__pycache__
+*.pyc
+*.o
+*.so
+.cache

Makefile

@@ -0,0 +1,11 @@
+all: libgdf.so
+
+libgdf.so: src/column.cpp src/unaryops.cu src/binaryops.cu src/errorhandling.cpp
+	nvcc -Iinclude -shared -o $@ $+
+
+test:
+	python setup.py build_ext --inplace
+	py.test -v
+
+clean:
+	rm -f libgdf.so

include/gdf/cffi/functions.h

@@ -0,0 +1,20 @@
+/* column operations */
+
+gdf_size_type gdf_column_sizeof();
+
+gdf_error gdf_column_view(gdf_column *column, void *data, gdf_valid_type *valid,
+                          gdf_size_type size, gdf_dtype dtype);
+
+/* error handling */
+
+const char * gdf_error_get_name(gdf_error errcode);
+
+/* unary operators */
+
+gdf_error gdf_sin_generic(gdf_column *input, gdf_column *output);
+
+gdf_error gdf_sin_f32(gdf_column *input, gdf_column *output);
+
+/* binary operators */
+
+gdf_error gdf_add_f32(gdf_column *lhs, gdf_column *rhs, gdf_column *output);

include/gdf/cffi/types.h

@@ -0,0 +1,26 @@
+typedef size_t gdf_size_type;
+typedef gdf_size_type gdf_index_type;
+typedef unsigned char gdf_valid_type;
+
+typedef enum {
+    GDF_invalid=0,
+    GDF_INT8,
+    GDF_INT16,
+    GDF_INT32,
+    GDF_FLOAT32,
+    GDF_FLOAT64,
+} gdf_dtype;
+
+typedef enum {
+    GDF_SUCCESS=0,
+    GDF_CUDA_ERROR,
+    GDF_UNSUPPORTED_DTYPE,
+    GDF_COLUMN_SIZE_MISMATCH,
+} gdf_error;
+
+typedef struct gdf_column_{
+    void *data;
+    gdf_valid_type *valid;
+    gdf_size_type size;
+    gdf_dtype dtype;
+} gdf_column;

include/gdf/errorutils.h

@@ -0,0 +1,8 @@
+#ifndef GDF_ERRORUTILS_H
+#define GDF_ERRORUTILS_H
+
+#define CUDA_TRY(x) if ((x)!=cudaSuccess) return GDF_CUDA_ERROR;
+
+#define CUDA_CHECK_LAST() CUDA_TRY(cudaGetLastError())
+
+#endif // GDF_ERRORUTILS_H

include/gdf/gdf.h

@@ -0,0 +1,14 @@
+#ifndef GDF_GDF_H
+#define GDF_GDF_H
+
+#include <cstdlib>
+
+#include "cffi/types.h"
+
+#define GDF_VALID_BITSIZE (sizeof(gdf_valid_type) * 8)
+
+extern "C" {
+#include "cffi/functions.h"
+}
+
+#endif /* GDF_GDF_H */

include/gdf/utils.h

@@ -0,0 +1,15 @@
+#ifndef GDF_UTILS_H
+#define GDF_UTILS_H
+
+#include <gdf/gdf.h>
+
+__device__
+static
+bool gdf_is_valid(const gdf_valid_type *valid, gdf_index_type pos) {
+    if ( valid )
+        return (valid[pos / GDF_VALID_BITSIZE] >> (pos % GDF_VALID_BITSIZE)) & 1;
+    else
+        return true;
+}
+
+#endif

libgdf_cffi/__init__.py

@@ -0,0 +1,14 @@
+from __future__ import absolute_import
+
+from .wrapper import _libgdf_wrapper
+from .wrapper import GDFError           # re-exported
+
+try:
+    from .libgdf_cffi import ffi
+except ImportError:
+    pass
+else:
+    libgdf_api = ffi.dlopen('libgdf.so')
+    libgdf = _libgdf_wrapper(ffi, libgdf_api)
+
+    del _libgdf_wrapper

libgdf_cffi/libgdf_build.py

@@ -0,0 +1,11 @@
+import cffi
+
+ffibuilder = cffi.FFI()
+ffibuilder.set_source("libgdf_cffi.libgdf_cffi", None)
+
+for fname in ['types.h', 'functions.h']:
+    with open('include/gdf/cffi/{}'.format(fname), 'r') as fin:
+        ffibuilder.cdef(fin.read())
+
+if __name__ == "__main__":
+    ffibuilder.compile()

libgdf_cffi/libgdf_cffi.py

@@ -0,0 +1,11 @@
+# auto-generated file
+import _cffi_backend
+
+ffi = _cffi_backend.FFI('libgdf_cffi.libgdf_cffi',
+    _version = 0x2601,
+    _types = b'\x00\x00\x15\x0D\x00\x00\x01\x0B\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x17\x03\x00\x00\x04\x11\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x04\x11\x00\x00\x04\x11\x00\x00\x04\x11\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x04\x11\x00\x00\x19\x03\x00\x00\x18\x03\x00\x00\x1C\x01\x00\x00\x00\x0B\x00\x00\x00\x0F\x00\x00\x10\x0D\x00\x00\x00\x0F\x00\x00\x16\x03\x00\x00\x02\x01\x00\x00\x00\x09\x00\x00\x04\x01\x00\x00\x00\x01',
+    _globals = (b'\xFF\xFF\xFF\x0BGDF_COLUMN_SIZE_MISMATCH',3,b'\xFF\xFF\xFF\x0BGDF_CUDA_ERROR',1,b'\xFF\xFF\xFF\x0BGDF_FLOAT32',4,b'\xFF\xFF\xFF\x0BGDF_FLOAT64',5,b'\xFF\xFF\xFF\x0BGDF_INT16',2,b'\xFF\xFF\xFF\x0BGDF_INT32',3,b'\xFF\xFF\xFF\x0BGDF_INT8',1,b'\xFF\xFF\xFF\x0BGDF_SUCCESS',0,b'\xFF\xFF\xFF\x0BGDF_UNSUPPORTED_DTYPE',2,b'\xFF\xFF\xFF\x0BGDF_invalid',0,b'\x00\x00\x07\x23gdf_add_f32',0,b'\x00\x00\x13\x23gdf_column_sizeof',0,b'\x00\x00\x0C\x23gdf_column_view',0,b'\x00\x00\x00\x23gdf_error_get_name',0,b'\x00\x00\x03\x23gdf_sin_f32',0,b'\x00\x00\x03\x23gdf_sin_generic',0),
+    _struct_unions = ((b'\x00\x00\x00\x17\x00\x00\x00\x02gdf_column_',b'\x00\x00\x0E\x11data',b'\x00\x00\x0F\x11valid',b'\x00\x00\x10\x11size',b'\x00\x00\x11\x11dtype'),),
+    _enums = (b'\x00\x00\x00\x11\x00\x00\x00\x16$gdf_dtype\x00GDF_invalid,GDF_INT8,GDF_INT16,GDF_INT32,GDF_FLOAT32,GDF_FLOAT64',b'\x00\x00\x00\x01\x00\x00\x00\x16$gdf_error\x00GDF_SUCCESS,GDF_CUDA_ERROR,GDF_UNSUPPORTED_DTYPE,GDF_COLUMN_SIZE_MISMATCH'),
+    _typenames = (b'\x00\x00\x00\x17gdf_column',b'\x00\x00\x00\x11gdf_dtype',b'\x00\x00\x00\x01gdf_error',b'\x00\x00\x00\x10gdf_index_type',b'\x00\x00\x00\x10gdf_size_type',b'\x00\x00\x00\x18gdf_valid_type'),
+)

libgdf_cffi/wrapper.py

@@ -0,0 +1,35 @@
+class GDFError(Exception):
+    def __init__(self, errcode, msg):
+        self.errcode = errcode
+        super(GDFError, self).__init__(msg)
+
+
+class _libgdf_wrapper(object):
+    def __init__(self, ffi, api):
+        self._ffi = ffi
+        self._api = api
+        self._cached = {}
+
+    def __getattr__(self, name):
+        try:
+            return self._cached[name]
+        except KeyError:
+            fn = getattr(self._api, name)
+
+            # hack to check the return type
+            textrepr = str(fn)
+            if 'gdf_error(*)' in textrepr:
+                def wrap(*args):
+                    # covert errcode to exception
+                    errcode = fn(*args)
+                    if errcode != self._api.GDF_SUCCESS:
+                        raw = self._api.gdf_error_get_name(errcode)
+                        errname = self._ffi.string(raw).decode('ascii')
+                        raise GDFError(errcode, errname)
+
+                wrap.__name__ = fn.__name__
+                self._cached[name] = wrap
+            else:
+                self._cached[name] = fn
+
+            return self._cached[name]

setup.py

@@ -0,0 +1,10 @@
+from setuptools import setup
+
+setup(name='libgdf_cffi',
+      package=["libgdf_cffi"],
+      setup_requires=["cffi>=1.0.0"],
+      cffi_modules=["libgdf_cffi/libgdf_build.py:ffibuilder"],
+      install_requires=["cffi>=1.0.0"],
+      )
+
+

src/binaryops.cu

@@ -0,0 +1,75 @@
+#include <gdf/gdf.h>
+#include <gdf/utils.h>
+#include <gdf/errorutils.h>
+
+
+template<typename T, typename F>
+__global__
+void gpu_binary_op(const T *lhs_data, const gdf_valid_type *lhs_valid,
+                   const T *rhs_data, const gdf_valid_type *rhs_valid,
+                   gdf_size_type size, T *results, F functor) {
+    int tid = threadIdx.x;
+    int blkid = blockIdx.x;
+    int blksz = blockDim.x;
+    int gridsz = gridDim.x;
+
+    int start = tid + blkid * blksz;
+    int step = blksz * gridsz;
+    if ( lhs_valid || rhs_valid ) {  // has valid mask
+        for (int i=start; i<size; i+=step) {
+            if (gdf_is_valid(lhs_valid, i) && gdf_is_valid(rhs_valid, i))
+                results[i] = functor.apply(lhs_data[i], rhs_data[i]);
+        }
+    } else {                         // no valid mask
+        for (int i=start; i<size; i+=step) {
+            results[i] = functor.apply(lhs_data[i], rhs_data[i]);
+        }
+    }
+}
+
+template<typename T, typename F>
+struct BinaryOp {
+    static
+    gdf_error launch(gdf_column *lhs, gdf_column *rhs, gdf_column *output) {
+        if (lhs->size != rhs->size || lhs->size != output->size) {
+            return GDF_COLUMN_SIZE_MISMATCH;
+        }
+
+        // find optimal blocksize
+        int mingridsize, blocksize;
+        CUDA_TRY(
+            cudaOccupancyMaxPotentialBlockSize(&mingridsize, &blocksize,
+                                               gpu_binary_op<T, F>)
+        );
+        // find needed gridsize
+        int gridsize = (lhs->size + blocksize - 1) / blocksize;
+
+        F functor;
+        gpu_binary_op<<<gridsize, blocksize>>>(
+            // inputs
+            (const T*)lhs->data, lhs->valid,
+            (const T*)rhs->data, rhs->valid,
+            lhs->size,
+            // output
+            (T*)output->data,
+            // action
+            functor
+        );
+
+        CUDA_CHECK_LAST();
+        return GDF_SUCCESS;
+    }
+};
+
+
+template<typename T>
+struct DeviceAdd {
+    __device__
+    T apply(T lhs, T rhs) {
+        return lhs + rhs;
+    }
+};
+
+gdf_error gdf_add_f32(gdf_column *lhs, gdf_column *rhs, gdf_column *output) {
+    return BinaryOp<float, DeviceAdd<float> >::launch(lhs, rhs, output);
+}

src/column.cpp

@@ -0,0 +1,16 @@
+#include <gdf/gdf.h>
+
+
+gdf_size_type gdf_column_sizeof() {
+    return sizeof(gdf_column);
+}
+
+gdf_error gdf_column_view(gdf_column *column, void *data, gdf_valid_type *valid,
+                    gdf_size_type size, gdf_dtype dtype) {
+    column->data = data;
+    column->valid = valid;
+    column->size = size;
+    column->dtype = dtype;
+    return GDF_SUCCESS;
+}
+

src/errorhandling.cpp

@@ -0,0 +1,11 @@
+#include <gdf/gdf.h>
+
+#define GETNAME(x) case x: return #x;
+const char * gdf_error_get_name(gdf_error errcode) {
+    switch (errcode) {
+    GETNAME(GDF_SUCCESS)
+    GETNAME(GDF_CUDA_ERROR)
+    GETNAME(GDF_UNSUPPORTED_DTYPE)
+    GETNAME(GDF_COLUMN_SIZE_MISMATCH)
+    }
+}

src/unaryops.cu

@@ -0,0 +1,84 @@
+#include <math.h>
+
+#include <gdf/gdf.h>
+#include <gdf/utils.h>
+#include <gdf/errorutils.h>
+
+template<typename T, typename F>
+__global__
+void gpu_unary_op(const T *data, const gdf_valid_type *valid,
+                    gdf_size_type size, T *results, F functor) {
+    int tid = threadIdx.x;
+    int blkid = blockIdx.x;
+    int blksz = blockDim.x;
+    int gridsz = gridDim.x;
+
+    int start = tid + blkid * blksz;
+    int step = blksz * gridsz;
+    if ( valid ) {  // has valid mask
+        for (int i=start; i<size; i+=step) {
+            if ( gdf_is_valid(valid, i) )
+                results[i] = functor.apply(data[i]);
+        }
+    } else {        // no valid mask
+        for (int i=start; i<size; i+=step) {
+            results[i] = functor.apply(data[i]);
+        }
+    }
+}
+
+template<typename T, typename F>
+struct UnaryOp {
+    static
+    gdf_error launch(gdf_column *input, gdf_column *output) {
+        /* check for size of the columns */
+        if (input->size != output->size) {
+            return GDF_COLUMN_SIZE_MISMATCH;
+        }
+
+        // find optimal blocksize
+        int mingridsize, blocksize;
+        CUDA_TRY(
+            cudaOccupancyMaxPotentialBlockSize(&mingridsize, &blocksize,
+                                               gpu_unary_op<T, F>)
+        );
+        // find needed gridsize
+        int gridsize = (input->size + blocksize - 1) / blocksize;
+
+        F functor;
+        gpu_unary_op<<<gridsize, blocksize>>>(
+            // input
+            (const T*)input->data, input->valid, input->size,
+            // output
+            (T*)output->data,
+            // action
+            functor
+        );
+
+        CUDA_CHECK_LAST();
+        return GDF_SUCCESS;
+    }
+};
+
+
+template<typename T>
+struct DeviceSin {
+    __device__
+    T apply(T data) {
+        return sin(data);
+    }
+};
+
+gdf_error gdf_sin_generic(gdf_column *input, gdf_column *output) {
+    switch ( input->dtype ) {
+    case GDF_FLOAT32:
+        return gdf_sin_f32(input, output);
+    default:
+        return GDF_UNSUPPORTED_DTYPE;
+    }
+}
+
+
+gdf_error gdf_sin_f32(gdf_column *input, gdf_column *output) {
+    return UnaryOp<float, DeviceSin<float> >::launch(input, output);
+}