add hello-opencl
This is a little hello world for OpenCL.
Change-Id: I9502407dab75694a19e97832bcfe33f47cbd97e9
Reviewed-on: https://skia-review.googlesource.com/136223
Commit-Queue: Allan MacKinnon <allanmac@google.com>
Reviewed-by: Allan MacKinnon <allanmac@google.com>
diff --git a/BUILD.gn b/BUILD.gn
index 3fa9344..f3f170e 100644
--- a/BUILD.gn
+++ b/BUILD.gn
@@ -26,6 +26,7 @@
skia_use_libpng = true
skia_use_libwebp = !is_fuchsia
skia_use_lua = is_skia_dev_build && !is_ios
+ skia_use_opencl = false
skia_use_piex = !is_win
skia_use_zlib = true
skia_use_metal = false
@@ -2177,4 +2178,16 @@
]
}
}
+
+ if (skia_use_opencl) {
+ test_app("hello-opencl") {
+ sources = [
+ "src/compute/common/cl/assert_cl.c",
+ "src/compute/common/cl/find_cl.c",
+ "tools/hello-opencl.cpp",
+ ]
+ include_dirs = [ "src/compute/common" ]
+ libs = [ "OpenCL" ]
+ }
+ }
}
diff --git a/src/compute/common/macros.h b/src/compute/common/macros.h
index 35f658f..d91a000 100644
--- a/src/compute/common/macros.h
+++ b/src/compute/common/macros.h
@@ -37,8 +37,11 @@
//
//
-#define ALLOCA(n) _alloca(n)
-
+#if defined(_MSC_VER)
+ #define ALLOCA(n) _alloca(n)
+#else
+ #define ALLOCA(n) alloca(n)
+#endif
//
//
//
diff --git a/tools/hello-opencl.cpp b/tools/hello-opencl.cpp
new file mode 100644
index 0000000..7e57f0b
--- /dev/null
+++ b/tools/hello-opencl.cpp
@@ -0,0 +1,118 @@
+/*
+ * Copyright 2018 Google Inc.
+ *
+ * Use of this source code is governed by a BSD-style license that can be
+ * found in the LICENSE file.
+ */
+
+// This is a simple OpenCL Hello World that tests you have a functioning OpenCL setup.
+
+#include <CL/cl.hpp>
+#include <initializer_list>
+
+extern "C" {
+ #include "cl/assert_cl.h" // for cl(), cl_ok() macros
+ #include "cl/find_cl.h" // for clFindIdsByName
+}
+
+int main(int argc, char** argv) {
+ // Find any OpenCL platform+device with these substrings.
+ const char* platform_match = argc > 1 ? argv[1] : "";
+ const char* device_match = argc > 2 ? argv[2] : "";
+
+ cl_platform_id platform_id;
+ cl_device_id device_id;
+
+ char device_name[256];
+ size_t device_name_len;
+
+ // clFindIdsByName will narrate what it's doing when this is set.
+ bool verbose = true;
+
+ // The cl() macro prepends cl to its argument, calls it, and asserts that it succeeded,
+ // printing out the file, line, and somewhat readable version of the error code on failure.
+ //
+ // It's generally used to call OpenCL APIs, but here we've written clFindIdsByName to match
+ // the convention, as its error conditions are just going to be passed along from OpenCL.
+ cl(FindIdsByName(platform_match, device_match,
+ &platform_id, &device_id,
+ sizeof(device_name), device_name, &device_name_len,
+ verbose));
+
+ printf("picked %.*s\n", (int)device_name_len, device_name);
+
+ // Allan's code is all C using OpenCL's C API,
+ // but we can mix that freely with the C++ API found in cl.hpp.
+ // cl_ok() comes in handy here, which is cl() without the extra cl- prefix.
+
+ cl::Device device(device_id);
+
+ std::string name,
+ vendor,
+ extensions;
+ cl_ok(device.getInfo(CL_DEVICE_NAME, &name));
+ cl_ok(device.getInfo(CL_DEVICE_VENDOR, &vendor));
+ cl_ok(device.getInfo(CL_DEVICE_EXTENSIONS, &extensions));
+
+ printf("name %s, vendor %s, extensions:\n%s\n",
+ name.c_str(), vendor.c_str(), extensions.c_str());
+
+ std::vector<cl::Device> devices = { device };
+
+ // Some APIs can't return their cl_int error but might still fail,
+ // so they take a pointer. cl_ok() is really handy here too.
+ cl_int ok;
+ cl::Context ctx(devices,
+ nullptr/*optional cl_context_properties*/,
+ nullptr/*optional error reporting callback*/,
+ nullptr/*context arguement for error reporting callback*/,
+ &ok);
+ cl_ok(ok);
+
+ cl::Program program(ctx,
+ "__kernel void mul(__global const float* a, "
+ " __global const float* b, "
+ " __global float* dst) {"
+ " int i = get_global_id(0); "
+ " dst[i] = a[i] * b[i]; "
+ "} ",
+ /*and build now*/true,
+ &ok);
+ cl_ok(ok);
+
+ std::vector<float> a,b,p;
+ for (int i = 0; i < 1000; i++) {
+ a.push_back(+i);
+ b.push_back(-i);
+ p.push_back( 0);
+ }
+
+ cl::Buffer A(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR , sizeof(float)*a.size(), a.data()),
+ B(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR , sizeof(float)*b.size(), b.data()),
+ P(ctx, CL_MEM_WRITE_ONLY| CL_MEM_HOST_READ_ONLY, sizeof(float)*p.size());
+
+ cl::Kernel mul(program, "mul", &ok);
+ cl_ok(ok);
+ cl_ok(mul.setArg(0, A));
+ cl_ok(mul.setArg(1, B));
+ cl_ok(mul.setArg(2, P));
+
+ cl::CommandQueue queue(ctx, device);
+
+ cl_ok(queue.enqueueNDRangeKernel(mul, cl::NDRange(0) /*offset*/
+ , cl::NDRange(1000) /*size*/));
+
+ cl_ok(queue.enqueueReadBuffer(P, true/*block until read is done*/
+ , 0 /*offset in bytes*/
+ , sizeof(float)*p.size() /*size in bytes*/
+ , p.data()));
+
+ for (int i = 0; i < 1000; i++) {
+ if (p[i] != a[i]*b[i]) {
+ return 1;
+ }
+ }
+
+ printf("OpenCL sez: %g x %g = %g\n", a[42], b[42], p[42]);
+ return 0;
+}