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>
This commit is contained in:
parent
c2a29fcc78
commit
f9ae670d05
13
BUILD.gn
13
BUILD.gn
@ -26,6 +26,7 @@ declare_args() {
|
||||
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_enable_tools) {
|
||||
]
|
||||
}
|
||||
}
|
||||
|
||||
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" ]
|
||||
}
|
||||
}
|
||||
}
|
||||
|
@ -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
|
||||
//
|
||||
//
|
||||
//
|
||||
|
118
tools/hello-opencl.cpp
Normal file
118
tools/hello-opencl.cpp
Normal file
@ -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;
|
||||
}
|
Loading…
Reference in New Issue
Block a user