Mike Klein | f9ae670 | 2018-06-20 14:05:05 -0400 | [diff] [blame] | 1 | /* |
| 2 | * Copyright 2018 Google Inc. |
| 3 | * |
| 4 | * Use of this source code is governed by a BSD-style license that can be |
| 5 | * found in the LICENSE file. |
| 6 | */ |
| 7 | |
| 8 | // This is a simple OpenCL Hello World that tests you have a functioning OpenCL setup. |
| 9 | |
Mike Klein | 8a1f15d | 2019-02-11 11:59:41 -0500 | [diff] [blame] | 10 | #include "cl.hpp" |
Mike Klein | f9ae670 | 2018-06-20 14:05:05 -0400 | [diff] [blame] | 11 | #include <initializer_list> |
Mike Klein | 8a1f15d | 2019-02-11 11:59:41 -0500 | [diff] [blame] | 12 | #include <stdio.h> |
| 13 | #include <stdlib.h> |
| 14 | #include <string> |
| 15 | #include <vector> |
Mike Klein | f9ae670 | 2018-06-20 14:05:05 -0400 | [diff] [blame] | 16 | |
Mike Klein | 8a1f15d | 2019-02-11 11:59:41 -0500 | [diff] [blame] | 17 | static inline void assert_cl(cl_int rc, const char* file, int line) { |
| 18 | if (rc != CL_SUCCESS) { |
| 19 | fprintf(stderr, "%s:%d, got OpenCL error code %d\n", file,line,rc); |
| 20 | exit(1); |
| 21 | } |
Mike Klein | f9ae670 | 2018-06-20 14:05:05 -0400 | [diff] [blame] | 22 | } |
Mike Klein | 8a1f15d | 2019-02-11 11:59:41 -0500 | [diff] [blame] | 23 | #define cl_ok(err) assert_cl(err, __FILE__, __LINE__) |
Mike Klein | f9ae670 | 2018-06-20 14:05:05 -0400 | [diff] [blame] | 24 | |
Mike Klein | 640654d | 2019-02-11 14:59:56 -0500 | [diff] [blame] | 25 | int main(int, char**) { |
| 26 | std::vector<cl::Platform> platforms; |
| 27 | cl_ok(cl::Platform::get(&platforms)); |
Mike Klein | f9ae670 | 2018-06-20 14:05:05 -0400 | [diff] [blame] | 28 | |
Mike Klein | 640654d | 2019-02-11 14:59:56 -0500 | [diff] [blame] | 29 | std::vector<cl::Device> devices; |
| 30 | for (cl::Platform platform : platforms) { |
| 31 | std::vector<cl::Device> platform_devices; |
| 32 | cl_ok(platform.getDevices(CL_DEVICE_TYPE_ALL, &platform_devices)); |
| 33 | devices.insert(devices.end(), platform_devices.begin(), platform_devices.end()); |
Mike Klein | 8a1f15d | 2019-02-11 11:59:41 -0500 | [diff] [blame] | 34 | } |
Mike Klein | f9ae670 | 2018-06-20 14:05:05 -0400 | [diff] [blame] | 35 | |
Mike Klein | 640654d | 2019-02-11 14:59:56 -0500 | [diff] [blame] | 36 | if (devices.empty()) { |
| 37 | fprintf(stderr, "No OpenCL devices available. :(\n"); |
| 38 | return 1; |
Mike Klein | 8a1f15d | 2019-02-11 11:59:41 -0500 | [diff] [blame] | 39 | } |
Mike Klein | f9ae670 | 2018-06-20 14:05:05 -0400 | [diff] [blame] | 40 | |
Mike Klein | 640654d | 2019-02-11 14:59:56 -0500 | [diff] [blame] | 41 | // To keep things simple we'll only create single-device cl::Contexts. |
| 42 | for (cl::Device device : devices) { |
| 43 | std::string name, |
| 44 | version, |
| 45 | driver, |
| 46 | vendor, |
| 47 | extensions; |
| 48 | cl_ok(device.getInfo(CL_DEVICE_NAME, &name)); |
| 49 | cl_ok(device.getInfo(CL_DEVICE_VERSION, &version)); |
| 50 | cl_ok(device.getInfo(CL_DEVICE_VENDOR, &vendor)); |
| 51 | cl_ok(device.getInfo(CL_DEVICE_EXTENSIONS, &extensions)); |
| 52 | cl_ok(device.getInfo(CL_DRIVER_VERSION, &driver)); |
Mike Klein | f9ae670 | 2018-06-20 14:05:05 -0400 | [diff] [blame] | 53 | |
Mike Klein | 640654d | 2019-02-11 14:59:56 -0500 | [diff] [blame] | 54 | fprintf(stdout, "Using %s%s, vendor %s, version %s, extensions:\n%s\n", |
| 55 | version.c_str(), name.c_str(), vendor.c_str(), driver.c_str(), extensions.c_str()); |
Mike Klein | f9ae670 | 2018-06-20 14:05:05 -0400 | [diff] [blame] | 56 | |
Mike Klein | 640654d | 2019-02-11 14:59:56 -0500 | [diff] [blame] | 57 | std::vector<cl::Device> devices = { device }; |
Mike Klein | f9ae670 | 2018-06-20 14:05:05 -0400 | [diff] [blame] | 58 | |
Mike Klein | 640654d | 2019-02-11 14:59:56 -0500 | [diff] [blame] | 59 | // Some APIs can't return their cl_int error but might still fail, |
| 60 | // so they take a pointer. cl_ok() is really handy here too. |
| 61 | cl_int ok; |
| 62 | cl::Context ctx(devices, |
| 63 | nullptr/*optional cl_context_properties*/, |
| 64 | nullptr/*optional error reporting callback*/, |
| 65 | nullptr/*context argument for error reporting callback*/, |
Mike Klein | f9ae670 | 2018-06-20 14:05:05 -0400 | [diff] [blame] | 66 | &ok); |
Mike Klein | 640654d | 2019-02-11 14:59:56 -0500 | [diff] [blame] | 67 | cl_ok(ok); |
Mike Klein | f9ae670 | 2018-06-20 14:05:05 -0400 | [diff] [blame] | 68 | |
Mike Klein | 640654d | 2019-02-11 14:59:56 -0500 | [diff] [blame] | 69 | cl::Program program(ctx, |
| 70 | "__kernel void mul(__global const float* a, " |
| 71 | " __global const float* b, " |
| 72 | " __global float* dst) {" |
| 73 | " int i = get_global_id(0); " |
| 74 | " dst[i] = a[i] * b[i]; " |
| 75 | "} ", |
| 76 | /*and build now*/true, |
| 77 | &ok); |
| 78 | cl_ok(ok); |
Mike Klein | f9ae670 | 2018-06-20 14:05:05 -0400 | [diff] [blame] | 79 | |
Mike Klein | 640654d | 2019-02-11 14:59:56 -0500 | [diff] [blame] | 80 | std::vector<float> a,b,p; |
| 81 | for (int i = 0; i < 1000; i++) { |
| 82 | a.push_back(+i); |
| 83 | b.push_back(-i); |
| 84 | p.push_back( 0); |
| 85 | } |
Mike Klein | f9ae670 | 2018-06-20 14:05:05 -0400 | [diff] [blame] | 86 | |
Mike Klein | 640654d | 2019-02-11 14:59:56 -0500 | [diff] [blame] | 87 | cl::Buffer |
| 88 | A(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR , sizeof(float)*a.size(), a.data()), |
| 89 | B(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR , sizeof(float)*b.size(), b.data()), |
| 90 | P(ctx, CL_MEM_WRITE_ONLY| CL_MEM_HOST_READ_ONLY, sizeof(float)*p.size()); |
Mike Klein | f9ae670 | 2018-06-20 14:05:05 -0400 | [diff] [blame] | 91 | |
Mike Klein | 640654d | 2019-02-11 14:59:56 -0500 | [diff] [blame] | 92 | cl::Kernel mul(program, "mul", &ok); |
| 93 | cl_ok(ok); |
| 94 | cl_ok(mul.setArg(0, A)); |
| 95 | cl_ok(mul.setArg(1, B)); |
| 96 | cl_ok(mul.setArg(2, P)); |
Mike Klein | f9ae670 | 2018-06-20 14:05:05 -0400 | [diff] [blame] | 97 | |
Mike Klein | 640654d | 2019-02-11 14:59:56 -0500 | [diff] [blame] | 98 | cl::CommandQueue queue(ctx, device); |
Mike Klein | f9ae670 | 2018-06-20 14:05:05 -0400 | [diff] [blame] | 99 | |
Mike Klein | 640654d | 2019-02-11 14:59:56 -0500 | [diff] [blame] | 100 | cl_ok(queue.enqueueNDRangeKernel(mul, cl::NDRange(0) /*offset*/ |
| 101 | , cl::NDRange(1000) /*size*/)); |
Mike Klein | f9ae670 | 2018-06-20 14:05:05 -0400 | [diff] [blame] | 102 | |
Mike Klein | 640654d | 2019-02-11 14:59:56 -0500 | [diff] [blame] | 103 | cl_ok(queue.enqueueReadBuffer(P, true/*block until read is done*/ |
| 104 | , 0 /*offset in bytes*/ |
| 105 | , sizeof(float)*p.size() /*size in bytes*/ |
| 106 | , p.data())); |
| 107 | |
| 108 | fprintf(stdout, "OpenCL sez: %g x %g = %g\n", a[42], b[42], p[42]); |
| 109 | for (int i = 0; i < 1000; i++) { |
| 110 | if (p[i] != a[i]*b[i]) { |
| 111 | return 1; |
| 112 | } |
Mike Klein | f9ae670 | 2018-06-20 14:05:05 -0400 | [diff] [blame] | 113 | } |
| 114 | } |
| 115 | |
Mike Klein | f9ae670 | 2018-06-20 14:05:05 -0400 | [diff] [blame] | 116 | return 0; |
| 117 | } |