Expose clCreateProgramWithIL functionality #61
Comments
|
I've been wanting a reason to start adding some 2.1 features so here you go: I'm not set up to test this so please let me know if it works (and possibly troubleshoot it if it doesn't). |
|
Wow thanks for adding this so quick! I've set everything up on my end to test, but now I'm getting I think you need to declare opencl_version_2_1 in your Cargo.toml? |
|
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. |
|
Alright, here's the ridiculousness you need in your Let me know if that works :) [EDIT]: Fixed ocl dependency ('path'->'git'). |
|
There must be some way to propagate that dependency with build scripts, I'll set that up eventually. |
|
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 Here's stdout: Here's my stacktrace: Here's the pretty-printed SPIR-V module I made if you're curious: |
|
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? |
|
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. |
|
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. |
|
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. |
|
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: Keep me updated on your progress :) |
|
I used the clang SPIR-V extension to build a SPIR-V binary out of this kernel: Producing a SPIR-V whose readable form is this: When I try to use |
|
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. |
|
Do you know which 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. |
|
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 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.: 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? :) |
|
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 |
|
Cool thank you. I'll play around with it. |
|
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. |
|
Futures, sweet! I'll give it a whirl when I get home :) |
|
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: 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])
} |
|
You've done well. I'm impressed. Are you interested at all in cortical learning algorithms? I have need of dreamers like you. |
|
Hmm you mean HTM? I looked into it before but couldn't find any enticing examples (like an MNIST model) |
|
Yeah HTM is a good example. |
I'd like to be able to use SPIR-V modules :)
The text was updated successfully, but these errors were encountered: