Expose clCreateProgramWithIL functionality

[cogciprocate]

Expose clCreateProgramWithIL functionality

[c0gent]

I've been wanting a reason to start adding some 2.1 features so here you go: Program::with_il.

// Add the following to the Cargo.toml of your crate(s):
[features]
opencl_version_2_1 = []

// You will also temporarily need to add the following until the new versions
// are published to crates.io (many new asynchronous processing features using
// futures-rs are also coming with the new versions so keep an eye out).
[replace]
"cl-sys:0.2.1" = { git = "https://github.com/cogciprocate/cl-sys" }
"ocl-core:0.3.2" = { git = "https://github.com/cogciprocate/ocl-core" }
"ocl-core:0.12.0" = { git = "https://github.com/cogciprocate/ocl" }


// There are a few different ways to create a program. You can make a builder
// first (you'd really only want to do this if passing the builder to
// something else such as a `ProQueBuilder`)...
let program_bldr = ocl::Program::builder()
    .il(il_byte_vec)
    .build(&context).unwrap();

// Or you can just create it directly:
let program = ocl::Program::new(il_byte_vec, &context);

I'm not set up to test this so please let me know if it works (and possibly troubleshoot it if it doesn't).

[tedsta]

Wow thanks for adding this so quick! I've set everything up on my end to test, but now I'm getting

`error: Package `ocl v0.12.0 (https://github.com/cogciprocate/ocl#0c830421)` does not have these features: `opencl_version_2_1`

I think you need to declare opencl_version_2_1 in your Cargo.toml?

[c0gent]

Yeah I'm not sure what the proper way to chain the features down through the libraries is... looking into it now.

My instructions were certainly wrong regardless though.

[c0gent]

Alright, here's the ridiculousness you need in your Cargo.toml:

[dependencies]
ocl = { git = "https://github.com/cogciprocate/ocl", features = ["opencl_version_2_1"] }
cl-sys = { git = "https://github.com/cogciprocate/cl-sys", features = ["opencl_version_2_1"] }
ocl-core = { git = "https://github.com/cogciprocate/ocl-core", features = ["opencl_version_2_1"] }

[replace]
"cl-sys:0.2.1" = { git = "https://github.com/cogciprocate/cl-sys", features = ["opencl_version_2_1"] }
"ocl-core:0.3.2" = { git = "https://github.com/cogciprocate/ocl-core", features = ["opencl_version_2_1"] }

Let me know if that works :)

[EDIT]: Fixed ocl dependency ('path'->'git').

[c0gent]

There must be some way to propagate that dependency with build scripts, I'll set that up eventually.

[tedsta]

My SPIR-V isn't valid :( I wrote it by hand. OpenCL driver is telling me I have an invalid program (error -44). I'll fiddle with it more tomorrow night. Thanks for adding this! I'll post back here if I get my module working.

For my purposes, this works. But I think it may not be bubbling the error all the way to the caller (me) properly. I call expect with a message on Program::with_il, and I never see that message.

Here's stdout:

thread 'tests::it_works' panicked at 'called `Result::unwrap()` on an `Err` value: 

################################ OPENCL ERROR ############################### 

Error executing function: clReleaseKernel  

Status error code: CL_INVALID_PROGRAM (-44)  

Please visit the following url for more information: 

https://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clReleaseKernel.html#errors  

############################################################################# 

Here's my stacktrace:

   1:     0x558ca0bdf28c - std::sys::imp::backtrace::tracing::imp::write::h9c41d2f69e5caabf
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libstd/sys/unix/backtrace/tracing/gcc_s.rs:42
   2:     0x558ca0be233e - std::panicking::default_hook::{{closure}}::hcc803c8663cda123
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libstd/panicking.rs:351
   3:     0x558ca0be1ee3 - std::panicking::default_hook::hd5bda4e453dfb4be
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libstd/panicking.rs:361
   4:     0x558ca0be27db - std::panicking::rust_panic_with_hook::hffbc74969c7b5d87
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libstd/panicking.rs:555
   5:     0x558ca0be2674 - std::panicking::begin_panic::hc4c5d184a1e3fb7c
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libstd/panicking.rs:517
   6:     0x558ca0be2599 - std::panicking::begin_panic_fmt::h34f5b320b0f94559
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libstd/panicking.rs:501
   7:     0x558ca0be2527 - rust_begin_unwind
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libstd/panicking.rs:477
   8:     0x558ca0c0d4fd - core::panicking::panic_fmt::h1016b85b51d1931f
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libcore/panicking.rs:69
   9:     0x558ca0b73bff - core::result::unwrap_failed::he6cb16427ffb30d0
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libcore/macros.rs:29
  10:     0x558ca0b6df05 - <core::result::Result<T, E>>::unwrap::h84b7d8df679838ac
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libcore/result.rs:745
  11:     0x558ca0b8aaf2 - <ocl_core::types::abs::Program as core::ops::Drop>::drop::h342c9efdf0fe5bfe
                        at /home/teddy/.cargo/git/checkouts/ocl-core-ea504f22c1693b63/11419c1/src/types/abs.rs:526
  12:     0x558ca0b73cd0 - drop::h158ebc264bf34ea0
  13:     0x558ca0b856a6 - ocl_core::functions::create_program_with_il::h720736d86fe35849
                        at /home/teddy/.cargo/git/checkouts/ocl-core-ea504f22c1693b63/11419c1/src/error.rs:88
                        at /home/teddy/.cargo/git/checkouts/ocl-core-ea504f22c1693b63/11419c1/src/functions.rs:66
                        at /home/teddy/.cargo/git/checkouts/ocl-core-ea504f22c1693b63/11419c1/src/functions.rs:1305
  14:     0x558ca0b156f6 - ocl::standard::program::Program::with_il::h18287c8ce92c5cd7
                        at /home/teddy/.cargo/git/checkouts/ocl-51b7c5264cbf134c/44c0ab7/src/standard/program.rs:394
  15:     0x558ca0af21e9 - tensor::tests::it_works::hdee36a75c8cfc8e0
                        at /home/teddy/code/rust/tensor-rs/src/lib.rs:101
  16:     0x558ca0b5360e - <F as test::FnBox<T>>::call_box::h0a98498b8201ff98
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libtest/lib.rs:1366
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libtest/lib.rs:140
  17:     0x558ca0be96ea - __rust_maybe_catch_panic
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libpanic_unwind/lib.rs:98
  18:     0x558ca0b47cca - std::panicking::try::do_call::h066c775502ca82e3
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libstd/panicking.rs:436
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libstd/panic.rs:361
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libtest/lib.rs:1311
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libstd/panic.rs:296
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libstd/panicking.rs:460
  19:     0x558ca0be96ea - __rust_maybe_catch_panic
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libpanic_unwind/lib.rs:98
  20:     0x558ca0b4e8f6 - <F as alloc::boxed::FnBox<A>>::call_box::h5882ce3c9c178b6d
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libstd/panicking.rs:436
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libstd/panic.rs:361
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libstd/thread/mod.rs:357
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/liballoc/boxed.rs:605
  21:     0x558ca0be1644 - std::sys::imp::thread::Thread::new::thread_start::h76badbf9b0ecaf58
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/liballoc/boxed.rs:615
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libstd/sys_common/thread.rs:21
                        at /buildslave/rust-buildbot/slave/nightly-dist-rustc-linux/build/src/libstd/sys/unix/thread.rs:84
  22:     0x7fc19e2c9183 - start_thread
  23:     0x7fc19dddf37c - clone
  24:                0x0 - <unknown>

Here's the pretty-printed SPIR-V module I made if you're curious:

; SPIR-V
; Version: 1.1
; Generator: Unknown
; Bound: 14
OpCapability Kernel
OpCapability Addresses
OpMemoryModel Logical OpenCL
OpEntryPoint Kernel %9 "test_kernel"
%1 = OpTypeVoid
%2 = OpTypeInt 64 0
%3 = OpTypeFloat 32
%4 = OpTypeVector %2 3
%5 = OpTypePointer CrossWorkgroup %3
%6 = OpTypeFunction %1 %5 %3
%7 = OpTypePointer UniformConstant %4
%8 = OpConstant  %2  0
%9 = OpFunction  %1  None %6
%10 = OpFunctionParameter  %5 
%11 = OpFunctionParameter  %3 
%12 = OpLabel
%13 = OpInBoundsPtrAccessChain  %5  %10 %8
OpStore %13 %11
OpReturn
OpFunctionEnd

[c0gent]

Alright I'm going to install the Intel 2.1 drivers and play around with it. Do you mind if I create a 2.1-specific example which includes your SPIR-V code?

[tedsta]

Don't mind a bit. I'll paste the code I have (uses the rspirv crate to generate the spir-v binary) tonight. Unfortunately the code is on my laptop at home and I am at work right now.

[c0gent]

Well after an hour of trying to get the 2.1 drivers to even install and link properly on my Windows laptop (bleh), I finally got this example to give an elaborate error message. I think it might be working now but I don't have time to find any other binary to test it with right now and I've reached the limits of my SPIR-V knowledge. Let me know how it goes for you.

[tedsta]

Here's what I have so far

I realize now I think I am building the IL byte array with the wrong endian-ness. I'll tinker more when I get home tonight.

[c0gent]

I forgot to mention that the signature to build programs has changed... It will also be changing again for v0.13.0 so just use the builder as shown here for now and you should be fine:

let program = Program::builder()
        .devices(device)
        .il(il_src)
        .build(&context).unwrap();

Keep me updated on your progress :)

[tedsta]

I used the clang SPIR-V extension to build a SPIR-V binary out of this kernel:

__kernel void multiply(__global const float *a,
                       __global const float *b,
                       __global float *c) {
    uintptr_t i = get_global_id(0);
    c[i] = a[i] * b[i];
}

Producing a SPIR-V whose readable form is this:

OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
%1 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical32 OpenCL
OpEntryPoint Kernel %10 "multiply"
OpSource OpenCL_C 200000
OpName %5 "__spirv_BuiltInGlobalInvocationId"
OpName %11 "a"
OpName %12 "b"
OpName %13 "c"
OpName %14 "entry"
OpName %16 "call"
OpName %17 "arrayidx"
OpName %19 "arrayidx1"
OpName %21 "mul"
OpName %22 "arrayidx2"
OpDecorate %23 FuncParamAttr NoWrite
%23 = OpDecorationGroup
OpDecorate %5 BuiltIn GlobalInvocationId
OpDecorate %5 Constant
OpDecorate %5 LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
OpGroupDecorate %23 %11 %12
%2 = OpTypeInt 32 0
%3 = OpTypeVector %2 3
%4 = OpTypePointer UniformConstant %3
%6 = OpTypeVoid
%7 = OpTypeFloat 32
%8 = OpTypePointer CrossWorkgroup %7
%9 = OpTypeFunction %6 %8 %8 %8
%5 = OpVariable  %4  UniformConstant
%10 = OpFunction  %6  None %9
%11 = OpFunctionParameter  %8 
%12 = OpFunctionParameter  %8 
%13 = OpFunctionParameter  %8 
%14 = OpLabel
%15 = OpLoad  %3  %5
%16 = OpCompositeExtract  %2  %15 0
%17 = OpInBoundsPtrAccessChain  %8  %11 %16
%18 = OpLoad  %7  %17 Aligned 4
%19 = OpInBoundsPtrAccessChain  %8  %12 %16
%20 = OpLoad  %7  %19 Aligned 4
%21 = OpFMul  %7  %18 %20
%22 = OpInBoundsPtrAccessChain  %8  %13 %16
OpStore %22 %21 Aligned 4
OpReturn
OpFunctionEnd

When I try to use libintelocl_2_1.so (from intel 2.1 experimental cpu only driver, sym linked as libOpenCL.so), I get an undefined reference to clEnqueueReleaseGLObjects. When I use opencl-1.2-base's libOpenCL.so from the same distribution (intel_sdk_for_opencl_2016_ubuntu_6.3.0.1904_x64), it builds and runs but I get "Invalid Kernel". That sounds better than "Invalid program" - progress!

[c0gent]

Perhaps I should put the OpenGL related stuff behind a feature gate. Seems odd though that I did not have a problem linking with the Intel driver. Perhaps I have a different version. I have the Feb 5 version if I remember correctly. I'll check later and try to run that kernel as well.

[tedsta]

Do you know which libOpenCL.so you are using? I admit I'm a little confused about these things. If we are able to choose which OpenCL implementation we use at runtime, shouldn't there just be a generic "libOpenCL finder" library that finds all the installed implementations and dynamically loads the selected one at runtime?

Like I said, when I link against the opencl-1.2-base libOpenCL.so, it seems to work. Because in the code I search for and use the experimental 2.1 driver.

[c0gent]

Yes I can understand your confusion. Here's how it works... OpenCL uses what's called an ICD loader which just stands for installable client driver loader. All that does is act as a front end and loads every driver listed in /etc/OpenCL/vendors (on Linux). It's what lets you choose between different platforms at runtime.

When you install any OpenCL drivers for any vendor, they will automatically install an ICD and it will be hooked up correctly. This is how you should leave things by default. Importantly, when using the ICD you have to link dynamically (which is what ocl does).

A side note if you're curious: a library or binary also has the option of linking directly to a driver statically. If you wanted to do that you or I would have to configure cl-sys properly before compiling for it to work though. When you mentioned that you were getting an error for an undefined reference, it was probably because you put a version of the driver in place (using symbolic links) of the ICD loader which had a different ABI meant for linking statically. Don't worry about any of that.

So...

I assume you're selecting your platform similarly to the way I do in the example I linked above, i.e.:

static PLATFORM_NAME: &'static str = "Experimental OpenCL 2.1 CPU Only Platform";

...

let platform = Platform::list().into_iter().find(|plat| plat.name() == PLATFORM_NAME)
        .unwrap_or(Platform::default());

If so, you're loading the correct platform from the ICD and everything is fine. Leave it as it is :)

Alright now here's my question to you. I need to put some working SPIR-V in binary form into the example so that I and others can test it. Do you have or know where I can find a known working 'hello world' type program binary? Do you perhaps have a link to the binary for the code you posted above? Or... Is there an easy way to compile it that will take me less than 10 minutes to download and figure out? :)

[tedsta]

Thanks for the explanation! It'd be much more intuitive if there was just one ICD loader maintained for each platform... But oh well.

Here's a link to the binary I produced with the multiply kernel I showed above: https://dl.dropboxusercontent.com/u/17256312/multiply.spirv

I followed the instructions here to build a clang that can produce spir-v binaries from OpenCL C: https://github.com/KhronosGroup/SPIR/tree/spirv-1.0

EDIT: Note I produced the binary with -triple=spir-unknown-unknown, not -triple=spir64-unknown-unknown

[c0gent]

Cool thank you. I'll play around with it.

[c0gent]

I just pushed a ton of new changes to master. Lots of new stuff and lots of breaking changes.

Since you are working off of it let me know if you run into any issues that aren't covered in the change log or that need more explaining.

[tedsta]

Futures, sweet! I'll give it a whirl when I get home :)

[tedsta]

I finally got a SPIR-V kernel working!!! :D

I needed to generate a 64 bit spir-v binary. You can find the one I'm using here: https://www.dropbox.com/s/81f8cqyp3zdhqr2/multiply64.spirv

It was generated with this command:

./clang -cc1 -emit-spirv -triple spir64-unknown-unknown -cl-std=CL2.0 -include opencl.h -x cl -o multiply.spirv multiply.cl

where multiply.cl is

_kernel void multiply(__global const float *a,
                       __global const float *b,
                       __global float *c) {
    uintptr_t i = get_global_id(0);
    c[i] = a[i] * b[i];
}

And here's my code:

extern crate ocl;

use ocl::{Platform, Device, Context, Queue, Buffer, Program, Kernel, Event, EventList};

pub fn find_platform() -> Option<Platform> {
    let platform_name = "Experimental OpenCL 2.1 CPU Only Platform";

    for platform in Platform::list() {
        if platform.name() == platform_name {
            return Some(platform);
        }
    }

    None
}

fn main() {
    use std::io::Read;
    use std::fs::File;
    use ocl::{self, Platform, Device, Context, Queue, Buffer, Program, Kernel, Event, EventList};
    use ::{find_platform, build_spirv_module};

    let platform = find_platform().unwrap();
    assert!(platform.name() == "Experimental OpenCL 2.1 CPU Only Platform");

    // Get first (and only) device
    let device = Device::first(platform);

    // Build context using the first device
    let context = Context::builder()
        .platform(platform)
        .devices(device)
        .build().expect("Failed to create context");

    let src = r#"
        __kernel void multiply(__global const float *a,
                               __global const float *b,
                               __global float *c) {
            int i = get_global_id(0);
            c[i] = a[i] * b[i];
        }
    "#;

    let mut f = File::open("multiply64.spirv").expect("Failed to open spir-v module");
    let mut il_byte_vec = Vec::new();
    f.read_to_end(&mut il_byte_vec).expect("Failed to read spir-v module");

    let queue = Queue::new(&context, device, Some(ocl::core::QUEUE_PROFILING_ENABLE)).expect("Failed to create queue");

    let dims = [10];
    let a_host = vec![1.0f32; dims[0]];
    let a = Buffer::<f32>::builder().queue(queue.clone())
                                    .dims(&dims)
                                    .flags(ocl::flags::MEM_READ_ONLY | ocl::flags::MEM_COPY_HOST_PTR)
                                    .host_data(&a_host)
                                    .build().expect("Failed to create buffer a");

    let b_host = vec![2.0f32; dims[0]];
    let b = Buffer::<f32>::builder().queue(queue.clone())
                                    .dims(&dims)
                                    .flags(ocl::flags::MEM_READ_ONLY | ocl::flags::MEM_COPY_HOST_PTR)
                                    .host_data(&b_host)
                                    .build().expect("Failed to create buffer b");
    
    let c = Buffer::<f32>::builder().queue(queue.clone())
                                    .dims(&dims)
                                    .flags(ocl::flags::MEM_READ_WRITE)
                                    .build().expect("Failed to create buffer c");
    let mut c_host = vec![0.0; dims[0]];

    let program = Program::builder()
        .devices(device)
        .il(il_byte_vec)
        //.src(src)
        .build(&context).expect("Failed to build program from SPIR-V module");
    let kernel = Kernel::new("multiply", &program).expect("Failed to create kernel")
        .queue(queue.clone())
        .gws(&dims)
        .arg_buf(&a)
        .arg_buf(&b)
        .arg_buf(&c);

    let mut event_list = EventList::new();
    kernel.cmd().enew(&mut event_list).enq().unwrap();
    event_list.wait_for().unwrap();

    let mut event = Event::empty();
    c.cmd().read(&mut c_host).enew(&mut event).enq().unwrap();
    event.wait_for().unwrap();

    assert!(c_host == [2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0])
}

[c0gent]

You've done well. I'm impressed.

Are you interested at all in cortical learning algorithms?

I have need of dreamers like you.

[tedsta]

Hmm you mean HTM? I looked into it before but couldn't find any enticing examples (like an MNIST model)

[c0gent]

Yeah HTM is a good example.