You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@singa.apache.org by wa...@apache.org on 2016/08/31 04:37:01 UTC
[1/4] incubator-singa git commit: SINGA-243 ViennaCL backend for
OpenCL support
Repository: incubator-singa
Updated Branches:
refs/heads/master 61ec2ec74 -> 1996e8a95
http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/595302a3/test/singa/test_opencl.cc
----------------------------------------------------------------------
diff --git a/test/singa/test_opencl.cc b/test/singa/test_opencl.cc
index 3ce1889..d464bc9 100644
--- a/test/singa/test_opencl.cc
+++ b/test/singa/test_opencl.cc
@@ -23,419 +23,286 @@
#include "singa/core/device.h"
#include "singa/core/tensor.h"
#include "singa/proto/core.pb.h"
+
using singa::CppCPU;
using singa::Block;
using singa::Shape;
using singa::Tensor;
#ifdef USE_OPENCL
+
using singa::OpenclDevice;
+
class OpenCL_TensorMath : public ::testing::Test {
protected:
OpenCL_TensorMath() {
- for (int i = 0; i < 4; i++) {
- float4[i] = (float)i;
- float4zero[i] = 0.0f;
- }
-
- for (int i = 0; i < 16; i++) {
- float16[i] = (float)i;
- float16zero[i] = 0.0f;
- }
-
auto ocl_dev = std::make_shared<OpenclDevice>();
-
- tf4in = Tensor(Shape{1, 4}, ocl_dev);
- tf4in.CopyDataFromHostPtr(float4, 4);
-
- tf4zin = Tensor(Shape{1, 4}, ocl_dev);
- tf4zin.CopyDataFromHostPtr(float4zero, 4);
-
- tf16in = Tensor(Shape{4, 4}, ocl_dev);
- tf16in.CopyDataFromHostPtr(float16, 16);
-
- tf16zin = Tensor(Shape{4, 4}, ocl_dev);
- tf16zin.CopyDataFromHostPtr(float16zero, 16);
-
- float empty[10000] = {};
+
+ a = Tensor(Shape{6}, ocl_dev);
+ b = Tensor(Shape{6}, ocl_dev);
+ c = Tensor(Shape{6, 1}, ocl_dev);
+ d = Tensor(Shape{3, 2}, ocl_dev);
+ e = Tensor(Shape{3, 2}, ocl_dev);
empty10k = Tensor(Shape{10000}, ocl_dev);
- empty10k.CopyDataFromHostPtr(empty, 10000);
+
+ a.CopyDataFromHostPtr<float>(dat1, 6);
+ b.CopyDataFromHostPtr<float>(dat2, 6);
+ e.CopyDataFromHostPtr<float>(dat1, 6);
}
-
- float float4[4];
- float float4zero[4];
- float float16[16];
- float float16zero[16];
-
- Tensor tf4in, tf16in;
- Tensor tf4zin, tf16zin;
+
+ Tensor a, b, c, d, e;
Tensor empty10k;
+ const float dat1[6] = {1.0f, 2.0f, 3.0f, 4.0f, 5.0f, 6.0f};
+ const float dat2[6] = {1.1f, 2.1f, 3.1f, 4.1f, 5.1f, 6.1f};
};
-// Makes a float array and fills it with increasing values from 0.
-float* MakeMatrix(const int size) {
- float* mat = new float[size];
- for (int i = 0; i < size; i++)
- mat[i] = i;
- return mat;
-}
-
-
-TEST(OpenclDevice, Constructor) {
- OpenclDevice dev;
- EXPECT_EQ(0, dev.id());
-}
-
-
-TEST(OpenclDevice, MemoryAllocFree) {
- OpenclDevice dev;
- Block* b = dev.NewBlock(4);
- EXPECT_NE(nullptr, b);
- EXPECT_EQ(4u, b->size());
- dev.FreeBlock(b);
-}
-
-// Tests for integrity of one round of data transfer to an OpenCL device and back.
-TEST(OpenclDevice, CopyDataToFrom) {
- OpenclDevice dev;
- CppCPU host;
-
- Block* a = host.NewBlock(4);
- Block* b = dev.NewBlock(4);
- Block* c = host.NewBlock(4);
-
- // Allocate the Block object on the host.
- char s[] = {'a', 'b', 'c', 'x'};
- host.CopyDataFromHostPtr(a, s, 4);
-
- // Copy back and forth.
- dev.CopyDataToFrom(b, a, 4, singa::kHostToDevice);
- dev.CopyDataToFrom(c, b, 4, singa::kDeviceToHost);
-
- const char* astr = static_cast<const char*>(c->data());
- EXPECT_EQ('a', astr[0]);
- EXPECT_EQ('b', astr[1]);
- EXPECT_EQ('x', astr[3]);
-}
-
-
-TEST(OpenclDevice, DuplicateDataOnDevice) {
- OpenclDevice dev;
- CppCPU host;
-
- Block* a = host.NewBlock(4);
- Block* b = dev.NewBlock(4);
- Block* c = dev.NewBlock(4);
- Block* d = host.NewBlock(4);
-
- // Allocate the Block object on the host.
- char s[] = {'a', 'b', 'c', 'x'};
- host.CopyDataFromHostPtr(a, s, 4);
-
- // Copy to device and duplicate.
- dev.CopyDataToFrom(b, a, 4, singa::kHostToDevice);
- dev.CopyDataToFrom(c, b, 4, singa::kDeviceToDevice);
- dev.CopyDataToFrom(d, c, 4, singa::kDeviceToHost);
-
- const char* astr = static_cast<const char*>(d->data());
- EXPECT_EQ('a', astr[0]);
- EXPECT_EQ('b', astr[1]);
- EXPECT_EQ('x', astr[3]);
-}
-
-// Tensor tests, uses OpenCL_TensorMath class defined above.
-
-TEST_F(OpenCL_TensorMath, CopyDataToDevice) {
- tf4in.ToHost();
- const float* out = tf4in.data<float>();
-
- EXPECT_EQ(1.0f, out[1]);
- EXPECT_EQ(3.0f, out[3]);
-}
-
-
TEST_F(OpenCL_TensorMath, MemberAbs) {
- tf4in = Abs(tf4in);
-
- tf4in.ToHost();
- const float* out = tf4in.data<float>();
-
- EXPECT_EQ(0.0f, out[0]);
- EXPECT_EQ(1.0f, out[1]);
- EXPECT_EQ(2.0f, out[2]);
- EXPECT_EQ(3.0f, out[3]);
+ Tensor aa = a.Clone();
+ Tensor bb = b.Clone();
+ Tensor cc = aa - bb;
+
+ cc.ToHost();
+ const float *dptr = cc.data<float>();
+ EXPECT_NEAR(-0.1, dptr[0], 1e-5);
+ EXPECT_NEAR(-0.1, dptr[1], 1e-5);
+ EXPECT_NEAR(-0.1, dptr[2], 1e-5);
+
+ Tensor p = Abs(cc);
+ const float *dptr1 = p.data<float>();
+ EXPECT_NEAR(0.1, dptr1[0], 1e-5);
+ EXPECT_NEAR(0.1, dptr1[1], 1e-5);
+ EXPECT_NEAR(0.1, dptr1[2], 1e-5);
}
-TEST_F(OpenCL_TensorMath, MemberExp) {
- tf4in = Exp(tf4in);
+//TEST_F(OpenCL_TensorMath, MemberClamp) { }
- tf4in.ToHost();
- const float* out = tf4in.data<float>();
- EXPECT_NEAR(exp(0.0f), out[0], 1e-5);
- EXPECT_NEAR(exp(1.0f), out[1], 1e-5);
- EXPECT_NEAR(exp(2.0f), out[2], 1e-5);
- EXPECT_NEAR(exp(3.0f), out[3], 1e-5);
+TEST_F(OpenCL_TensorMath, MemberExp) {
+ Tensor p = Exp(a);
+
+ p.ToHost();
+ const float *dptr1 = p.data<float>();
+ EXPECT_NEAR(exp(1.0f), dptr1[0], 1e-5);
+ EXPECT_NEAR(exp(2.0f), dptr1[1], 1e-5);
+ EXPECT_NEAR(exp(3.0f), dptr1[2], 1e-5);
}
TEST_F(OpenCL_TensorMath, MemberLog) {
- tf4in = Log(tf4in);
-
- tf4in.ToHost();
- const float* out = tf4in.data<float>();
-
-// EXPECT_NEAR(log(0.0f), out[0], 1e-5); // Evaluates to neg infinity.
- EXPECT_NEAR(log(1.0f), out[1], 1e-5);
- EXPECT_NEAR(log(2.0f), out[2], 1e-5);
- EXPECT_NEAR(log(3.0f), out[3], 1e-5);
+ Tensor p = Log(a);
+
+ p.ToHost();
+ const float *dptr1 = p.data<float>();
+ EXPECT_NEAR(log(1.0f), dptr1[0], 1e-5);
+ EXPECT_NEAR(log(2.0f), dptr1[1], 1e-5);
+ EXPECT_NEAR(log(3.0f), dptr1[2], 1e-5);
}
TEST_F(OpenCL_TensorMath, MemberReLU) {
- tf4in -= 2.0f;
- Tensor result = ReLU(tf4in);
-
- result.ToHost();
- const float* out = result.data<float>();
-
- EXPECT_NEAR(0.0f, out[0], 1e-5);
- EXPECT_NEAR(0.0f, out[1], 1e-5);
- EXPECT_NEAR(0.0f, out[2], 1e-5);
- EXPECT_NEAR(1.0f, out[3], 1e-5);
+ Tensor aa = a.Clone();
+ Tensor cc = aa - 2.0f;
+
+ cc.ToHost();
+ const float *dptr = cc.data<float>();
+ EXPECT_NEAR(-1.0f, dptr[0], 1e-5);
+ EXPECT_NEAR(0.0f, dptr[1], 1e-5);
+ EXPECT_NEAR(1.0f, dptr[2], 1e-5);
+
+ Tensor p = ReLU(cc);
+ p.ToHost();
+ const float *dptr1 = p.data<float>();
+ EXPECT_NEAR(0.0f, dptr1[0], 1e-5);
+ EXPECT_NEAR(0.0f, dptr1[1], 1e-5);
+ EXPECT_NEAR(1.0f, dptr1[2], 1e-5);
}
TEST_F(OpenCL_TensorMath, MemberSigmoid) {
- tf4in = Sigmoid(tf4in);
-
- tf4in.ToHost();
- const float* out = tf4in.data<float>();
-
- EXPECT_NEAR(1.0f / (1.0f + exp(-0.0f)), out[0], 1e-5);
- EXPECT_NEAR(1.0f / (1.0f + exp(-1.0f)), out[1], 1e-5);
- EXPECT_NEAR(1.0f / (1.0f + exp(-2.0f)), out[2], 1e-5);
+ Tensor p = Sigmoid(a);
+ p.ToHost();
+ const float *dptr1 = p.data<float>();
+ EXPECT_NEAR(1.0f / (1.0f + exp(-1.0f)), dptr1[0], 1e-5);
+ EXPECT_NEAR(1.0f / (1.0f + exp(-2.0f)), dptr1[1], 1e-5);
+ EXPECT_NEAR(1.0f / (1.0f + exp(-3.0f)), dptr1[2], 1e-5);
}
TEST_F(OpenCL_TensorMath, MemberSign) {
- tf4in -= 1.0f;
-
- tf4in.ToHost();
- const float* out = tf4in.data<float>();
-
- EXPECT_NEAR(-1.0f, out[0], 1e-5);
- EXPECT_NEAR(0.0f, out[1], 1e-5);
- EXPECT_NEAR(1.0f, out[2], 1e-5);
- EXPECT_NEAR(2.0f, out[3], 1e-5);
+ Tensor aa = a.Clone();
+ Tensor cc = aa - 2.0f;
+ cc.ToHost();
+ const float *dptr = cc.data<float>();
+ EXPECT_NEAR(-1.0f, dptr[0], 1e-5);
+ EXPECT_NEAR(0.0f, dptr[1], 1e-5);
+ EXPECT_NEAR(1.0f, dptr[2], 1e-5);
+
+ Tensor p = Sign(cc);
+ p.ToHost();
+ const float *dptr1 = p.data<float>();
+ EXPECT_EQ(0.0f, dptr1[0]);
+ EXPECT_EQ(0.0f, dptr1[1]);
+ EXPECT_EQ(1.0f, dptr1[2]);
}
TEST_F(OpenCL_TensorMath, MemberSqrt) {
- tf4in = Sqrt(tf4in);
-
- tf4in.ToHost();
- const float* out = tf4in.data<float>();
-
- EXPECT_NEAR(0.0f, out[0], 1e-5);
- EXPECT_NEAR(1.0f, out[1], 1e-5);
- EXPECT_NEAR(sqrt(2.0f), out[2], 1e-5);
- EXPECT_NEAR(sqrt(3.0f), out[3], 1e-5);
+ Tensor p = Sqrt(a);
+ p.ToHost();
+ const float *dptr1 = p.data<float>();
+ EXPECT_NEAR(sqrt(1.0), dptr1[0], 1e-5);
+ EXPECT_NEAR(sqrt(2.0), dptr1[1], 1e-5);
+ EXPECT_NEAR(sqrt(3.0), dptr1[2], 1e-5);
}
TEST_F(OpenCL_TensorMath, MemberSquare) {
- tf4in = Square(tf4in);
-
- tf4in.ToHost();
- const float* out = tf4in.data<float>();
-
- EXPECT_NEAR(0.0f, out[0], 1e-5);
- EXPECT_NEAR(1.0f, out[1], 1e-5);
- EXPECT_NEAR(4.0f, out[2], 1e-5);
- EXPECT_NEAR(9.0f, out[3], 1e-5);
+ Tensor p = Square(a);
+ p.ToHost();
+ const float *dptr1 = p.data<float>();
+ EXPECT_NEAR(1.0, dptr1[0], 1e-5);
+ EXPECT_NEAR(4.0, dptr1[1], 1e-5);
+ EXPECT_NEAR(9.0, dptr1[2], 1e-5);
}
TEST_F(OpenCL_TensorMath, MemberTanh) {
- tf4in = Tanh(tf4in);
-
- tf4in.ToHost();
- const float* out = tf4in.data<float>();
-
- EXPECT_NEAR(0.0f, out[0], 1e-5);
- EXPECT_NEAR(tanh(1.0f), out[1], 1e-5);
- EXPECT_NEAR(tanh(2.0f), out[2], 1e-5);
- EXPECT_NEAR(tanh(3.0f), out[3], 1e-5);
+ Tensor p = Tanh(a);
+ p.ToHost();
+ const float *dptr1 = p.data<float>();
+ EXPECT_NEAR(tanh(1.0), dptr1[0], 1e-5);
+ EXPECT_NEAR(tanh(2.0), dptr1[1], 1e-5);
+ EXPECT_NEAR(tanh(3.0), dptr1[2], 1e-5);
}
-
TEST_F(OpenCL_TensorMath, Sum) {
- Tensor result = Sum(tf4in, 0);
-
- result.ToHost();
- const float* out = result.data<float>();
+ float result = Sum(a);
+ EXPECT_EQ(21.0f, result);
+}
- EXPECT_NEAR(0.0f, out[0], 1e-5);
- EXPECT_NEAR(1.0f, out[1], 1e-5);
- EXPECT_NEAR(2.0f, out[2], 1e-5);
- EXPECT_NEAR(3.0f, out[3], 1e-5);
+/*
+TEST_F(OpenCL_TensorMath, SoftMax) {
+ Tensor p1 = SoftMax(Reshape(e, Shape{1, 6}));
+ p1.ToHost();
+ const float *dptr1 = p1.data<float>();
+ float sum = 0;
+ for (int i = 0; i < 6; i++) sum += exp(i + 1);
+ EXPECT_NEAR(exp(1) / sum, dptr1[0], 1e-5);
+ EXPECT_NEAR(exp(3) / sum, dptr1[2], 1e-5);
+ EXPECT_NEAR(exp(5) / sum, dptr1[4], 1e-5);
+ EXPECT_NEAR(exp(2) / sum, dptr1[1], 1e-5);
+ EXPECT_NEAR(exp(4) / sum, dptr1[3], 1e-5);
+ EXPECT_NEAR(exp(6) / sum, dptr1[5], 1e-5);
+
+ Tensor p2 = SoftMax(e);
+ p2.ToHost();
+ const float *dptr2 = p2.data<float>();
+ EXPECT_NEAR(exp(1) / (exp(1) + exp(2)), dptr2[0], 1e-5);
+ EXPECT_NEAR(exp(2) / (exp(1) + exp(2)), dptr2[1], 1e-5);
}
+*/
TEST_F(OpenCL_TensorMath, MemberLT) {
- Tensor result = tf4in < 2.0f;
-
- result.ToHost();
- const float* out = result.data<float>();
-
- EXPECT_FLOAT_EQ(1.0f, out[0]);
- EXPECT_FLOAT_EQ(1.0f, out[1]);
- EXPECT_FLOAT_EQ(0.0f, out[2]);
- EXPECT_FLOAT_EQ(0.0f, out[3]);
+ Tensor p1 = a < 2.0f;
+ p1.ToHost();
+ const float *dptr1 = p1.data<float>();
+ EXPECT_FLOAT_EQ(1.0f, dptr1[0]);
+ EXPECT_FLOAT_EQ(0.0f, dptr1[1]);
+ EXPECT_FLOAT_EQ(0.0f, dptr1[2]);
}
TEST_F(OpenCL_TensorMath, MemberLE) {
- Tensor result = tf4in <= 2.0f;
-
- result.ToHost();
- const float* out = result.data<float>();
-
- EXPECT_FLOAT_EQ(1.0f, out[0]);
- EXPECT_FLOAT_EQ(1.0f, out[1]);
- EXPECT_FLOAT_EQ(1.0f, out[2]);
- EXPECT_FLOAT_EQ(0.0f, out[3]);
+ Tensor p1 = a <= 2.0f;
+ p1.ToHost();
+ const float *dptr1 = p1.data<float>();
+ EXPECT_FLOAT_EQ(1.0f, dptr1[0]);
+ EXPECT_FLOAT_EQ(1.0f, dptr1[1]);
+ EXPECT_FLOAT_EQ(0.0f, dptr1[2]);
}
TEST_F(OpenCL_TensorMath, MemberGT) {
- Tensor result = tf4in > 2.0f;
-
- result.ToHost();
- const float* out = result.data<float>();
-
- EXPECT_FLOAT_EQ(0.0f, out[0]);
- EXPECT_FLOAT_EQ(0.0f, out[1]);
- EXPECT_FLOAT_EQ(0.0f, out[2]);
- EXPECT_FLOAT_EQ(1.0f, out[3]);
+ Tensor p1 = a > 2.0f;
+ p1.ToHost();
+ const float *dptr1 = p1.data<float>();
+ EXPECT_FLOAT_EQ(0.0f, dptr1[0]);
+ EXPECT_FLOAT_EQ(0.0f, dptr1[1]);
+ EXPECT_FLOAT_EQ(1.0f, dptr1[2]);
}
TEST_F(OpenCL_TensorMath, MemberGE) {
- Tensor result = tf4in >= 2.0f;
-
- result.ToHost();
- const float* out = result.data<float>();
-
- EXPECT_FLOAT_EQ(0.0f, out[0]);
- EXPECT_FLOAT_EQ(0.0f, out[1]);
- EXPECT_FLOAT_EQ(1.0f, out[2]);
- EXPECT_FLOAT_EQ(1.0f, out[3]);
+ Tensor p1 = a >= 2.0f;
+ p1.ToHost();
+ const float *dptr1 = p1.data<float>();
+ EXPECT_FLOAT_EQ(0.0f, dptr1[0]);
+ EXPECT_FLOAT_EQ(1.0f, dptr1[1]);
+ EXPECT_FLOAT_EQ(1.0f, dptr1[2]);
}
TEST_F(OpenCL_TensorMath, MemberPow) {
- Tensor result = Pow(tf4in, 2.0f);
-
- result.ToHost();
- const float* out = result.data<float>();
-
- EXPECT_FLOAT_EQ(0.0f, out[0]);
- EXPECT_FLOAT_EQ(1.0f, out[1]);
- EXPECT_FLOAT_EQ(4.0f, out[2]);
- EXPECT_FLOAT_EQ(9.0f, out[3]);
-
- result = Pow(tf4in, tf4in);
-
- result.ToHost();
- const float* out1 = result.data<float>();
-
- EXPECT_FLOAT_EQ(1.0f, out1[0]); // 0 ^ 0 is 1, apparently.
- EXPECT_FLOAT_EQ(1.0f, out1[1]);
- EXPECT_FLOAT_EQ(4.0f, out1[2]);
- EXPECT_FLOAT_EQ(27.0f, out1[3]);
+ Tensor p1 = Pow(b, 3.0f);
+ p1.ToHost();
+ const float *dptr1 = p1.data<float>();
+ EXPECT_FLOAT_EQ(pow(1.1f, 3.0f), dptr1[0]);
+ EXPECT_FLOAT_EQ(pow(2.1f, 3.0f), dptr1[1]);
+ EXPECT_FLOAT_EQ(pow(3.1f, 3.0f), dptr1[2]);
+
+ Tensor p2 = Pow(a,b);
+ p2.ToHost();
+ const float *dptr2 = p2.data<float>();
+ EXPECT_FLOAT_EQ(pow(1.0f,1.1f), dptr2[0]);
+ EXPECT_FLOAT_EQ(pow(2.0f,2.1f), dptr2[1]);
+ EXPECT_FLOAT_EQ(pow(3.0f,3.1f), dptr2[2]);
}
TEST_F(OpenCL_TensorMath, MemberSub) {
- Tensor result = tf4in - tf4zin;
-
- result.ToHost();
- const float* out = result.data<float>();
-
- EXPECT_FLOAT_EQ(0.0f, out[0]);
- EXPECT_FLOAT_EQ(1.0f, out[1]);
- EXPECT_FLOAT_EQ(2.0f, out[2]);
- EXPECT_FLOAT_EQ(3.0f, out[3]);
-
- result = tf4in - 0.0f;
-
- result.ToHost();
- const float* out1 = result.data<float>();
-
- EXPECT_FLOAT_EQ(0.0f, out1[0]);
- EXPECT_FLOAT_EQ(1.0f, out1[1]);
- EXPECT_FLOAT_EQ(2.0f, out1[2]);
- EXPECT_FLOAT_EQ(3.0f, out1[3]);
+ Tensor p1 = a - b;
+ p1.ToHost();
+ const float *dptr1 = p1.data<float>();
+ EXPECT_NEAR(-0.1, dptr1[0], 1e-5);
+ EXPECT_NEAR(-0.1, dptr1[1], 1e-5);
+ EXPECT_NEAR(-0.1, dptr1[2], 1e-5);
}
TEST_F(OpenCL_TensorMath, MemberEltwiseMult) {
- Tensor result = tf4in * tf4zin;
-
- result.ToHost();
- const float* out = result.data<float>();
-
- EXPECT_FLOAT_EQ(0.0f, out[0]);
- EXPECT_FLOAT_EQ(0.0f, out[1]);
- EXPECT_FLOAT_EQ(0.0f, out[2]);
- EXPECT_FLOAT_EQ(0.0f, out[3]);
-
- result = tf4in * 10.0f;
-
- result.ToHost();
- const float* out1 = result.data<float>();
-
- EXPECT_FLOAT_EQ(0.0f, out1[0]);
- EXPECT_FLOAT_EQ(10.0f, out1[1]);
- EXPECT_FLOAT_EQ(20.0f, out1[2]);
- EXPECT_FLOAT_EQ(30.0f, out1[3]);
+ Tensor p1 = a * b;
+ p1.ToHost();
+ const float *dptr1 = p1.data<float>();
+ EXPECT_NEAR(1.0 * 1.1, dptr1[0], 1e-5);
+ EXPECT_NEAR(2.0 * 2.1, dptr1[1], 1e-5);
+ EXPECT_NEAR(3.0 * 3.1, dptr1[2], 1e-5);
}
TEST_F(OpenCL_TensorMath, MemberDiv) {
- Tensor result = tf4in / tf4in;
-
- result.ToHost();
- const float* out = result.data<float>();
-
-// EXPECT_FLOAT_EQ(0.0f, out[0]); // Divide by zero.
- EXPECT_FLOAT_EQ(1.0f, out[1]);
- EXPECT_FLOAT_EQ(1.0f, out[2]);
- EXPECT_FLOAT_EQ(1.0f, out[3]);
-
- result = tf4in / 10.0f;
-
- result.ToHost();
- const float* out1 = result.data<float>();
-
- EXPECT_FLOAT_EQ(0.0f, out1[0]);
- EXPECT_FLOAT_EQ(0.1f, out1[1]);
- EXPECT_FLOAT_EQ(0.2f, out1[2]);
- EXPECT_FLOAT_EQ(0.3f, out1[3]);
-
- result = Div(10.0f, tf4in);
-
- result.ToHost();
- const float* out2 = result.data<float>();
-
-// EXPECT_FLOAT_EQ(0.0f, out[0]); // Divide by 0.
- EXPECT_FLOAT_EQ(10.0f, out2[1]);
- EXPECT_FLOAT_EQ(5.0f, out2[2]);
- EXPECT_NEAR((10.0f / 3.0f), out2[3], 1e-5);
+ Tensor p1 = a / b;
+ p1.ToHost();
+ const float *dptr1 = p1.data<float>();
+ EXPECT_NEAR(1.0 / 1.1, dptr1[0], 1e-5);
+ EXPECT_NEAR(2.0 / 2.1, dptr1[1], 1e-5);
+ EXPECT_NEAR(3.0 / 3.1, dptr1[2], 1e-5);
+
+ Tensor p2 = Div(10.0f, b);
+ p2.ToHost();
+ const float *dptr2 = p2.data<float>();
+ EXPECT_NEAR(10.0 / 1.1, dptr2[0], 1e-5);
+ EXPECT_NEAR(10.0 / 2.1, dptr2[1], 1e-5);
+ EXPECT_NEAR(10.0 / 3.1, dptr2[2], 1e-5);
+
+ Tensor p3 = a / 8.0f;
+ p3.ToHost();
+ const float *dptr3 = p3.data<float>();
+ EXPECT_NEAR(1.0 / 8.0, dptr3[0], 1e-5);
+ EXPECT_NEAR(2.0 / 8.0, dptr3[1], 1e-5);
+ EXPECT_NEAR(3.0 / 8.0, dptr3[2], 1e-5);
}
// **************************************
@@ -444,63 +311,49 @@ TEST_F(OpenCL_TensorMath, MemberDiv) {
TEST_F(OpenCL_TensorMath, Bernoulli) {
const float p = 0.3f;
-
Bernoulli(p, &empty10k);
-
empty10k.ToHost();
const float* out = empty10k.data<float>();
-
float sum = 0.0f;
for (int i = 0; i < 10000; i++) sum += out[i];
-
float mean = sum / 10000;
-
EXPECT_NEAR(mean, p, 1e-2);
sum = 0.0f;
for (int i = 0; i < 10000; i++) sum += (out[i] - mean) * (out[i] - mean);
float variance = sum / 9999;
-
EXPECT_NEAR(variance, p * (1 - p), 1e-2);
}
TEST_F(OpenCL_TensorMath, Gaussian) {
Gaussian(0.0f, 1.0f, &empty10k);
-
empty10k.ToHost();
const float* out = empty10k.data<float>();
-
float sum = 0.0f;
for (int i = 0; i < 10000; i++) sum += out[i];
float mean = sum / 10000;
-
EXPECT_NEAR(mean, 0.0f, 1e-2);
sum = 0.0f;
for (int i = 0; i < 10000; i++) sum += (out[i] - mean) * (out[i] - mean);
float variance = sum / 9999;
-
EXPECT_NEAR(variance, 1.0f, 1e-2);
}
TEST_F(OpenCL_TensorMath, Uniform) {
Uniform(0.1f, 0.2f, &empty10k);
-
empty10k.ToHost();
const float* out = empty10k.data<float>();
-
float sum = 0.0f;
for (int i = 0; i < 10000; i++) sum += out[i];
float mean = sum / 10000;
-
EXPECT_NEAR(mean, 0.15f, 1e-2);
sum = 0.0f;
for (int i = 0; i < 10000; i++) sum += (out[i] - mean) * (out[i] - mean);
float variance = sum / 9999;
-
EXPECT_NEAR(variance, 0.01f, 1e-2);
}
@@ -509,121 +362,96 @@ TEST_F(OpenCL_TensorMath, Uniform) {
// *********************************************************
-TEST_F(OpenCL_TensorMath, EltwiseAdd) {
- Tensor result = tf4in + tf4in;
-
- result.ToHost();
- const float* out = result.data<float>();
-
- EXPECT_EQ(0.0f, out[0]);
- EXPECT_EQ(2.0f, out[1]);
- EXPECT_EQ(4.0f, out[2]);
- EXPECT_EQ(6.0f, out[3]);
-
- result = tf4in + tf4zin;
-
- result.ToHost();
- const float* out1 = result.data<float>();
-
- EXPECT_EQ(0.0f, out1[0]);
- EXPECT_EQ(1.0f, out1[1]);
- EXPECT_EQ(2.0f, out1[2]);
- EXPECT_EQ(3.0f, out1[3]);
-
- result = Tensor(tf4in.shape(), tf4in.device(), tf4in.data_type());
- Add(tf4in, tf4in, &result);
-
- result.ToHost();
- const float* out2 = result.data<float>();
-
- EXPECT_EQ(0.0f, out2[0]);
- EXPECT_EQ(2.0f, out2[1]);
- EXPECT_EQ(4.0f, out2[2]);
- EXPECT_EQ(6.0f, out2[3]);
-
- result = tf4in + 1.0f;
+TEST_F(OpenCL_TensorMath, MemberAddTensor) {
+ Tensor aa = a.Clone();
+ aa += a;
+ aa.ToHost();
+ const float *dptr = aa.data<float>();
+ EXPECT_FLOAT_EQ(2.0f, dptr[0]);
+ EXPECT_FLOAT_EQ(4.0f, dptr[1]);
+ EXPECT_FLOAT_EQ(6.0f, dptr[2]);
+
+ // check p is initialized to 0
+ Tensor p(Shape{6});
+ p += aa;
+ p.ToHost();
+ const float *dptr1 = p.data<float>();
+ EXPECT_FLOAT_EQ(2.0f, dptr1[0]);
+ EXPECT_FLOAT_EQ(4.0f, dptr1[1]);
+ EXPECT_FLOAT_EQ(6.0f, dptr1[2]);
+
+ a += b;
+ a.ToHost();
+ const float *dptr2 = a.data<float>();
+ EXPECT_FLOAT_EQ(2.1f, dptr2[0]);
+ EXPECT_FLOAT_EQ(4.1f, dptr2[1]);
+ EXPECT_FLOAT_EQ(6.1f, dptr2[2]);
+ EXPECT_FLOAT_EQ(12.1f, dptr2[5]);
+}
- result.ToHost();
- const float* out3 = result.data<float>();
- EXPECT_EQ(1.0f, out3[0]);
- EXPECT_EQ(2.0f, out3[1]);
- EXPECT_EQ(3.0f, out3[2]);
- EXPECT_EQ(4.0f, out3[3]);
+TEST_F(OpenCL_TensorMath, AddTensors) {
+ Tensor ret(a.shape(), a.device(), a.data_type());
+ Add(a, b, &ret);
+ ret.ToHost();
+ const float *dptr = ret.data<float>();
+ EXPECT_FLOAT_EQ(2.1f, dptr[0]);
+ EXPECT_FLOAT_EQ(4.1f, dptr[1]);
+ EXPECT_FLOAT_EQ(6.1f, dptr[2]);
+ EXPECT_FLOAT_EQ(12.1f, dptr[5]);
+
+ Tensor d = a + b;
+ d.ToHost();
+ const float *dptr2 = d.data<float>();
+ EXPECT_FLOAT_EQ(2.1f, dptr2[0]);
+ EXPECT_FLOAT_EQ(4.1f, dptr2[1]);
+ EXPECT_FLOAT_EQ(6.1f, dptr2[2]);
+ EXPECT_FLOAT_EQ(12.1f, dptr2[5]);
+
+ Add(a, b, &a);
+ a.ToHost();
+ const float *dptr1 = a.data<float>();
+ EXPECT_FLOAT_EQ(2.1f, dptr1[0]);
+ EXPECT_FLOAT_EQ(4.1f, dptr1[1]);
+ EXPECT_FLOAT_EQ(6.1f, dptr1[2]);
+ EXPECT_FLOAT_EQ(12.1f, dptr1[5]);
}
TEST_F(OpenCL_TensorMath, SetValue) {
- const float one_third = 1.0f / 3.0f;
- empty10k.SetValue(one_third);
-
- empty10k.ToHost();
- const float* out = empty10k.data<float>();
-
- EXPECT_EQ(one_third, out[0]);
- EXPECT_EQ(one_third, out[1]);
- EXPECT_EQ(one_third, out[1024]);
- EXPECT_EQ(one_third, out[4096]);
- EXPECT_EQ(one_third, out[9998]);
- EXPECT_EQ(one_third, out[9999]);
+ Tensor t(Shape{4});
+ t.SetValue(0.3f);
+ t.ToHost();
+ const float *ptr = t.data<float>();
+ for (int i = 0; i < 4; i++) EXPECT_FLOAT_EQ(ptr[i], 0.3f);
}
TEST_F(OpenCL_TensorMath, Axpy) {
- Axpy(10.0f, tf4in, &tf4in);
-
- tf4in.ToHost();
- const float* out = tf4in.data<float>();
-
- EXPECT_EQ(0.0f, out[0]); // 0 * 10 + 0 = 0
- EXPECT_EQ(11.0f, out[1]); // 1 * 10 + 1 = 11
- EXPECT_EQ(22.0f, out[2]); // 2 * 10 + 2 = 22
- EXPECT_EQ(33.0f, out[3]); // 3 * 10 + 3 = 33
+ Tensor ret(b.shape(), b.device(), b.data_type());
+ const float zero[6] = {0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f};
+ ret.CopyDataFromHostPtr<float>(zero, 6);
+ Axpy(10.0f, b, &ret);
+ ret.ToHost();
+ const float* out = ret.data<float>();
+
+ EXPECT_EQ(11.0f, out[0]); // 1.1 * 10 + 0 = 11
+ EXPECT_EQ(21.0f, out[1]); // 2.1 * 10 + 1 = 22
+ EXPECT_EQ(31.0f, out[2]); // 3.1 * 10 + 2 = 33
+ EXPECT_EQ(41.0f, out[3]); // 4.1 * 10 + 3 = 44
}
-TEST_F(OpenCL_TensorMath, Mult) {
- Tensor result = Mult(tf4in, tf4zin.T()); // Multiply with zero.
+TEST_F(OpenCL_TensorMath, GEMM) {
+ a.Reshape(Shape{6, 1});
+ Tensor result = Mult(a.T(), a);
result.ToHost();
const float* out = result.data<float>();
- EXPECT_EQ(0.0f, out[0]); // 1x4 * 4x1 = 1x1.
-
- result = Mult(tf4in, tf4in.T());
-
- result.ToHost();
- const float* out0 = result.data<float>();
-
- EXPECT_EQ(14.0f, out0[0]); // 1x4 * 4x1 = 1x1.
-
- tf16zin.SetValue(10.0f); // Multiply with 10.0.
- result = Mult(tf16in, tf16zin); // 4x4 * 4x4 = 4x4.
-
- result.ToHost();
- const float* out1 = result.data<float>();
- EXPECT_EQ(240.0f, out1[0]);
- EXPECT_EQ(280.0f, out1[1]);
- EXPECT_EQ(320.0f, out1[2]);
- EXPECT_EQ(360.0f, out1[3]);
-
- EXPECT_EQ(240.0f, out1[4]);
- EXPECT_EQ(280.0f, out1[5]);
- EXPECT_EQ(320.0f, out1[6]);
- EXPECT_EQ(360.0f, out1[7]);
-
- EXPECT_EQ(240.0f, out1[8]);
- EXPECT_EQ(280.0f, out1[9]);
- EXPECT_EQ(320.0f, out1[10]);
- EXPECT_EQ(360.0f, out1[11]);
-
- EXPECT_EQ(240.0f, out1[12]);
- EXPECT_EQ(280.0f, out1[13]);
- EXPECT_EQ(320.0f, out1[14]);
- EXPECT_EQ(360.0f, out1[15]);
+ EXPECT_EQ(91.0f, out[0]);
}
-
// TODO: ComputeCrossEntropy, SoftmaxCrossEntropy
//
#endif // USE_OPENCL
http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/595302a3/test/singa/test_opencl_device.cc
----------------------------------------------------------------------
diff --git a/test/singa/test_opencl_device.cc b/test/singa/test_opencl_device.cc
new file mode 100644
index 0000000..2797a40
--- /dev/null
+++ b/test/singa/test_opencl_device.cc
@@ -0,0 +1,108 @@
+/************************************************************
+*
+* Licensed to the Apache Software Foundation (ASF) under one
+* or more contributor license agreements. See the NOTICE file
+* distributed with this work for additional information
+* regarding copyright ownership. The ASF licenses this file
+* to you under the Apache License, Version 2.0 (the
+* "License"); you may not use this file except in compliance
+* with the License. You may obtain a copy of the License at
+*
+* http://www.apache.org/licenses/LICENSE-2.0
+*
+* Unless required by applicable law or agreed to in writing,
+* software distributed under the License is distributed on an
+* "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+* KIND, either express or implied. See the License for the
+* specific language governing permissions and limitations
+* under the License.
+*
+*************************************************************/
+
+#include "gtest/gtest.h"
+#include "singa/core/device.h"
+#include "singa/proto/core.pb.h"
+
+#ifdef USE_OPENCL
+
+using singa::CppCPU;
+using singa::OpenclDevice;
+using singa::Block;
+
+TEST(OpenclDevice, Constructor) {
+ OpenclDevice dev;
+ EXPECT_EQ(0, dev.id());
+}
+
+
+TEST(OpenclDevice, MemoryMallocFree) {
+ OpenclDevice dev;
+ Block* b = dev.NewBlock(4);
+ EXPECT_NE(nullptr, b);
+ EXPECT_EQ(4u, b->size());
+ dev.FreeBlock(b);
+}
+
+
+TEST(OpenclDevice, Exec) {
+ OpenclDevice dev;
+ Block* b = dev.NewBlock(4);
+ int x = 1, y =3, z = 0;
+ dev.Exec([x, y, &z](singa::Context *ctx) {
+ z = x + y;
+ }, {b}, {b}, false);
+ EXPECT_EQ(x + y, z);
+ dev.FreeBlock(b);
+}
+
+// Tests for integrity of one round of data transfer to an OpenCL device and back.
+TEST(OpenclDevice, CopyDataToFrom) {
+ OpenclDevice dev;
+ CppCPU host;
+
+ Block* a = host.NewBlock(4);
+ Block* b = dev.NewBlock(4);
+ Block* c = host.NewBlock(4);
+
+ // Allocate the Block object on the host.
+ char s[] = {'a', 'b', 'c', 'x'};
+ host.CopyDataFromHostPtr(a, s, 4);
+
+ // Copy back and forth.
+ dev.CopyDataToFrom(b, a, 4, singa::kHostToDevice);
+ dev.CopyDataToFrom(c, b, 4, singa::kDeviceToHost);
+
+ const char* astr = static_cast<const char*>(c->data());
+ EXPECT_EQ('a', astr[0]);
+ EXPECT_EQ('b', astr[1]);
+ EXPECT_EQ('c', astr[2]);
+ EXPECT_EQ('x', astr[3]);
+}
+
+
+TEST(OpenclDevice, DuplicateDataOnDevice) {
+ OpenclDevice dev;
+ CppCPU host;
+
+ Block* a = host.NewBlock(4);
+ Block* b = dev.NewBlock(4);
+ Block* c = dev.NewBlock(4);
+ Block* d = host.NewBlock(4);
+
+ // Allocate the Block object on the host.
+ char s[] = {'a', 'b', 'c', 'x'};
+ host.CopyDataFromHostPtr(a, s, 4);
+
+ // Copy to device and duplicate.
+ dev.CopyDataToFrom(b, a, 4, singa::kHostToDevice);
+ dev.CopyDataToFrom(c, b, 4, singa::kDeviceToDevice);
+ dev.CopyDataToFrom(d, c, 4, singa::kDeviceToHost);
+
+ const char* astr = static_cast<const char*>(d->data());
+ EXPECT_EQ('a', astr[0]);
+ EXPECT_EQ('b', astr[1]);
+ EXPECT_EQ('c', astr[2]);
+ EXPECT_EQ('x', astr[3]);
+}
+
+#endif // USE_OPENCL
[3/4] incubator-singa git commit: Update installation.md
Posted by wa...@apache.org.
Update installation.md
Added build instructions for OpenCL support.
Added instructions for getting the OpenCL ICD.
Project: http://git-wip-us.apache.org/repos/asf/incubator-singa/repo
Commit: http://git-wip-us.apache.org/repos/asf/incubator-singa/commit/5380b3a9
Tree: http://git-wip-us.apache.org/repos/asf/incubator-singa/tree/5380b3a9
Diff: http://git-wip-us.apache.org/repos/asf/incubator-singa/diff/5380b3a9
Branch: refs/heads/master
Commit: 5380b3a952accf6d9d5474088942be8411ccfbe6
Parents: 595302a
Author: Tan Li Boon <un...@users.noreply.github.com>
Authored: Wed Aug 31 11:16:50 2016 +0800
Committer: Tan Li Boon <ta...@u.nus.edu>
Committed: Wed Aug 31 11:39:16 2016 +0800
----------------------------------------------------------------------
doc/en/docs/installation.md | 23 +++++++++++++++++++++++
1 file changed, 23 insertions(+)
----------------------------------------------------------------------
http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/5380b3a9/doc/en/docs/installation.md
----------------------------------------------------------------------
diff --git a/doc/en/docs/installation.md b/doc/en/docs/installation.md
index 3fb6ba9..ecf3e97 100755
--- a/doc/en/docs/installation.md
+++ b/doc/en/docs/installation.md
@@ -14,6 +14,7 @@
* lmdb (tested with 0.9)
* cuda (tested with 6.5, 7.0 and 7.5)
* cudnn (v4 and v5)
+* opencl-headers and viennacl (version 1.7.1 or newer) for OpenCL support
PySINGA has additional dependencies
@@ -146,6 +147,28 @@ After compiling SINGA, you can run the unit tests by
You can see all the testing cases with testing results. If SINGA passes all
tests, then you have successfully installed SINGA. Please proceed to try the examples!
+### OpenCL support (Linux)
+
+Install `opencl-headers` and `libviennacl-dev` (Ubuntu 16) or `viennacl` (Fedora).
+
+Additionally, you will need the OpenCL Installable Client Driver (ICD) for the platforms that you want to run OpenCL on.
+
+* For AMD and nVidia GPUs, the driver package should also install the correct OpenCL ICD.
+* For Intel CPUs and/or GPUs, get the driver from the [Intel website.](https://software.intel.com/en-us/articles/opencl-drivers) Note that the drivers provided on that website only supports recent CPUs and Iris GPUs.
+* For older Intel CPUs, you can use the `beignet-opencl-icd` package.
+
+Note that running OpenCL on CPUs is not currently recommended because it is slow. Memory transfer is on the order of whole seconds (1000's of ms on CPUs as compared to 1's of ms on GPUs).
+
+More information on setting up a working OpenCL environment may be found [here](https://wiki.tiker.net/OpenCLHowTo).
+
+If the package version of ViennaCL is not at least 1.7.1, you will need to build it from source:
+
+Clone [the repository from here](https://github.com/viennacl/viennacl-dev), checkout the `release-1.7.1` tag and build it.
+Remember to add its directory to `PATH` and the built libraries to `LD_LIBRARY_PATH`.
+
+To build SINGA with OpenCL support, you need to pass the flag during cmake:
+
+ cmake .. -DUSE_OPENCL=ON
### Windows
To be added.
[4/4] incubator-singa git commit: Merge PR #251 for updating opencl
support using ViennaCL
Posted by wa...@apache.org.
Merge PR #251 for updating opencl support using ViennaCL
Conflicts:
doc/en/docs/installation.md
Project: http://git-wip-us.apache.org/repos/asf/incubator-singa/repo
Commit: http://git-wip-us.apache.org/repos/asf/incubator-singa/commit/1996e8a9
Tree: http://git-wip-us.apache.org/repos/asf/incubator-singa/tree/1996e8a9
Diff: http://git-wip-us.apache.org/repos/asf/incubator-singa/diff/1996e8a9
Branch: refs/heads/master
Commit: 1996e8a950f6c204cbc418ca1737cbfe752998ad
Parents: 61ec2ec 5380b3a
Author: Wei Wang <wa...@comp.nus.edu.sg>
Authored: Wed Aug 31 12:34:17 2016 +0800
Committer: Wei Wang <wa...@comp.nus.edu.sg>
Committed: Wed Aug 31 12:34:17 2016 +0800
----------------------------------------------------------------------
.travis.yml | 6 +-
CMakeLists.txt | 3 +-
LICENSE | 4 +-
cmake/Dependencies.cmake | 25 +-
cmake/Thirdparty/FindViennaCL.cmake | 45 ++
doc/en/docs/installation.md | 23 +
include/singa/core/common.h | 10 +-
include/singa/core/device.h | 56 +-
include/singa/utils/opencl_utils.h | 142 +---
src/core/device/opencl_device.cc | 182 ++---
src/core/tensor/tensor_math_opencl.cl | 19 +-
src/core/tensor/tensor_math_opencl.h | 1007 ++++++++--------------------
src/utils/opencl_utils.cc | 63 --
test/singa/test_opencl.cc | 700 ++++++++-----------
test/singa/test_opencl_device.cc | 108 +++
15 files changed, 845 insertions(+), 1548 deletions(-)
----------------------------------------------------------------------
http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/1996e8a9/cmake/Dependencies.cmake
----------------------------------------------------------------------
diff --cc cmake/Dependencies.cmake
index 5fb7578,aa2212b..a412151
--- a/cmake/Dependencies.cmake
+++ b/cmake/Dependencies.cmake
@@@ -55,23 -55,26 +55,26 @@@ ENDIF(
IF(USE_OPENCL)
FIND_PACKAGE(OpenCL REQUIRED)
- IF(NOT OPENCL_FOUND)
+ IF(NOT OpenCL_FOUND)
MESSAGE(SEND_ERROR "OpenCL was requested, but not found.")
ELSE()
- INCLUDE_DIRECTORIES(SYSTEM ${OpenCL_INCPATH})
- LIST(APPEND SINGA_LINKER_LIBS ${OPENCL_LIBRARIES})
- MESSAGE(STATUS "Found OpenCL at ${OPENCL_INCLUDE_DIRS}")
- IF(NOT OPENCL_HAS_CPP_BINDINGS)
- MESSAGE(SEND_ERROR "OpenCL C++ bindings cl2.hpp was not found.")
- ELSE()
- MESSAGE(STATUS "Found OpenCL C++ bindings.")
- ENDIF()
- ENDIF()
+ MESSAGE(STATUS "Found OpenCL headers at ${OpenCL_INCLUDE_DIRS}")
+ INCLUDE_DIRECTORIES(SYSTEM ${OpenCL_INCLUDE_DIR})
+ LIST(APPEND SINGA_LINKER_LIBS ${OpenCL_LIBRARIES})
+ FIND_PACKAGE(ViennaCL REQUIRED)
+ IF(NOT ViennaCL_FOUND)
+ MESSAGE(SEND_ERROR "ViennaCL is required if OpenCL is enabled.")
+ ELSE()
+ MESSAGE(STATUS "Found ViennaCL headers at ${ViennaCL_INCLUDE_DIR}")
+ INCLUDE_DIRECTORIES(SYSTEM ${ViennaCL_INCLUDE_DIR})
+ LIST(APPEND SINGA_LINKER_LIBS ${ViennaCL_LIBRARIES})
+ ENDIF()
+ ENDIF()
ENDIF()
-FIND_PACKAGE(Glog REQUIRED)
-INCLUDE_DIRECTORIES(SYSTEM ${GLOG_INCLUDE_DIRS})
-LIST(APPEND SINGA_LINKER_LIBS ${GLOG_LIBRARIES})
+#FIND_PACKAGE(Glog REQUIRED)
+#INCLUDE_DIRECTORIES(SYSTEM ${GLOG_INCLUDE_DIRS})
+#LIST(APPEND SINGA_LINKER_LIBS ${GLOG_LIBRARIES})
#MESSAGE(STATUS "Found glog at ${GLOG_INCLUDE_DIRS}")
IF(USE_OPENCV)
http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/1996e8a9/doc/en/docs/installation.md
----------------------------------------------------------------------
diff --cc doc/en/docs/installation.md
index 08508dc,ecf3e97..a8b0530
--- a/doc/en/docs/installation.md
+++ b/doc/en/docs/installation.md
@@@ -9,13 -9,14 +9,14 @@@
### Optional
-* glog
-* opencv (tested with 2.4.8)
-* lmdb (tested with 0.9)
* cuda (tested with 6.5, 7.0 and 7.5)
* cudnn (v4 and v5)
+* opencv (tested with 2.4.8)
+* lmdb (tested with 0.9)
+* glog
+ * opencl-headers and viennacl (version 1.7.1 or newer) for OpenCL support
-PySINGA has additional dependencies
+PySINGA (the Python binding) has additional dependencies
* python(==2.7)
* pip(>=1.5)
@@@ -191,6 -147,28 +192,28 @@@ After compiling SINGA, you can run the
You can see all the testing cases with testing results. If SINGA passes all
tests, then you have successfully installed SINGA. Please proceed to try the examples!
+ ### OpenCL support (Linux)
+
+ Install `opencl-headers` and `libviennacl-dev` (Ubuntu 16) or `viennacl` (Fedora).
+
+ Additionally, you will need the OpenCL Installable Client Driver (ICD) for the platforms that you want to run OpenCL on.
+
+ * For AMD and nVidia GPUs, the driver package should also install the correct OpenCL ICD.
+ * For Intel CPUs and/or GPUs, get the driver from the [Intel website.](https://software.intel.com/en-us/articles/opencl-drivers) Note that the drivers provided on that website only supports recent CPUs and Iris GPUs.
+ * For older Intel CPUs, you can use the `beignet-opencl-icd` package.
+
+ Note that running OpenCL on CPUs is not currently recommended because it is slow. Memory transfer is on the order of whole seconds (1000's of ms on CPUs as compared to 1's of ms on GPUs).
+
+ More information on setting up a working OpenCL environment may be found [here](https://wiki.tiker.net/OpenCLHowTo).
+
+ If the package version of ViennaCL is not at least 1.7.1, you will need to build it from source:
+
+ Clone [the repository from here](https://github.com/viennacl/viennacl-dev), checkout the `release-1.7.1` tag and build it.
+ Remember to add its directory to `PATH` and the built libraries to `LD_LIBRARY_PATH`.
+
+ To build SINGA with OpenCL support, you need to pass the flag during cmake:
-
++
+ cmake .. -DUSE_OPENCL=ON
### Windows
To be added.
[2/4] incubator-singa git commit: SINGA-243 ViennaCL backend for
OpenCL support
Posted by wa...@apache.org.
SINGA-243 ViennaCL backend for OpenCL support
- Replaced current OpenCL backend with ViennaCL libs
- Removed need for OpenCL C++ headers
- Updated unit tests files to match changes
- Updated Dependencies.cmake to match changes
- Added license information.
Project: http://git-wip-us.apache.org/repos/asf/incubator-singa/repo
Commit: http://git-wip-us.apache.org/repos/asf/incubator-singa/commit/595302a3
Tree: http://git-wip-us.apache.org/repos/asf/incubator-singa/tree/595302a3
Diff: http://git-wip-us.apache.org/repos/asf/incubator-singa/diff/595302a3
Branch: refs/heads/master
Commit: 595302a36e9fee2ef63b57ddf81fe75a5fe00a40
Parents: 65bf582
Author: Tan Li Boon <un...@users.noreply.github.com>
Authored: Wed Aug 17 12:48:18 2016 +0800
Committer: Tan Li Boon <ta...@u.nus.edu>
Committed: Wed Aug 31 10:53:21 2016 +0800
----------------------------------------------------------------------
.travis.yml | 6 +-
CMakeLists.txt | 3 +-
LICENSE | 4 +-
cmake/Dependencies.cmake | 25 +-
cmake/Thirdparty/FindViennaCL.cmake | 45 ++
include/singa/core/common.h | 10 +-
include/singa/core/device.h | 56 +-
include/singa/utils/opencl_utils.h | 142 +---
src/core/device/opencl_device.cc | 182 ++---
src/core/tensor/tensor_math_opencl.cl | 19 +-
src/core/tensor/tensor_math_opencl.h | 1007 ++++++++--------------------
src/utils/opencl_utils.cc | 63 --
test/singa/test_opencl.cc | 700 ++++++++-----------
test/singa/test_opencl_device.cc | 108 +++
14 files changed, 822 insertions(+), 1548 deletions(-)
----------------------------------------------------------------------
http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/595302a3/.travis.yml
----------------------------------------------------------------------
diff --git a/.travis.yml b/.travis.yml
index 8b1f89c..d9bba51 100644
--- a/.travis.yml
+++ b/.travis.yml
@@ -7,13 +7,13 @@ before_install:
- sudo apt-get -qq update
- sudo apt-get install -qq -y libopenblas-dev libgoogle-glog-dev libprotobuf-dev protobuf-compiler
- sudo apt-get install -qq -y opencl-headers ocl-icd-*
- - wget https://github.com/KhronosGroup/OpenCL-CLHPP/releases/download/v2.0.9/cl2.hpp
- - sudo mv cl2.hpp /usr/include/CL/
+#- wget https://github.com/KhronosGroup/OpenCL-CLHPP/releases/download/v2.0.9/cl2.hpp
+#- sudo mv cl2.hpp /usr/include/CL/
#- sudo apt-get install -qq libgtest-dev
before_script:
- mkdir build && cd build
- - cmake .. -DUSE_CUDA=OFF -DUSE_CUDNN=OFF -DUSE_PYTHON=OFF -DBUILD_OPENCL_TESTS=OFF
+ - cmake .. -DUSE_CUDA=OFF -DUSE_CUDNN=OFF -DUSE_PYTHON=OFF
script:
- make
http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/595302a3/CMakeLists.txt
----------------------------------------------------------------------
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 3f6bea2..611cee4 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -43,9 +43,8 @@ OPTION(USE_CUDNN "Use Cudnn libs" ON)
OPTION(USE_OPENCV "Use opencv" OFF)
OPTION(USE_LMDB "Use LMDB libs" OFF)
OPTION(USE_PYTHON "Generate py wrappers" OFF)
-#OPTION(USE_OPENCL "Use OpenCL" OFF)
+OPTION(USE_OPENCL "Use OpenCL" OFF)
OPTION(ENABLE_DIST "enable distributed training" OFF)
-#OPTION(BUILD_OPENCL_TESTS "Build OpenCL tests" OFF)
INCLUDE("cmake/Dependencies.cmake")
INCLUDE("cmake/Utils.cmake")
http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/595302a3/LICENSE
----------------------------------------------------------------------
diff --git a/LICENSE b/LICENSE
index f658def..4f9d1e7 100644
--- a/LICENSE
+++ b/LICENSE
@@ -303,7 +303,9 @@ SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
===========================================================================
SINGA bundles the following under BSD 2-clause license:
-include/singa/utils/cuda_utils.h, src/core/tensor/distribution.cl
+include/singa/utils/cuda_utils.h
+src/core/tensor/distribution.cl
+cmake/ThirdParty/FindViennaCL.cmake
COPYRIGHT
http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/595302a3/cmake/Dependencies.cmake
----------------------------------------------------------------------
diff --git a/cmake/Dependencies.cmake b/cmake/Dependencies.cmake
index b5fda6d..aa2212b 100644
--- a/cmake/Dependencies.cmake
+++ b/cmake/Dependencies.cmake
@@ -50,23 +50,26 @@ IF(USE_CBLAS)
FIND_PACKAGE(CBLAS REQUIRED)
INCLUDE_DIRECTORIES(SYSTEM ${CBLAS_INCLUDE_DIR})
LIST(APPEND SINGA_LINKER_LIBS ${CBLAS_LIBRARIES})
- MESSAGE(STATUS "FOUND cblas at ${CBLAS_LIBRARIES}")
+ MESSAGE(STATUS "Found cblas at ${CBLAS_LIBRARIES}")
ENDIF()
IF(USE_OPENCL)
FIND_PACKAGE(OpenCL REQUIRED)
- IF(NOT OPENCL_FOUND)
+ IF(NOT OpenCL_FOUND)
MESSAGE(SEND_ERROR "OpenCL was requested, but not found.")
ELSE()
- INCLUDE_DIRECTORIES(SYSTEM ${OpenCL_INCPATH})
- LIST(APPEND SINGA_LINKER_LIBS ${OPENCL_LIBRARIES})
- MESSAGE(STATUS "Found OpenCL at ${OPENCL_INCLUDE_DIRS}")
- IF(NOT OPENCL_HAS_CPP_BINDINGS)
- MESSAGE(SEND_ERROR "OpenCL C++ bindings cl2.hpp was not found.")
- ELSE()
- MESSAGE(STATUS "Found OpenCL C++ bindings.")
- ENDIF()
- ENDIF()
+ MESSAGE(STATUS "Found OpenCL headers at ${OpenCL_INCLUDE_DIRS}")
+ INCLUDE_DIRECTORIES(SYSTEM ${OpenCL_INCLUDE_DIR})
+ LIST(APPEND SINGA_LINKER_LIBS ${OpenCL_LIBRARIES})
+ FIND_PACKAGE(ViennaCL REQUIRED)
+ IF(NOT ViennaCL_FOUND)
+ MESSAGE(SEND_ERROR "ViennaCL is required if OpenCL is enabled.")
+ ELSE()
+ MESSAGE(STATUS "Found ViennaCL headers at ${ViennaCL_INCLUDE_DIR}")
+ INCLUDE_DIRECTORIES(SYSTEM ${ViennaCL_INCLUDE_DIR})
+ LIST(APPEND SINGA_LINKER_LIBS ${ViennaCL_LIBRARIES})
+ ENDIF()
+ ENDIF()
ENDIF()
FIND_PACKAGE(Glog REQUIRED)
http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/595302a3/cmake/Thirdparty/FindViennaCL.cmake
----------------------------------------------------------------------
diff --git a/cmake/Thirdparty/FindViennaCL.cmake b/cmake/Thirdparty/FindViennaCL.cmake
new file mode 100644
index 0000000..263c80f
--- /dev/null
+++ b/cmake/Thirdparty/FindViennaCL.cmake
@@ -0,0 +1,45 @@
+# This file is retrieved from caffe/cmake/Modules/FindViennaCL.cmake.
+
+SET(ViennaCL_WITH_OPENCL TRUE)
+
+SET(VIENNACL_INCLUDE_SEARCH_PATHS
+ ..
+ /usr/include
+ /usr/local/include
+ /opt/ViennaCL/include
+ $ENV{VIENNACL_HOME}
+ $ENV{VIENNACL_HOME}/include
+)
+
+FIND_PATH(ViennaCL_INCLUDE_DIR NAMES viennacl/forwards.h PATHS ${VIENNACL_INCLUDE_SEARCH_PATHS})
+
+SET(ViennaCL_FOUND ON)
+
+# Check include files
+IF(NOT ViennaCL_INCLUDE_DIR)
+ SET(ViennaCL_FOUND OFF)
+ MESSAGE(STATUS "Could not find ViennaCL include. Turning ViennaCL_FOUND off")
+ENDIF()
+
+IF (ViennaCL_FOUND)
+ IF (NOT ViennaCL_FIND_QUIETLY)
+ MESSAGE(STATUS "Found ViennaCL include: ${ViennaCL_INCLUDE_DIR}")
+ ENDIF (NOT ViennaCL_FIND_QUIETLY)
+ELSE (ViennaCL_FOUND)
+ IF (ViennaCL_FIND_REQUIRED)
+ MESSAGE(FATAL_ERROR "Could not find ViennaCL")
+ ENDIF (ViennaCL_FIND_REQUIRED)
+ENDIF (ViennaCL_FOUND)
+
+IF(ViennaCL_WITH_OPENCL)
+ find_package(OpenCL REQUIRED)
+ENDIF(ViennaCL_WITH_OPENCL)
+
+set(ViennaCL_INCLUDE_DIRS ${ViennaCL_INCLUDE_DIR} ${OPENCL_INCLUDE_DIRS})
+set(ViennaCL_LIBRARIES ${OPENCL_LIBRARIES})
+
+MARK_AS_ADVANCED(
+ ViennaCL_INCLUDE_DIR
+ ViennaCL_INCLUDE_DIRS
+ ViennaCL_LIBRARIES
+)
http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/595302a3/include/singa/core/common.h
----------------------------------------------------------------------
diff --git a/include/singa/core/common.h b/include/singa/core/common.h
index 53a9726..dc552c1 100644
--- a/include/singa/core/common.h
+++ b/include/singa/core/common.h
@@ -36,10 +36,7 @@
#ifdef USE_OPENCL
-#define CL_HPP_MINIMUM_OPENCL_VERSION 120
-#define CL_HPP_TARGET_OPENCL_VERSION 120
-#include <CL/cl2.hpp>
-#include <unordered_map>
+#include "singa/utils/opencl_utils.h"
#endif // USE_OPENCL
using std::atomic;
@@ -110,9 +107,8 @@ typedef struct _Context {
#endif // USE_CUDA
#ifdef USE_OPENCL
- std::shared_ptr<std::unordered_map<std::string, cl::Kernel>> kernels;
- cl::CommandQueue ocl_cmdq;
- cl::Context ocl_ctx;
+ // This stores the context ID of the OpenCL context controlled by ViennaCL.
+ long vcl_ctx_id;
#endif
} Context;
http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/595302a3/include/singa/core/device.h
----------------------------------------------------------------------
diff --git a/include/singa/core/device.h b/include/singa/core/device.h
index 810d41f..62fa250 100644
--- a/include/singa/core/device.h
+++ b/include/singa/core/device.h
@@ -36,12 +36,6 @@
#endif // USE_CUDA
#ifdef USE_OPENCL
-// http://github.khronos.org/OpenCL-CLHPP/
-// cl2.hpp includes cl.h, do not re-include.
-#define CL_HPP_MINIMUM_OPENCL_VERSION 120
-#define CL_HPP_TARGET_OPENCL_VERSION 120
-#include <unordered_map>
-#include <CL/cl2.hpp>
#include "singa/utils/opencl_utils.h"
#endif // USE_OPENCL
@@ -217,50 +211,26 @@ public:
OpenclDevice(int id = 0, int num_executors = 1);
~OpenclDevice();
- /// Get the specified kernel.
- cl::Kernel GetKernel(const std::string& kname, cl_int* status = nullptr);
-
- /// Get the command queue associated with this device.
- cl::CommandQueue GetCmdQ() { return cmdq; }
-
- /// Prints information about all Devices in each Platform.
- void PrintAllDeviceInfo();
-
- /// Prints status about CL source code builds.
- void PrintClBuildInfo(cl::Program &p);
-
// Overridden, inherited methods
void SetRandSeed(unsigned seed) override;
void CopyDataToFrom(Block* dst, Block* src, size_t nBytes,
CopyDirection direction, int dst_offset = 0,
int src_offset = 0);
-/*
- void CopyDataFromHostPtr(Block* dst, const void* src, size_t nBytes = 0,
- size_t dst_offset = 0) override;*/
protected:
/// The OpenCL device that this object represents.
/// Each OpenclDevice contains exactly one cl::Device for the lifetime of the
/// object.
- cl::Device this_device;
+ viennacl::ocl::device this_device;
/// Each OpenclDevice has one OpenCL context. It is created along with the
/// creation of this object.
- cl::Context ocl_ctx;
-
- /// The CommandQueue that is associated with this device.
- /// Since each OpenclDevice contains only one cl::Device and one cl::Context,
- /// it naturally also contains one cl::CommandQueue that is associated
- /// with said Device and Context.
- cl::CommandQueue cmdq;
-
- /// A list of kernels that has been compiled on this device.
- std::shared_ptr<std::unordered_map<std::string, cl::Kernel>> kernels;
+ viennacl::ocl::context vcl_ctx;
/// Searches the given paths for all .cl files and builds
/// OpenCL programs, then stores them in the Kernels map.
- void BuildPrograms(const std::string &kdir = cl_src_path);
+ void BuildPrograms(const std::string &kdir);
// Overridden, inherited methods.
@@ -280,21 +250,6 @@ protected:
private:
- /// Copies a data block from host to device.
- /// src: a pointer to an array of data.
- /// dst: a pointer to a cl::Buffer object.
- void WriteToDevice(cl::Buffer* dst, const void* src, const size_t size);
-
- /// Reads a data block from device to host.
- /// src: a pointer to an cl::Buffer object.
- /// dst: a pointer to an malloc'ed empty array.
- void ReadFromDevice(void* dst, const cl::Buffer* src, const size_t size);
-
- /// Duplicates a block of data on the device.
- /// src: a pointer to the original cl::Buffer object.
- /// dst: a pointer to the new cl::Buffer object to copy the data into.
- void CopyDeviceBuffer(cl::Buffer* dst, const cl::Buffer* src, const size_t size);
-
static const std::string cl_src_path;
};
#endif // USE_OPENCL
@@ -368,11 +323,6 @@ public:
/// except the context initialization.
static bool CheckDevice(const int device_id);
-
-private:
-#ifdef USE_OPENCL
- cl::Platform clPlatform;
-#endif // USE_OPENCL
};
http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/595302a3/include/singa/utils/opencl_utils.h
----------------------------------------------------------------------
diff --git a/include/singa/utils/opencl_utils.h b/include/singa/utils/opencl_utils.h
index 664a9e1..8c05643 100644
--- a/include/singa/utils/opencl_utils.h
+++ b/include/singa/utils/opencl_utils.h
@@ -24,121 +24,47 @@
#ifdef USE_OPENCL
-#include <iostream>
+#define CL_USE_DEPRECATED_OPENCL_1_2_APIS
-// http://github.khronos.org/OpenCL-CLHPP/
-// cl2.hpp includes cl.h, do not re-include.
#define CL_HPP_MINIMUM_OPENCL_VERSION 120
#define CL_HPP_TARGET_OPENCL_VERSION 120
-#include <CL/cl2.hpp>
-#define CL_BREAK_ON_FAILURE if (status != CL_SUCCESS) return;
-
-
-inline const char* clGetBuildInfoString(const cl_build_status status) {
- switch (status) {
- case CL_BUILD_NONE: return "CL_BUILD_NONE";
- case CL_BUILD_ERROR: return "CL_BUILD_ERROR";
- case CL_BUILD_SUCCESS: return "CL_BUILD_SUCCESS";
- case CL_BUILD_IN_PROGRESS: return "CL_BUILD_IN_PROGRESS";
- default: return "";
- }
-}
-
-
-inline const char* clGetErrorString(const cl_int status) {
-
- switch(status) {
-
- // Run-time and JIT compiler errors
- case 0: return "CL_SUCCESS";
- case -1: return "CL_DEVICE_NOT_FOUND";
- case -2: return "CL_DEVICE_NOT_AVAILABLE";
- case -3: return "CL_COMPILER_NOT_AVAILABLE";
- case -4: return "CL_MEM_OBJECT_ALLOCATION_FAILURE";
- case -5: return "CL_OUT_OF_RESOURCES";
- case -6: return "CL_OUT_OF_HOST_MEMORY";
- case -7: return "CL_PROFILING_INFO_NOT_AVAILABLE";
- case -8: return "CL_MEM_COPY_OVERLAP";
- case -9: return "CL_IMAGE_FORMAT_MISMATCH";
- case -10: return "CL_IMAGE_FORMAT_NOT_SUPPORTED";
- case -11: return "CL_BUILD_PROGRAM_FAILURE";
- case -12: return "CL_MAP_FAILURE";
- case -13: return "CL_MISALIGNED_SUB_BUFFER_OFFSET";
- case -14: return "CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST";
- case -15: return "CL_COMPILE_PROGRAM_FAILURE";
- case -16: return "CL_LINKER_NOT_AVAILABLE";
- case -17: return "CL_LINK_PROGRAM_FAILURE";
- case -18: return "CL_DEVICE_PARTITION_FAILED";
- case -19: return "CL_KERNEL_ARG_INFO_NOT_AVAILABLE";
-
- // Compile-time errors
- case -30: return "CL_INVALID_VALUE";
- case -31: return "CL_INVALID_DEVICE_TYPE";
- case -32: return "CL_INVALID_PLATFORM";
- case -33: return "CL_INVALID_DEVICE";
- case -34: return "CL_INVALID_CONTEXT";
- case -35: return "CL_INVALID_QUEUE_PROPERTIES";
- case -36: return "CL_INVALID_COMMAND_QUEUE";
- case -37: return "CL_INVALID_HOST_PTR";
- case -38: return "CL_INVALID_MEM_OBJECT";
- case -39: return "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR";
- case -40: return "CL_INVALID_IMAGE_SIZE";
- case -41: return "CL_INVALID_SAMPLER";
- case -42: return "CL_INVALID_BINARY";
- case -43: return "CL_INVALID_BUILD_OPTIONS";
- case -44: return "CL_INVALID_PROGRAM";
- case -45: return "CL_INVALID_PROGRAM_EXECUTABLE";
- case -46: return "CL_INVALID_KERNEL_NAME";
- case -47: return "CL_INVALID_KERNEL_DEFINITION";
- case -48: return "CL_INVALID_KERNEL";
- case -49: return "CL_INVALID_ARG_INDEX";
- case -50: return "CL_INVALID_ARG_VALUE";
- case -51: return "CL_INVALID_ARG_SIZE";
- case -52: return "CL_INVALID_KERNEL_ARGS";
- case -53: return "CL_INVALID_WORK_DIMENSION";
- case -54: return "CL_INVALID_WORK_GROUP_SIZE";
- case -55: return "CL_INVALID_WORK_ITEM_SIZE";
- case -56: return "CL_INVALID_GLOBAL_OFFSET";
- case -57: return "CL_INVALID_EVENT_WAIT_LIST";
- case -58: return "CL_INVALID_EVENT";
- case -59: return "CL_INVALID_OPERATION";
- case -60: return "CL_INVALID_GL_OBJECT";
- case -61: return "CL_INVALID_BUFFER_SIZE";
- case -62: return "CL_INVALID_MIP_LEVEL";
- case -63: return "CL_INVALID_GLOBAL_WORK_SIZE";
- case -64: return "CL_INVALID_PROPERTY";
- case -65: return "CL_INVALID_IMAGE_DESCRIPTOR";
- case -66: return "CL_INVALID_COMPILER_OPTIONS";
- case -67: return "CL_INVALID_LINKER_OPTIONS";
- case -68: return "CL_INVALID_DEVICE_PARTITION_COUNT";
-
- // Extension errors
- case -1000: return "CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR";
- case -1001: return "CL_PLATFORM_NOT_FOUND_KHR";
- case -1002: return "CL_INVALID_D3D10_DEVICE_KHR";
- case -1003: return "CL_INVALID_D3D10_RESOURCE_KHR";
- case -1004: return "CL_D3D10_RESOURCE_ALREADY_ACQUIRED_KHR";
- case -1005: return "CL_D3D10_RESOURCE_NOT_ACQUIRED_KHR";
-
- default: return "Unknown OpenCL status";
+#ifndef VIENNACL_WITH_OPENCL
+ #define VIENNACL_WITH_OPENCL
+#endif
+
+#ifndef __APPLE__
+ #include "CL/cl.h"
+#else
+ #include "OpenCL/cl.h"
+#endif
+
+#include <viennacl/backend/opencl.hpp>
+
+#include <viennacl/ocl/device.hpp>
+#include <viennacl/ocl/platform.hpp>
+#include <viennacl/ocl/backend.hpp>
+#include <viennacl/ocl/device_utils.hpp>
+#include <viennacl/ocl/utils.hpp>
+#include <viennacl/ocl/program.hpp>
+#include <viennacl/ocl/kernel.hpp>
+
+
+inline viennacl::ocl::handle<cl_mem>
+WrapHandle(cl_mem in, viennacl::ocl::context *ctx) {
+ if (in != nullptr) {
+ viennacl::ocl::handle<cl_mem> memhandle(in, *ctx);
+ memhandle.inc();
+ return memhandle;
+ } else {
+ cl_int err;
+ cl_mem dummy = clCreateBuffer(ctx->handle().get(), CL_MEM_READ_WRITE, 0,
+ nullptr, &err);
+ viennacl::ocl::handle<cl_mem> memhandle(dummy, *ctx);
+ return memhandle;
}
}
-
-/// Special function used to perform error checking and logging.
-inline bool OCL_CHECK(const cl_int status, const char* what) {
- if (status == CL_SUCCESS) return true; // Nothing wrong.
- LOG(ERROR) << status << ": " << clGetErrorString(status) << " " << what << std::endl;
- return false;
-}
-
-/// Prints information about the specified Platform.
-void PrintPlatformInfo(const cl::Platform &p);
-
-/// Prints information about the specified Device.
-void PrintDeviceInfo(const cl::Device &dev);
-
#endif // USE_OPENCL
#endif // SINGA_UTILS_OPENCL_UTILS_H_
http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/595302a3/src/core/device/opencl_device.cc
----------------------------------------------------------------------
diff --git a/src/core/device/opencl_device.cc b/src/core/device/opencl_device.cc
index b941cd2..6b371c4 100644
--- a/src/core/device/opencl_device.cc
+++ b/src/core/device/opencl_device.cc
@@ -23,89 +23,41 @@
#include "singa/core/device.h"
#include "singa/utils/tinydir.h"
+#include "singa/utils/opencl_utils.h"
#ifdef USE_OPENCL
-using std::string;
+using namespace viennacl;
+using namespace viennacl::backend::opencl;
namespace singa {
-const string OpenclDevice::cl_src_path = "../src/core/tensor";
+const std::string OpenclDevice::cl_src_path = "../src/core/tensor";
OpenclDevice::OpenclDevice(int id, int num_executors)
: Device(id, num_executors) {
+ CHECK_GE(id, 0);
lang_ = kOpencl;
- this->kernels = std::make_shared<std::unordered_map<string, cl::Kernel>>();
-
- // Create the OpenCL Device, Context, and CommandQueue.
- /// TODO: This merely chooses the first device on the first platform.
- cl_int status = CL_SUCCESS;
-
- std::vector<cl::Platform> platforms;
- status = cl::Platform::get(&platforms);
- OCL_CHECK(status, "Failed to find any OpenCL platforms!");
-
- std::vector<cl::Device> devices;
- status = platforms[0].getDevices(CL_DEVICE_TYPE_ALL, &devices);
- OCL_CHECK(status, "Failed to get list of devices from platform!");
-
- this->this_device = cl::Device(devices[0]);
- this->ocl_ctx = cl::Context(this_device, nullptr, nullptr, nullptr, &status);
- OCL_CHECK(status, "Failed to create context!");
-
- this->cmdq = cl::CommandQueue(ocl_ctx, this_device, CL_QUEUE_PROFILING_ENABLE, &status);
- OCL_CHECK(status, "Failed to create a command queue!");
-
- BuildPrograms();
-
- ctx_.kernels = kernels;
- ctx_.ocl_cmdq = cmdq;
- ctx_.ocl_ctx = ocl_ctx;
+
+ ocl::current_context().build_options("-cl-std=CL1.2");
+
+ ctx_.vcl_ctx_id = 0;
+ this->this_device = ocl::current_device();
+
+ BuildPrograms(cl_src_path);
}
OpenclDevice::~OpenclDevice() {
// Flush and finish the command queue.
+ auto cmdq = ocl::current_context().get_queue();
+
cmdq.flush();
cmdq.finish();
}
-cl::Kernel OpenclDevice::GetKernel(const std::string& kname, cl_int* status) {
- if (!status) *status = CL_SUCCESS;
- if (kernels->find(kname) == kernels->end()) {
- // TODO: Not found
- LOG(ERROR) << "Error: Kernel " << kname << " could not be found!";
- if (!status) *status = CL_INVALID_KERNEL;
- }
- return kernels->at(kname);
-}
-
-/*
-void OpenclDevice::PrintAllDeviceInfo() {
- cl_int status = CL_SUCCESS;
-
- for (auto dev : devices) {
- PrintDeviceInfo(d);
- }
-}
-*/
-
-
-void OpenclDevice::PrintClBuildInfo(cl::Program &p) {
- cl_int status = CL_SUCCESS;
-
- auto buildStatus = p.getBuildInfo<CL_PROGRAM_BUILD_STATUS>(&status);
- for (auto pair : buildStatus)
- std::cout << clGetBuildInfoString(pair.second) << std::endl;
-
- auto buildLog = p.getBuildInfo<CL_PROGRAM_BUILD_LOG>(&status);
- for (auto pair : buildLog)
- std::cout << pair.second << std::endl;
-}
-
-
void OpenclDevice::SetRandSeed(unsigned seed) { seed = seed; }
@@ -113,19 +65,33 @@ void OpenclDevice::CopyDataToFrom(Block* dst, Block* src, size_t nBytes,
CopyDirection direction, int dst_offset, int src_offset) {
// Pointers must be valid.
if (!dst || !src) return;
+
+ auto ocl_ctx = viennacl::ocl::get_context(ctx_.vcl_ctx_id);
- CopyToFrom(dst->mutable_data(), src->data(), nBytes, direction);
+ switch(direction) {
+ case kHostToDevice: {
+ auto dst_handle = WrapHandle(static_cast<cl_mem>(dst->mutable_data()), &ocl_ctx);
+ memory_write(dst_handle, dst_offset, nBytes, src->data());
+ return;
+ }
+ case kDeviceToHost: {
+ auto src_handle = WrapHandle(static_cast<cl_mem>(src->mutable_data()), &ocl_ctx);
+ memory_read(src_handle, src_offset, nBytes, dst->mutable_data());
+ return;
+ }
+ case kDeviceToDevice: {
+ auto src_handle = WrapHandle(static_cast<cl_mem>(src->mutable_data()), &ocl_ctx);
+ auto dst_handle = WrapHandle(static_cast<cl_mem>(dst->mutable_data()), &ocl_ctx);
+ memory_copy(src_handle, dst_handle, src_offset, dst_offset, nBytes);
+ return;
+ }
+ default:
+ return;
+ }
}
-/*
-void OpenclDevice::CopyDataFromHostPtr(Block* dst, const void* src, size_t nBytes, size_t dst_offset) {
- CopyToFrom(dst->mutable_data(), src, 4, kHostToDevice);
-}
-*/
void OpenclDevice::BuildPrograms(const std::string &kdir) {
- cl_int status = CL_SUCCESS;
-
tinydir_dir dir;
tinydir_open(&dir, kdir.c_str());
@@ -137,63 +103,47 @@ void OpenclDevice::BuildPrograms(const std::string &kdir) {
tinydir_next(&dir);
continue;
}
-
+
std::ifstream clFile(file.path, std::ios_base::binary);
std::stringstream buffer;
buffer << clFile.rdbuf();
std::string clSrc(buffer.str());
- cl::Program program(this->ocl_ctx, clSrc, false, &status);
- OCL_CHECK(status, "Program creation failed.");
- status = program.build({this_device}, "-cl-std=CL1.2");
- if (status == CL_SUCCESS) {
- std::vector<cl::Kernel> built_kernels;
- status = program.createKernels(&built_kernels);
- OCL_CHECK(status, "Failed to create kernels in built program.");
-
- for (auto k : built_kernels) {
- std::string name = k.getInfo<CL_KERNEL_FUNCTION_NAME>(&status);
- this->kernels->insert(std::make_pair(name, k));
- }
- } else {
- OCL_CHECK(status, "Build failed on source path");
- LOG(ERROR) << file.path << std::endl;
- PrintClBuildInfo(program);
- }
+ std::string name(file.name);
+ ocl::current_context().add_program(clSrc, name);
tinydir_next(&dir);
}
}
-// Device IO functions.
-// TODO:
-// Research - MapBuffers can improve performance when the device uses shared memory
-// but is more complex to understand. http://stackoverflow.com/questions/22057692/whats-the-difference-between-clenqueuemapbuffer-and-clenqueuewritebuffer
-// Intel graphics (and possibly AMD APUs) should use MapBuffers?
-// https://software.intel.com/en-us/articles/getting-the-most-from-opencl-12-how-to-increase-performance-by-minimizing-buffer-copies-on-intel-processor-graphics
-
void OpenclDevice::DoExec(function<void(Context*)>&& fn, int executor) {
fn(&ctx_);
}
-// NOTE: ASSUMES dst AND/OR src POINTERS CAN BE CAST TO cl::Buffer POINTERS!
+
void OpenclDevice::CopyToFrom(void* dst, const void* src, size_t nBytes,
CopyDirection direction, Context* ctx) {
// Pointers must be valid.
if (!dst || !src) return;
+
+ auto ocl_ctx = viennacl::ocl::get_context(ctx->vcl_ctx_id);
switch(direction) {
case kHostToDevice: {
- WriteToDevice(static_cast<cl::Buffer*>(dst), src, nBytes);
+ auto dst_handle = WrapHandle(static_cast<cl_mem>(dst), &ocl_ctx);
+ memory_write(dst_handle, 0, nBytes, src);
return;
}
case kDeviceToHost: {
- ReadFromDevice(dst, static_cast<const cl::Buffer*>(src), nBytes);
+ auto src_handle = WrapHandle((const cl_mem)src, &ocl_ctx);
+ memory_read(src_handle, 0, nBytes, dst);
return;
}
case kDeviceToDevice: {
- CopyDeviceBuffer(static_cast<cl::Buffer*>(dst), static_cast<const cl::Buffer*>(src), nBytes);
+ auto src_handle = WrapHandle((const cl_mem)src, &ocl_ctx);
+ auto dst_handle = WrapHandle(static_cast<cl_mem>(dst), &ocl_ctx);
+ memory_copy(src_handle, dst_handle, 0, 0, nBytes);
return;
}
default:
@@ -203,10 +153,7 @@ void OpenclDevice::CopyToFrom(void* dst, const void* src, size_t nBytes,
void* OpenclDevice::Malloc(int size) {
- cl_int status = CL_SUCCESS;
-
- cl::Buffer* buffer = new cl::Buffer(ocl_ctx, CL_MEM_READ_WRITE, size, nullptr, &status);
- OCL_CHECK(status, "Unable to allocate memory in OpenCL device.");
+ cl_mem buffer = memory_create(ocl::current_context(), size, nullptr);
return static_cast<void*>(buffer);
}
@@ -214,33 +161,8 @@ void* OpenclDevice::Malloc(int size) {
void OpenclDevice::Free(void* p) {
if (!p) return;
- cl::Buffer* buffer = static_cast<cl::Buffer*>(p);
- delete buffer;
-}
-
-
-void OpenclDevice::WriteToDevice(cl::Buffer* dst, const void* src, const size_t size) {
- cl_int status = CL_SUCCESS;
-
- status = cmdq.enqueueWriteBuffer(*dst, CL_TRUE, 0, size, src);
- OCL_CHECK(status, "Unable to write data to OpenCL device.");
-}
-
-
-void OpenclDevice::ReadFromDevice(void* dst, const cl::Buffer* src, const size_t size) {
- cl_int status = CL_SUCCESS;
-
- status = cmdq.enqueueReadBuffer(*src, CL_TRUE, 0, size, dst);
- OCL_CHECK(status, "Unable to read data from OpenCL device.");
-}
-
-
-// dst: cl::Buffer pointer src: cl::Buffer pointer
-void OpenclDevice::CopyDeviceBuffer(cl::Buffer* dst, const cl::Buffer* src, const size_t size) {
- cl_int status = CL_SUCCESS;
-
- status = cmdq.enqueueCopyBuffer(*src, *dst, 0, 0, size);
- OCL_CHECK(status, "Unable to copy buffer in OpenCL device.");
+ cl_mem buffer = static_cast<cl_mem>(p);
+ clReleaseMemObject(buffer);
}
} // namespace singa
http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/595302a3/src/core/tensor/tensor_math_opencl.cl
----------------------------------------------------------------------
diff --git a/src/core/tensor/tensor_math_opencl.cl b/src/core/tensor/tensor_math_opencl.cl
index f9cf96e..7b89970 100644
--- a/src/core/tensor/tensor_math_opencl.cl
+++ b/src/core/tensor/tensor_math_opencl.cl
@@ -24,7 +24,7 @@
// This reduction code is serial reduction modified from AMD's example.
// http://developer.amd.com/resources/documentation-articles/articles-whitepapers/opencl-optimization-case-study-simple-reductions/
__kernel
-void clkernel_abs(const int num, __global const float* in, __global float* out) {
+void clkernel_fabs(const int num, __global const float* in, __global float* out) {
const int i = get_global_id(0);
if (i >= num) return;
out[i] = fabs(in[i]);
@@ -462,7 +462,7 @@ void clkernel_crossentropy(const uint batchsize, const uint dim,
int truth_idx = t[gidx];
if (truth_idx <= 0) return;
- float prob_of_truth = p[gidx + truth_idx];
+ float prob_of_truth = p[gidx * dim + truth_idx];
loss[gidx] = -log(fmax(prob_of_truth, -FLT_MIN));
}
@@ -480,6 +480,21 @@ void clkernel_softmaxentropy(const uint batchsize, const uint dim,
}
+__kernel
+void clkernel_rowmax(const uint nrow, const uint ncol,
+ __global const float* in, __global float* out) {
+ const uint row_id = get_global_id(0);
+ if (row_id >= nrow) return;
+
+ float row_max_val = -FLT_MAX;
+ for (uint i = 0; i < ncol; i++) {
+ row_max_val = fmax(row_max_val, in[row_id * ncol + i]);
+ }
+
+ out[row_id] = row_max_val;
+}
+
+
// **************************************
// Matrix functions
// **************************************
http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/595302a3/src/core/tensor/tensor_math_opencl.h
----------------------------------------------------------------------
diff --git a/src/core/tensor/tensor_math_opencl.h b/src/core/tensor/tensor_math_opencl.h
index c289a56..c387031 100644
--- a/src/core/tensor/tensor_math_opencl.h
+++ b/src/core/tensor/tensor_math_opencl.h
@@ -19,17 +19,27 @@
#ifndef SINGA_CORE_TENSOR_TENSOR_MATH_OPENCL_H_
#ifdef USE_OPENCL
-#include <limits>
-#include "singa/utils/opencl_utils.h"
#include "tensor_math.h"
+#include "singa/utils/opencl_utils.h"
-namespace singa {
+#include <viennacl/scalar.hpp>
+#include <viennacl/vector.hpp>
+#include <viennacl/matrix.hpp>
+
+#include <viennacl/linalg/inner_prod.hpp>
+#include <viennacl/linalg/norm_2.hpp>
+#include <viennacl/linalg/sum.hpp>
+#include <viennacl/linalg/scalar_operations.hpp>
+#include <viennacl/linalg/vector_operations.hpp>
+#include <viennacl/linalg/matrix_operations.hpp>
-// Some forward declarations of utility functions that only exist here.
-void Transpose(const size_t nrow, const size_t ncol, cl::Buffer& in, cl::Buffer& out, Context* ctx);
-void DiagVec_Left(const size_t size, cl::Buffer& in, cl::Buffer& out, Context* ctx);
-void DiagVec_Right(const size_t size, cl::Buffer& in, cl::Buffer& out, Context* ctx);
+#include <viennacl/ocl/kernel.hpp>
+
+using viennacl::ocl::get_context;
+using viennacl::ocl::enqueue;
+
+namespace singa {
// **************************************
// Element-wise functions
@@ -37,436 +47,250 @@ void DiagVec_Right(const size_t size, cl::Buffer& in, cl::Buffer& out, Context*
template<>
void Abs<float, lang::Opencl>(const size_t num, const Block* in, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_abs";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer inbuf = *(static_cast<cl::Buffer*>(in->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, inbuf);
- kernel.setArg(2, outbuf);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ auto ocl_ctx = get_context(ctx->vcl_ctx_id);
+ auto kernel = ocl_ctx.get_kernel("tensor_math_opencl.cl", "clkernel_fabs");
+
+ viennacl::vector<float> v_in((const cl_mem)in->data(), num);
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), num);
+
+ v_out = v_in;
+ enqueue(kernel((cl_int)num, v_in, v_out));
}
template<>
void Add<float, lang::Opencl>(const size_t num, const Block* in, const float x, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_add_scalar";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer inbuf = *(static_cast<cl::Buffer*>(in->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, x);
- kernel.setArg(2, inbuf);
- kernel.setArg(3, outbuf);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ auto ocl_ctx = get_context(ctx->vcl_ctx_id);
+
+ viennacl::vector<float> x_in = viennacl::scalar_vector<float>(num, x, ocl_ctx);
+ viennacl::vector<float> v_in((const cl_mem)in->data(), num);
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), num);
+
+ v_out = v_in + x_in;
}
template<>
void Add<float, lang::Opencl>(const size_t num, const Block* in1, const Block* in2, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_add";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer in1buf = *(static_cast<cl::Buffer*>(in1->mutable_data()));
- cl::Buffer in2buf = *(static_cast<cl::Buffer*>(in2->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, in1buf);
- kernel.setArg(2, in2buf);
- kernel.setArg(3, outbuf);
+ viennacl::vector<float> v_in1((const cl_mem)in1->data(), num);
+ viennacl::vector<float> v_in2((const cl_mem)in2->data(), num);
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), num);
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ v_out = v_in1 + v_in2;
}
template<>
-void Clamp<float, lang::Opencl>(const size_t num, const float low, const float high, const Block* in, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_clamp";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer inbuf = *(static_cast<cl::Buffer*>(in->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, low);
- kernel.setArg(2, high);
- kernel.setArg(3, inbuf);
- kernel.setArg(4, outbuf);
+void Clamp<float, lang::Opencl>(const size_t num, const float low, const float high,
+ const Block* in, Block* out, Context* ctx) {
+ auto ocl_ctx = get_context(ctx->vcl_ctx_id);
+ auto kernel = ocl_ctx.get_kernel("tensor_math_opencl.cl", "clkernel_clamp");
+
+ viennacl::vector<float> v_in((const cl_mem)in->data(), num);
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), num);
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ enqueue(kernel((cl_int)num, low, high, v_in, v_out));
}
template<>
void Div<float, lang::Opencl>(const size_t num, const Block* in, const float x, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_divide_scalar_matx";
- auto kernel = ctx->kernels->at(kname);
+ auto ocl_ctx = get_context(ctx->vcl_ctx_id);
- cl::Buffer inbuf = *(static_cast<cl::Buffer*>(in->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
+ viennacl::vector<float> x_in = viennacl::scalar_vector<float>(num, x, ocl_ctx);
+ viennacl::vector<float> v_in((const cl_mem)in->data(), num);
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), num);
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, inbuf);
- kernel.setArg(2, x);
- kernel.setArg(3, outbuf);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ v_out = viennacl::linalg::element_div(v_in, x_in);
}
template<>
void Div<float, lang::Opencl>(const size_t num, const float x, const Block* in, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_divide_scalar_xmat";
- auto kernel = ctx->kernels->at(kname);
+ auto ocl_ctx = get_context(ctx->vcl_ctx_id);
- cl::Buffer inbuf = *(static_cast<cl::Buffer*>(in->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, x);
- kernel.setArg(2, inbuf);
- kernel.setArg(3, outbuf);
+ viennacl::vector<float> x_in = viennacl::scalar_vector<float>(num, x, ocl_ctx);
+ viennacl::vector<float> v_in((const cl_mem)in->data(), num);
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), num);
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ v_out = viennacl::linalg::element_div(x_in, v_in);
}
template<>
void Div<float, lang::Opencl>(const size_t num, const Block* in1, const Block* in2, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_divide";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer in1buf = *(static_cast<cl::Buffer*>(in1->mutable_data()));
- cl::Buffer in2buf = *(static_cast<cl::Buffer*>(in2->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, in1buf);
- kernel.setArg(2, in2buf);
- kernel.setArg(3, outbuf);
+ viennacl::vector<float> v_in1((const cl_mem)in1->data(), num);
+ viennacl::vector<float> v_in2((const cl_mem)in2->data(), num);
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), num);
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ v_out = viennacl::linalg::element_div(v_in1, v_in2);
}
template<>
void EltwiseMult<float, lang::Opencl>(const size_t num, const Block* in, const float x, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_eltmult_scalar";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer inbuf = *(static_cast<cl::Buffer*>(in->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, x);
- kernel.setArg(2, inbuf);
- kernel.setArg(3, outbuf);
+ auto ocl_ctx = get_context(ctx->vcl_ctx_id);
+
+ viennacl::vector<float> x_in = viennacl::scalar_vector<float>(num, x, ocl_ctx);
+ viennacl::vector<float> v_in((const cl_mem)in->data(), num);
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), num);
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ v_out = viennacl::linalg::element_prod(v_in, x_in);
}
template<>
void EltwiseMult<float, lang::Opencl>(const size_t num, const Block* in1, const Block* in2, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
+ viennacl::vector<float> v_in1((const cl_mem)in1->data(), num);
+ viennacl::vector<float> v_in2((const cl_mem)in2->data(), num);
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), num);
- std::string kname = "clkernel_eltmult";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer in1buf = *(static_cast<cl::Buffer*>(in1->mutable_data()));
- cl::Buffer in2buf = *(static_cast<cl::Buffer*>(in2->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, in1buf);
- kernel.setArg(2, in2buf);
- kernel.setArg(3, outbuf);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ v_out = viennacl::linalg::element_prod(v_in1, v_in2);
}
template<>
void Exp<float, lang::Opencl>(const size_t num, const Block* in, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_exp";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer inbuf = *(static_cast<cl::Buffer*>(in->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, inbuf);
- kernel.setArg(2, outbuf);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ viennacl::vector<float> v_in((const cl_mem)in->data(), num);
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), num);
+
+ v_out = viennacl::linalg::element_exp(v_in);
}
template<>
void LE<float, lang::Opencl>(const size_t num, const Block *in, const float x, Block *out, Context *ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_le";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer inbuf = *(static_cast<cl::Buffer*>(in->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, inbuf);
- kernel.setArg(2, x);
- kernel.setArg(3, outbuf);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ auto ocl_ctx = get_context(ctx->vcl_ctx_id);
+ auto kernel = ocl_ctx.get_kernel("tensor_math_opencl.cl", "clkernel_le");
+
+ viennacl::vector<float> in_buf((const cl_mem)in->data(), num);
+ viennacl::vector<float> out_buf(static_cast<cl_mem>(out->mutable_data()), num);
+
+ enqueue(kernel((cl_int)num, in_buf, x, out_buf));
}
template<>
void Log<float, lang::Opencl>(const size_t num, const Block* in, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_log";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer inbuf = *(static_cast<cl::Buffer*>(in->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, inbuf);
- kernel.setArg(2, outbuf);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ viennacl::vector<float> v_in((const cl_mem)in->data(), num);
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), num);
+
+ v_out = viennacl::linalg::element_log(v_in);
}
template<>
void LT<float, lang::Opencl>(const size_t num, const Block *in, const float x, Block *out, Context *ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_lt";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer inbuf = *(static_cast<cl::Buffer*>(in->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, inbuf);
- kernel.setArg(2, x);
- kernel.setArg(3, outbuf);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ auto ocl_ctx = get_context(ctx->vcl_ctx_id);
+ auto kernel = ocl_ctx.get_kernel("tensor_math_opencl.cl", "clkernel_lt");
+
+ viennacl::vector<float> in_buf((const cl_mem)in->data(), num);
+ viennacl::vector<float> out_buf(static_cast<cl_mem>(out->mutable_data()), num);
+
+ enqueue(kernel((cl_int)num, in_buf, x, out_buf));
}
template<>
void GE<float, lang::Opencl>(const size_t num, const Block *in, const float x, Block *out, Context *ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_ge";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer inbuf = *(static_cast<cl::Buffer*>(in->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, inbuf);
- kernel.setArg(2, x);
- kernel.setArg(3, outbuf);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ auto ocl_ctx = get_context(ctx->vcl_ctx_id);
+ auto kernel = ocl_ctx.get_kernel("tensor_math_opencl.cl", "clkernel_ge");
+
+ viennacl::vector<float> in_buf((const cl_mem)in->data(), num);
+ viennacl::vector<float> out_buf(static_cast<cl_mem>(out->mutable_data()), num);
+
+ enqueue(kernel((cl_int)num, in_buf, x, out_buf));
}
template<>
void GT<float, lang::Opencl>(const size_t num, const Block *in, const float x, Block *out, Context *ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_gt";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer inbuf = *(static_cast<cl::Buffer*>(in->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, inbuf);
- kernel.setArg(2, x);
- kernel.setArg(3, outbuf);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ auto ocl_ctx = get_context(ctx->vcl_ctx_id);
+ auto kernel = ocl_ctx.get_kernel("tensor_math_opencl.cl", "clkernel_gt");
+
+ viennacl::vector<float> in_buf((const cl_mem)in->data(), num);
+ viennacl::vector<float> out_buf(static_cast<cl_mem>(out->mutable_data()), num);
+
+ enqueue(kernel((cl_int)num, in_buf, x, out_buf));
}
template<>
void Pow<float, lang::Opencl>(const size_t num, const Block* in, float x, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_pow_scalar";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer inbuf = *(static_cast<cl::Buffer*>(in->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, x);
- kernel.setArg(2, inbuf);
- kernel.setArg(3, outbuf);
+ auto ocl_ctx = get_context(ctx->vcl_ctx_id);
+
+ viennacl::vector<float> x_in = viennacl::scalar_vector<float>(num, x, ocl_ctx);
+ viennacl::vector<float> v_in((const cl_mem)in->data(), num);
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), num);
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ v_out = viennacl::linalg::element_pow(v_in, x_in);
}
template<>
void Pow<float, lang::Opencl>(const size_t num, const Block* in1, const Block* in2, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_pow";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer in1buf = *(static_cast<cl::Buffer*>(in1->mutable_data()));
- cl::Buffer in2buf = *(static_cast<cl::Buffer*>(in2->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, in1buf);
- kernel.setArg(2, in2buf);
- kernel.setArg(3, outbuf);
+ viennacl::vector<float> v_in1((const cl_mem)in1->data(), num);
+ viennacl::vector<float> v_in2((const cl_mem)in2->data(), num);
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), num);
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ v_out = viennacl::linalg::element_pow(v_in1, v_in2);
}
template<>
void ReLU<float, lang::Opencl>(const size_t num, const Block* in, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_relu";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer inbuf = *(static_cast<cl::Buffer*>(in->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, inbuf);
- kernel.setArg(2, outbuf);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ auto ocl_ctx = get_context(ctx->vcl_ctx_id);
+ auto kernel = ocl_ctx.get_kernel("tensor_math_opencl.cl", "clkernel_relu");
+
+ viennacl::vector<float> in_buf((const cl_mem)in->data(), num);
+ viennacl::vector<float> out_buf(static_cast<cl_mem>(out->mutable_data()), num);
+
+ enqueue(kernel((cl_int)num, in_buf, out_buf));
}
+
template<>
void Set<float, lang::Opencl>(const size_t num, const float x, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
+ auto ocl_ctx = get_context(ctx->vcl_ctx_id);
- std::string kname = "clkernel_set";
- auto kernel = ctx->kernels->at(kname);
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), num);
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, x);
- kernel.setArg(2, outbuf);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ v_out = viennacl::scalar_vector<float>(num, x, ocl_ctx);
}
template<>
void Sigmoid<float, lang::Opencl>(const size_t num, const Block* in, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_sigmoid";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer inbuf = *(static_cast<cl::Buffer*>(in->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, inbuf);
- kernel.setArg(2, outbuf);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ auto ocl_ctx = get_context(ctx->vcl_ctx_id);
+
+ const viennacl::vector<float> zero = viennacl::zero_vector<float>(num, ocl_ctx);
+ const viennacl::vector<float> one = viennacl::scalar_vector<float>(num, 1.0f, ocl_ctx);
+
+ viennacl::vector<float> v_in((const cl_mem)in->data(), num);
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), num);
+
+ v_out = viennacl::linalg::element_div(one, viennacl::linalg::element_exp(zero - v_in) + one);
}
template<>
void Sign<float, lang::Opencl>(const size_t num, const Block* in, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_sign";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer inbuf = *(static_cast<cl::Buffer*>(in->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, inbuf);
- kernel.setArg(2, outbuf);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ auto ocl_ctx = get_context(ctx->vcl_ctx_id);
+ auto kernel = ocl_ctx.get_kernel("tensor_math_opencl.cl", "clkernel_abs");
+
+ viennacl::vector<float> in_buf((const cl_mem)in->data(), num);
+ viennacl::vector<float> out_buf(static_cast<cl_mem>(out->mutable_data()), num);
+
+ enqueue(kernel(num, in_buf, out_buf));
}
template<>
void Sqrt<float, lang::Opencl>(const size_t num, const Block* in, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_sqrt";
- auto kernel = ctx->kernels->at(kname);
+ viennacl::vector<float> v_in((const cl_mem)in->data(), num);
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), num);
- cl::Buffer inbuf = *(static_cast<cl::Buffer*>(in->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, inbuf);
- kernel.setArg(2, outbuf);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ v_out = viennacl::linalg::element_sqrt(v_in);
}
@@ -478,168 +302,85 @@ void Square<float, lang::Opencl>(const size_t num, const Block* in, Block* out,
template<>
void Sub<float, lang::Opencl>(const size_t num, const Block* in, const float x, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_subtract_scalar";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer inbuf = *(static_cast<cl::Buffer*>(in->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, inbuf);
- kernel.setArg(2, x);
- kernel.setArg(3, outbuf);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ Add<float, lang::Opencl>(num, in, -x, out, ctx);
}
template<>
void Sub<float, lang::Opencl>(const size_t num, const Block* in1, const Block* in2, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
+ viennacl::vector<float> v_in1((const cl_mem)in1->data(), num);
+ viennacl::vector<float> v_in2((const cl_mem)in2->data(), num);
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), num);
- std::string kname = "clkernel_subtract";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer in1buf = *(static_cast<cl::Buffer*>(in1->mutable_data()));
- cl::Buffer in2buf = *(static_cast<cl::Buffer*>(in2->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, in1buf);
- kernel.setArg(2, in2buf);
- kernel.setArg(3, outbuf);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ v_out = v_in1 - v_in2;
}
template<>
void Sum<float, lang::Opencl>(const size_t num, const Block* in, float* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
+ viennacl::vector<float> v_in((const cl_mem)in->data(), num);
- std::string kname = "clkernel_reduce";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer inbuf = *(static_cast<cl::Buffer*>(in->mutable_data()));
-
- size_t size = sizeof(float) * num;
- cl::Buffer outval(ctx->ocl_ctx, CL_MEM_WRITE_ONLY, size, nullptr, &status);
- OCL_CHECK(status, "Failed to create buffer!");
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, inbuf);
- kernel.setArg(2, outval);
- kernel.setArg(3, cl::Local(size));
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
-
- float* temp = new float[num];
- status = ctx->ocl_cmdq.enqueueReadBuffer(outval, CL_TRUE, 0, size, temp);
- OCL_CHECK(status, "Failed to read from buffer!");
- out[0] = temp[0];
- delete temp;
+ out[0] = viennacl::linalg::sum(v_in);
}
template<>
void Tanh<float, lang::Opencl>(const size_t num, const Block* in, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_tanh";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer inbuf = *(static_cast<cl::Buffer*>(in->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
+ viennacl::vector<float> v_in((const cl_mem)in->data(), num);
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), num);
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, inbuf);
- kernel.setArg(2, outbuf);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ v_out = viennacl::linalg::element_tanh(v_in);
}
// **************************************
// Random functions
// **************************************
-/// Seed value required for generating distributions.
-static unsigned int seed[4] = {0, 32, 42, 888};
/// Number of generation rounds used in the current algorithm.
static cl_uint rounds = 8;
template<>
void Bernoulli<float, lang::Opencl>(const size_t num, const float p, Block* out, Context *ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "PRNG_threefry4x32_bernoulli";
- auto kernel = ctx->kernels->at(kname);
+ auto ocl_ctx = get_context(ctx->vcl_ctx_id);
+ auto kernel = ocl_ctx.get_kernel("distribution.cl", "PRNG_threefry4x32_bernoulli");
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, outbuf);
- kernel.setArg(1, seed);
- kernel.setArg(2, 0.0f); // inf
- kernel.setArg(3, 1.0f); // sup
- kernel.setArg(4, p); // threshold
- kernel.setArg(5, rounds);
- kernel.setArg(6, cl_uint(num) / 4);
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), num);
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num/4));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ viennacl::ocl::packed_cl_uint seed = {0, 32, 42, 888};
+
+ enqueue(kernel(v_out, seed, 0.0f, 1.0f, p, rounds, cl_uint(num / 4)));
}
template<>
void Gaussian<float, lang::Opencl>(const size_t num, const float mean, const float std, Block* out, Context *ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "PRNG_threefry4x32_gaussian";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, outbuf);
- kernel.setArg(1, seed);
- kernel.setArg(2, mean); // E
- kernel.setArg(3, std); // V
- kernel.setArg(4, rounds);
- kernel.setArg(5, cl_uint(num) / 4);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num/4));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ auto ocl_ctx = get_context(ctx->vcl_ctx_id);
+ auto kernel = ocl_ctx.get_kernel("distribution.cl", "PRNG_threefry4x32_gaussian");
+
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), num);
+
+ viennacl::ocl::packed_cl_uint seed = {0, 32, 42, 888};
+
+ enqueue(kernel(v_out, seed, mean, std, rounds, cl_uint(num/4)));
}
template<>
void Uniform<float, lang::Opencl>(const size_t num, const float low, const float high, Block* out, Context *ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "PRNG_threefry4x32_uniform";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
+ auto ocl_ctx = get_context(ctx->vcl_ctx_id);
+ auto kernel = ocl_ctx.get_kernel("distribution.cl", "PRNG_threefry4x32_uniform");
- status = kernel.setArg(0, outbuf); OCL_CHECK(status, "kernel arg 0");
- status = kernel.setArg(1, seed); OCL_CHECK(status, "kernel arg 1");
- status = kernel.setArg(2, low); OCL_CHECK(status, "kernel arg 2");
- status = kernel.setArg(3, high); OCL_CHECK(status, "kernel arg 3");
- status = kernel.setArg(4, rounds); OCL_CHECK(status, "kernel arg 4");
- status = kernel.setArg(5, cl_uint(num) / 4); OCL_CHECK(status, "kernel arg 5");
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num/4));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ viennacl::ocl::packed_cl_uint seed = {0, 32, 42, 888};
+
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), num);
+
+ enqueue(kernel(v_out, seed, low, high, rounds, cl_uint(num/4)));
}
// *********************************************************
// BLAS functions, ref to http://docs.nvidia.com/cuda/cublas
// *********************************************************
-
+/*
template<>
void Amax<float, lang::Opencl>(const size_t num, const Block* in, size_t* out, Context* ctx) {
cl_int status = CL_SUCCESS;
@@ -699,7 +440,7 @@ void Amin<float, lang::Opencl>(const size_t num, const Block* in, size_t* out, C
delete temp;
}
-
+
template<>
void Asum<float, lang::Opencl>(const size_t num, const Block* in, float* out, Context* ctx) {
cl_int status = CL_SUCCESS;
@@ -727,256 +468,141 @@ void Asum<float, lang::Opencl>(const size_t num, const Block* in, float* out, Co
out[0] = temp[0];
delete temp;
}
-
-
+*/
+/// out = alpha * in + out
template<>
void Axpy<float, lang::Opencl>(const size_t num, const float alpha, const Block* in, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_axpy";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer inbuf = *(static_cast<cl::Buffer*>(in->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, alpha);
- kernel.setArg(2, inbuf);
- kernel.setArg(3, outbuf);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ viennacl::vector<float> inbuf((const cl_mem)in->data(), num);
+ viennacl::vector<float> outbuf(static_cast<cl_mem>(out->mutable_data()), num);
+
+ outbuf += alpha * inbuf;
}
-
+/// out = ||in||_2^2, i.e, L2 norm.
template<>
void Nrm2<float, lang::Opencl>(const size_t num, const Block* in, float* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_nrm2";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer inbuf = *(static_cast<cl::Buffer*>(in->mutable_data()));
-
- size_t size = sizeof(float) * num;
- cl::Buffer outval(ctx->ocl_ctx, CL_MEM_WRITE_ONLY, size, nullptr, &status);
- OCL_CHECK(status, "Failed to create buffer!");
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, inbuf);
- kernel.setArg(2, outval);
- kernel.setArg(3, cl::Local(sizeof(float) * (std::pow(2, num))));
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
-
- float* temp = new float[num];
- status = ctx->ocl_cmdq.enqueueReadBuffer(outval, CL_TRUE, 0, size, temp);
- OCL_CHECK(status, "Failed to read from buffer!");
- out[0] = temp[0];
- delete temp;
+ viennacl::vector<float> inbuf((const cl_mem)in->data(), num);
+
+ out[0] = viennacl::linalg::norm_2(inbuf);
}
template<>
void Scale<float, lang::Opencl>(const size_t num, const float x, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_scale";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, x);
- kernel.setArg(2, outbuf);
+ auto ocl_ctx = get_context(ctx->vcl_ctx_id);
+
+ viennacl::vector<float> x_in = viennacl::scalar_vector<float>(num, x, ocl_ctx);
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), num);
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ v_out = viennacl::linalg::element_prod(v_out, x_in);
}
template<>
void Dot<float, lang::Opencl>(const size_t num, const Block *in1, const Block *in2, float *out, Context *ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_dot";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer in1buf = *(static_cast<cl::Buffer*>(in1->mutable_data()));
- cl::Buffer in2buf = *(static_cast<cl::Buffer*>(in2->mutable_data()));
-
- size_t size = sizeof(float) * num;
- cl::Buffer outval(ctx->ocl_ctx, CL_MEM_WRITE_ONLY, size, nullptr, &status);
- OCL_CHECK(status, "Failed to create buffer!");
-
- kernel.setArg(0, (cl_int)num);
- kernel.setArg(1, in1buf);
- kernel.setArg(2, in2buf);
- kernel.setArg(3, outval);
- kernel.setArg(4, cl::Local(size));
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(num));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
-
- float* temp = new float[num];
- status = ctx->ocl_cmdq.enqueueReadBuffer(outval, CL_TRUE, 0, size, temp);
- OCL_CHECK(status, "Failed to read from buffer!");
- out[0] = temp[0];
- delete temp;
+ viennacl::vector<float> in1_buf((const cl_mem)in1->data(), num);
+ viennacl::vector<float> in2_buf((const cl_mem)in2->data(), num);
+
+ out[0] = viennacl::linalg::inner_prod(in1_buf, in2_buf);
}
-
+/// out = alpha * A * v + beta * out.
template<>
void GEMV<float, lang::Opencl>(bool trans, const size_t m, const size_t n, const float alpha,
const Block *A, const Block *v, const float beta, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_gemv";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer Abuf = *(static_cast<cl::Buffer*>(A->mutable_data()));
- cl::Buffer vbuf = *(static_cast<cl::Buffer*>(v->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)m);
- kernel.setArg(1, (cl_int)n);
- kernel.setArg(2, alpha);
- kernel.setArg(3, Abuf);
- kernel.setArg(4, vbuf);
- kernel.setArg(5, beta);
- kernel.setArg(6, outbuf);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(m, n));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+
+ viennacl::matrix<float> A_in((const cl_mem)A->data(), m, n);
+ viennacl::vector<float> v_in((const cl_mem)v->data(), trans ? m : n);
+ viennacl::vector<float> o_in(static_cast<cl_mem>(out->mutable_data()), trans ? n : m);
+
+ if (trans) viennacl::trans(A_in);
+
+ o_in *= beta;
+ o_in += alpha * viennacl::linalg::prod(A_in, v_in);
}
+/// multiply a matrix with a diagnoal matrix constructed using values from 'v'.
+/// if matrix_lef_side is true, do M*v; else do v*M
template<>
void DGMM<float, lang::Opencl>(bool side_right,
const size_t nrow, const size_t ncol,
const Block *M, const Block *v, Block *out, Context *ctx) {
- cl_int status = CL_SUCCESS;
-
- cl::Buffer Mbuf = *(static_cast<cl::Buffer*>(M->mutable_data()));
- cl::Buffer vbuf = *(static_cast<cl::Buffer*>(v->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
- std::string kname;
+ viennacl::matrix<float> M_buf((const cl_mem)M->data(), nrow, ncol);
+ viennacl::vector<float> v_buf((const cl_mem)v->data(), nrow);
+ viennacl::matrix<float> out_buf(static_cast<cl_mem>(out->mutable_data()), nrow, ncol);
+
+ auto diag = viennacl::diag(v_buf);
+
if (side_right) {
- DiagVec_Right(ncol, vbuf, vbuf, ctx);
- kname = "clkernel_dgmm_right";
+ out_buf = viennacl::linalg::prod(diag, M_buf);
} else {
- DiagVec_Left(nrow, vbuf, vbuf, ctx);
- kname = "clkernel_dgmm_left";
+ out_buf = viennacl::linalg::prod(M_buf, diag);
}
-
- auto kernel = ctx->kernels->at(kname);
-
- kernel.setArg(0, (cl_int)nrow);
- kernel.setArg(1, (cl_int)ncol);
- kernel.setArg(2, Mbuf);
- kernel.setArg(3, vbuf);
- kernel.setArg(4, outbuf);
- kernel.setArg(5, cl::Local(sizeof(float) * nrow * ncol));
-
- cl::NDRange global(nrow); // Only nrow because current implementation is 1 dimensional
-// cl::NDRange local();
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), global);
- OCL_CHECK(status, "Failed to enqueue kernel function!");
}
-
+/// C = alpha * A * B + beta * C.
template<>
void GEMM<float, lang::Opencl>(const bool transA, const bool transB,
const size_t nrowA, const size_t ncolB, const size_t ncolA,
const float alpha, const Block *A, const Block *B, const float beta,
Block *C, Context *ctx) {
- cl_int status = CL_SUCCESS;
- std::string kname = "clkernel_gemm";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer Abuf = *(static_cast<cl::Buffer*>(A->mutable_data()));
- cl::Buffer Bbuf = *(static_cast<cl::Buffer*>(B->mutable_data()));
- cl::Buffer Cbuf = *(static_cast<cl::Buffer*>(C->mutable_data()));
-
- // If matrix A needs to be transposed, do it.
- if (transA)
- Transpose(nrowA, ncolA, Abuf, Abuf, ctx);
-
- // If vector B needs to be transposed, do it.
- if (transB)
- Transpose(nrowA, ncolB, Bbuf, Bbuf, ctx);
-
- kernel.setArg(0, (cl_int)nrowA);
- kernel.setArg(1, (cl_int)ncolB);
- kernel.setArg(2, (cl_int)ncolA);
- kernel.setArg(3, alpha);
- kernel.setArg(4, Abuf);
- kernel.setArg(5, Bbuf);
- kernel.setArg(6, beta);
- kernel.setArg(7, Cbuf);
- kernel.setArg(8, cl::Local(sizeof(float) * nrowA * ncolB));
- kernel.setArg(9, cl::Local(sizeof(float) * nrowA * ncolB));
-
-// TODO: Try to make the work group size a power of 2 given an arbitrary matrix.
- cl::NDRange global(nrowA, ncolB);
- cl::NDRange local(nrowA, ncolB);
+ viennacl::matrix<float> A_buf((const cl_mem)A->data(), nrowA, ncolA);
+ viennacl::matrix<float> B_buf((const cl_mem)B->data(), ncolA, ncolB);
+ viennacl::matrix<float> C_buf(static_cast<cl_mem>(C->mutable_data()), nrowA, ncolB);
+
+ if (transA) viennacl::trans(A_buf);
+ if (transB) viennacl::trans(B_buf);
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), global, local);
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ C_buf *= beta;
+ C_buf += alpha * viennacl::linalg::prod(A_buf, B_buf);
}
+
template <>
void ComputeCrossEntropy<float, lang::Opencl>(const size_t batchsize, const size_t dim,
const Block *p, const Block *t, Block *loss,
Context *ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_crossentropy";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer pbuf = *(static_cast<cl::Buffer*>(p->mutable_data()));
- cl::Buffer tbuf = *(static_cast<cl::Buffer*>(t->mutable_data()));
- cl::Buffer lossbuf = *(static_cast<cl::Buffer*>(loss->mutable_data()));
-
- kernel.setArg(0, (cl_uint)batchsize);
- kernel.setArg(1, (cl_uint)dim);
- kernel.setArg(2, pbuf);
- kernel.setArg(3, tbuf);
- kernel.setArg(4, lossbuf);
-
- cl::NDRange global(batchsize);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), global);
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ auto ocl_ctx = get_context(ctx->vcl_ctx_id);
+ auto kernel = ocl_ctx.get_kernel("tensor_math_opencl.cl", "clkernel_crossentropy");
+
+ viennacl::vector<float> p_buf((const cl_mem)p->data(), batchsize);
+ viennacl::vector<float> t_buf((const cl_mem)t->data(), batchsize);
+ viennacl::vector<float> loss_buf(static_cast<cl_mem>(loss->mutable_data()), batchsize);
+
+ enqueue(kernel((cl_uint)batchsize, (cl_uint)dim, p_buf, t_buf, loss_buf));
}
+
template <>
void SoftmaxCrossEntropyBwd<float, lang::Opencl>(const size_t batchsize, const size_t dim,
const Block *p, const Block *t, Block *grad,
Context *ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_softmaxentropy";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer pbuf = *(static_cast<cl::Buffer*>(p->mutable_data()));
- cl::Buffer tbuf = *(static_cast<cl::Buffer*>(t->mutable_data()));
- cl::Buffer gradbuf = *(static_cast<cl::Buffer*>(grad->mutable_data()));
-
- kernel.setArg(0, (cl_uint)batchsize);
- kernel.setArg(1, (cl_uint)dim);
- kernel.setArg(2, pbuf);
- kernel.setArg(3, tbuf);
- kernel.setArg(4, gradbuf);
+ auto ocl_ctx = get_context(ctx->vcl_ctx_id);
+ auto kernel = ocl_ctx.get_kernel("tensor_math_opencl.cl", "clkernel_softmaxentropy");
+
+ viennacl::vector<float> p_buf((const cl_mem)p->data(), batchsize);
+ viennacl::vector<float> t_buf((const cl_mem)t->data(), batchsize);
+ viennacl::vector<float> grad_buf(static_cast<cl_mem>(grad->mutable_data()), batchsize);
+
+ enqueue(kernel((cl_uint)batchsize, (cl_uint)dim, p_buf, t_buf, grad_buf));
+}
- cl::NDRange global(batchsize);
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), global);
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+template<>
+void RowMax<float, lang::Opencl>(const size_t nrow, const size_t ncol,
+ const Block *in, Block *out, Context *ctx) {
+ auto ocl_ctx = get_context(ctx->vcl_ctx_id);
+ auto kernel = ocl_ctx.get_kernel("tensor_math_opencl.cl", "clkernel_rowmax");
+
+// kernel.global_work_size(0, nrow);
+
+ viennacl::matrix<float> in_buf((const cl_mem)in->data(), nrow, ncol);
+ viennacl::vector<float> outbuf(static_cast<cl_mem>(out->mutable_data()), nrow);
+
+ enqueue(kernel((cl_uint)nrow, (cl_uint)ncol, in_buf, outbuf));
}
// **************************************
@@ -985,129 +611,46 @@ void SoftmaxCrossEntropyBwd<float, lang::Opencl>(const size_t batchsize, const s
/*
template<>
void AddCol<float, lang::Opencl>(const size_t nrow, const size_t ncol, const Block* A, const Block* v, Block* out, Context* ctx) {
- std::string kname = "clkernel_addcol";
- auto kernel = ctx->kernels->at(kname);
- kernel.setArg(0, (cl_int)nrow);
- kernel.setArg(1, (cl_int)ncol);
- kernel.setArg(2, static_cast<const float*>(A->mutable_data()));
- kernel.setArg(3, static_cast<const float*>(v->mutable_data()));
- kernel.setArg(3, static_cast<float*>(out->mutable_data()));
- ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(nrow, ncol));
}
+
template<>
void AddRow<float, lang::Opencl>(const size_t nrow, const size_t ncol, const Block* A, const Block* v, Block* out, Context* ctx) {
- std::string kname = "clkernel_addrow";
- auto kernel = ctx->kernels->at(kname);
- kernel.setArg(0, (cl_int)nrow);
- kernel.setArg(1, (cl_int)ncol);
- kernel.setArg(2, static_cast<const float*>(A->mutable_data()));
- kernel.setArg(3, static_cast<const float*>(v->mutable_data()));
- kernel.setArg(3, static_cast<float*>(out->mutable_data()));
- ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(nrow, ncol));
}
+
template<>
void Outer<float, lang::Opencl>(const size_t m, const size_t n, const Block* lhs, const Block* rhs, Block* out, Context* ctx) {
- std::string kname = "clkernel_outerproduct";
- auto kernel = ctx->kernels->at(kname);
- kernel.setArg(0, (cl_int)m);
- kernel.setArg(1, (cl_int)n);
- kernel.setArg(2, static_cast<const float*>(lhs->data()));
- kernel.setArg(3, static_cast<const float*>(rhs->data()));
- kernel.setArg(4, static_cast<float*>(out->mutable_data()));
-
- ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(m, n));
+ viennacl::vector<float> lhs_in((const cl_mem)lhs->data(), m);
+ viennacl::vector<float> rhs_in((const cl_mem)rhs->data(), n);
+ viennacl::matrix<float> out_buf(static_cast<cl_mem>(out->mutable_data()), m, n);
+
+ out_buf = viennacl::linalg::outer_prod(lhs_in, rhs_in);
}
-template<>
-void SumColumns<float, lang::Opencl>(const size_t nrow, const size_t ncol, const Block* in, Block* out, Context* ctx) {
- std::string kname = "clkernel_sumcol";
- auto kernel = ctx->kernels->at(kname);
- kernel.setArg(0, (cl_int)nrow);
- kernel.setArg(1, (cl_int)ncol);
- kernel.setArg(2, static_cast<const float*>(in->mutable_data()));
- kernel.setArg(3, static_cast<float*>(out->mutable_data()));
- ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(nrow, ncol));
-}*/
-/*
template<>
-void SumRows<float, lang::Opencl>(const size_t nrow, const size_t ncol, const Block* in, Block* out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_sumrow";
- auto kernel = ctx->kernels->at(kname);
-
- cl::Buffer inbuf = *(static_cast<cl::Buffer*>(in->mutable_data()));
- cl::Buffer outbuf = *(static_cast<cl::Buffer*>(out->mutable_data()));
-
- kernel.setArg(0, (cl_int)nrow);
- kernel.setArg(1, (cl_int)ncol);
- kernel.setArg(2, inbuf);
- kernel.setArg(3, outbuf);
- kernel.setArg(4, cl::Local(sizeof(float) * nrow * ncol));
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(nrow, ncol));
-}
-*/
-
-
-#define BLOCK_DIM 16
-
-void Transpose(const size_t nrow, const size_t ncol, cl::Buffer& in, cl::Buffer& out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_transpose";
- auto kernel = ctx->kernels->at(kname);
-
- kernel.setArg(0, (cl_uint)nrow);
- kernel.setArg(1, (cl_uint)ncol);
- kernel.setArg(2, in);
- kernel.setArg(3, out);
- kernel.setArg(4, cl::Local((BLOCK_DIM + 1) * BLOCK_DIM));
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(nrow, ncol));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
-}
-
-#undef BLOCK_DIM
-
-
-/// This is a utility function that transforms a single-row vector into a diagonal matrix.
-/// For example, a vector of size n will become a matrix of size n*n where only the positions nx == ny will have values.
-void DiagVec_Left(const size_t size, cl::Buffer& in, cl::Buffer& out, Context* ctx) {
- cl_int status = CL_SUCCESS;
-
- std::string kname = "clkernel_diagvec_left";
- auto kernel = ctx->kernels->at(kname);
+void SumColumns<float, lang::Opencl>(const size_t nrow, const size_t ncol, const Block* in, Block* out, Context* ctx) {
+ viennacl::matrix<float> m_in((const cl_mem)in->data(), nrow, ncol);
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), nrow);
- kernel.setArg(0, (cl_uint)size);
- kernel.setArg(1, in);
- kernel.setArg(2, out);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(size));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+ v_out = viennacl::linalg::column_sum(m_in);
}
-void DiagVec_Right(const size_t size, cl::Buffer& in, cl::Buffer& out, Context* ctx) {
- cl_int status = CL_SUCCESS;
- std::string kname = "clkernel_diagvec_right";
- auto kernel = ctx->kernels->at(kname);
-
- kernel.setArg(0, (cl_uint)size);
- kernel.setArg(1, in);
- kernel.setArg(2, out);
-
- status = ctx->ocl_cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), cl::NDRange(size));
- OCL_CHECK(status, "Failed to enqueue kernel function!");
+template<>
+void SumRows<float, lang::Opencl>(const size_t nrow, const size_t ncol, const Block* in, Block* out, Context* ctx) {
+ viennacl::matrix<float> m_in((const cl_mem)in->data(), nrow, ncol);
+ viennacl::vector<float> v_out(static_cast<cl_mem>(out->mutable_data()), ncol);
+
+ v_out = viennacl::linalg::column_sum(m_in);
}
+*/
} // namespace singa
#endif // USE_OPENCL
-#endif // SINGA_CORE_TENSOR_TENSOR_MATH_OPENCL_H_
+#endif // SINGA_CORE_TENSOR_TENSOR_MATH_OPENCL_H_v_in + x;
http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/595302a3/src/utils/opencl_utils.cc
----------------------------------------------------------------------
diff --git a/src/utils/opencl_utils.cc b/src/utils/opencl_utils.cc
deleted file mode 100644
index e4fe69b..0000000
--- a/src/utils/opencl_utils.cc
+++ /dev/null
@@ -1,63 +0,0 @@
-/************************************************************
-*
-* Licensed to the Apache Software Foundation (ASF) under one
-* or more contributor license agreements. See the NOTICE file
-* distributed with this work for additional information
-* regarding copyright ownership. The ASF licenses this file
-* to you under the Apache License, Version 2.0 (the
-* "License"); you may not use this file except in compliance
-* with the License. You may obtain a copy of the License at
-*
-* http://www.apache.org/licenses/LICENSE-2.0
-*
-* Unless required by applicable law or agreed to in writing,
-* software distributed under the License is distributed on an
-* "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-* KIND, either express or implied. See the License for the
-* specific language governing permissions and limitations
-* under the License.
-*
-*************************************************************/
-
-#include "singa/utils/opencl_utils.h"
-
-#ifdef USE_OPENCL
-
-void PrintDeviceInfo(const cl::Device &dev) {
- cl_int status = CL_SUCCESS;
-
- LOG(INFO) << "\tDevice type: " << dev.getInfo<CL_DEVICE_TYPE>(&status);
- LOG(INFO) << "\tUnified memory: " << dev.getInfo<CL_DEVICE_HOST_UNIFIED_MEMORY>(&status);
- LOG(INFO) << "\tClock speed (MHz): " << dev.getInfo<CL_DEVICE_MAX_CLOCK_FREQUENCY>(&status);
- LOG(INFO) << "\tECC memory: " << dev.getInfo<CL_DEVICE_ERROR_CORRECTION_SUPPORT>(&status);
- LOG(INFO) << "\tLittle endian: " << dev.getInfo<CL_DEVICE_ENDIAN_LITTLE>(&status);
- LOG(INFO) << "\tCompute units: " << dev.getInfo<CL_DEVICE_MAX_COMPUTE_UNITS>(&status);
- LOG(INFO) << "\tMax work grp size: " << dev.getInfo<CL_DEVICE_MAX_WORK_GROUP_SIZE>(&status);
-//LOG(INFO) << "\tMax work item size: " << dev.getInfo<CL_DEVICE_MAX_WORK_ITEM_SIZES>(&status);
- LOG(INFO) << "\tMax item dimension: " << dev.getInfo<CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS>(&status);
- LOG(INFO) << "\tQueue properties: " << dev.getInfo<CL_DEVICE_QUEUE_PROPERTIES>(&status);
- LOG(INFO) << "\tExecution capabilities: " << dev.getInfo<CL_DEVICE_EXECUTION_CAPABILITIES>(&status);
- LOG(INFO) << "\tMax mem alloc size: " << dev.getInfo<CL_DEVICE_MAX_MEM_ALLOC_SIZE>(&status);
- LOG(INFO) << "\tGlobal mem size: " << dev.getInfo<CL_DEVICE_GLOBAL_MEM_SIZE>(&status);
- LOG(INFO) << "\tLocal mem size: " << dev.getInfo<CL_DEVICE_LOCAL_MEM_SIZE>(&status);
- LOG(INFO) << "\n";
-
- OCL_CHECK(status, "Failed to retrieve device information!");
-}
-
-
-void PrintPlatformInfo(const cl::Platform &p) {
- cl_int status = CL_SUCCESS;
-
- LOG(INFO) << "\tName: " << p.getInfo<CL_PLATFORM_NAME>(&status);
- LOG(INFO) << "\tProfile: " << p.getInfo<CL_PLATFORM_PROFILE>(&status);
- LOG(INFO) << "\tVersion: " << p.getInfo<CL_PLATFORM_VERSION>(&status);
- LOG(INFO) << "\tVendor: " << p.getInfo<CL_PLATFORM_VENDOR>(&status);
- LOG(INFO) << "\tExtensions: " << p.getInfo<CL_PLATFORM_EXTENSIONS>(&status);
- LOG(INFO) << "\n";
-
- OCL_CHECK(status, "Failed to retrieve platform information!");
-}
-
-
-#endif // USE_OPENCL