Fixes to compile with most recent rust stable, as well as added features

Cargo.toml

@@ -27,4 +27,5 @@ path = "examples/platform/main.rs"
 
 [dependencies]
 log = "^0.3.1"
-libc = "^0.1.8"
+libc = "^0.1.8"
+lazy_static = "0.2.1"

examples/platform/main.rs

@@ -15,6 +15,9 @@ fn main() {
             println!("   Type: {}", device.device_type());
             println!("   Profile: {}", device.profile());
             println!("   Compute Units: {}", device.compute_units());
+            println!("   Global Memory Size: {} MB", device.global_mem_size() / (1024 * 1024));
+            println!("   Local Memory Size: {} MB", device.local_mem_size() / (1024 * 1024));
+            println!("   Max Alloc Size: {} MB", device.max_mem_alloc_size() / (1024 * 1024));
         }
     }
 }

src/cl.rs

@@ -1,7 +1,5 @@
 #![allow(non_camel_case_types, dead_code)]
 
-extern crate std;
-
 use libc;
 use std::fmt;
 

src/ext.rs

@@ -1,6 +1,5 @@
 #![allow(unused,
          unused_attributes,
-         raw_pointer_derive,
          non_camel_case_types,
          non_snake_case)]
 

src/hl.rs

@@ -1,14 +1,13 @@
 //! A higher level API.
-
 use libc;
-use std;
 use std::ffi::CString;
 use std::iter::repeat;
 use std::marker::PhantomData;
 use std::mem;
 use std::ptr;
 use std::string::String;
 use std::vec::Vec;
+use std::sync::Mutex;
 
 use cl;
 use cl::*;
@@ -22,6 +21,7 @@ pub enum DeviceType {
       CPU, GPU
 }
 
+
 fn convert_device_type(device: DeviceType) -> cl_device_type {
     match device {
         DeviceType::CPU => CL_DEVICE_TYPE_CPU,
@@ -131,7 +131,9 @@ impl Platform {
 // This mutex is used to work around weak OpenCL implementations.
 // On some implementations concurrent calls to clGetPlatformIDs
 // will cause the implantation to return invalid status.
-static mut platforms_mutex: std::sync::StaticMutex = std::sync::MUTEX_INIT;
+lazy_static! {
+    static ref platforms_mutex : Mutex<i32> = Mutex::new(0);
+}
 
 pub fn get_platforms() -> Vec<Platform>
 {
@@ -249,6 +251,47 @@ impl Device {
 		}
 	}
 
+    pub fn global_mem_size(&self) -> usize {
+        unsafe {
+            let mut ct: usize = 0;
+            let status = clGetDeviceInfo(
+                self.id,
+                CL_DEVICE_GLOBAL_MEM_SIZE,
+                16,
+                (&mut ct as *mut usize) as *mut libc::c_void,
+                ptr::null_mut());
+            check(status, "Could not get size of global memory.");
+            return ct;
+        }
+    }
+
+    pub fn local_mem_size(&self) -> usize {
+        unsafe {
+            let mut ct: usize = 0;
+            let status = clGetDeviceInfo(
+                self.id,
+                CL_DEVICE_LOCAL_MEM_SIZE,
+                16,
+                (&mut ct as *mut usize) as *mut libc::c_void,
+                ptr::null_mut());
+            check(status, "Could not get size of local memory.");
+            return ct;
+        }
+    }
+
+    pub fn max_mem_alloc_size(&self) -> usize {
+        unsafe {
+            let mut ct: usize = 0;
+            let status = clGetDeviceInfo(
+                self.id,
+                CL_DEVICE_MAX_MEM_ALLOC_SIZE,
+                16,
+                (&mut ct as *mut usize) as *mut libc::c_void,
+                ptr::null_mut());
+            check(status, "Could not get size of local memory.");
+            return ct;
+        }
+    }
 
     pub fn create_context(&self) -> Context
     {
@@ -651,6 +694,12 @@ impl Kernel {
     {
         set_kernel_arg(self, i as cl::cl_uint, x)
     }
+
+    pub fn alloc_local<T>(&self, i: usize, l: usize)
+    {
+        alloc_kernel_local::<T>(self, i as cl::cl_uint, 
+            l as libc::size_t)
+    }
 }
 
 pub fn create_kernel(program: &Program, kernel: & str) -> Kernel
@@ -721,6 +770,18 @@ pub fn set_kernel_arg<T: KernelArg>(kernel: & Kernel,
     }
 }
 
+pub fn alloc_kernel_local<T>(kernel: &Kernel,
+                       position: cl_uint,
+                       // size: libc::size_t,
+                       length: libc::size_t){
+    unsafe
+    {
+        let tsize = mem::size_of::<T>() as libc::size_t;
+        let ret = clSetKernelArg(kernel.kernel, position,
+                tsize*length, ptr::null());
+        check(ret, "Failed to set kernel arg!");
+    }
+}
 
 pub struct Event
 {

src/lib.rs

@@ -2,14 +2,15 @@
 #![allow(missing_copy_implementations)]
 #![allow(non_upper_case_globals)]
 
-#![feature(static_mutex)]
-
 //! OpenCL bindings for Rust.
 
 extern crate libc;
 #[macro_use]
 extern crate log;
 
+#[macro_use]
+extern crate lazy_static;
+
 #[link(name = "OpenCL", kind = "framework")]
 #[cfg(target_os = "macos")]
 extern { }

tests/test.rs

@@ -1,5 +1,3 @@
-#![feature(slice_bytes)]
-
 #[macro_use]
 extern crate log;
 
@@ -64,7 +62,7 @@ mod mem {
             unsafe {
                 let ptr = ptr as *const u8;
                 let src = slice::from_raw_parts(ptr, len);
-                slice::bytes::copy_memory(src, target);
+                target.clone_from_slice(src);
             }
         });
 
@@ -77,7 +75,7 @@ mod mem {
             unsafe {
                 let ptr = ptr as *mut u8;
                 let mut dst = slice::from_raw_parts_mut(ptr, len);
-                slice::bytes::copy_memory(src, dst);
+                dst.copy_from_slice(src);
             }
         })
     }
@@ -370,6 +368,34 @@ mod hl {
         })
     }
 
+    #[test]
+    fn memory_alloc_local()
+    {
+        let src = "__kernel void test(__global int *i, \
+                                    __local int *t) { \
+                   *t = *i; \
+                   *t += 1; \
+                   *i = *t; \
+                   }";
+        ::test_all_platforms_devices(&mut |device, ctx, queue| {
+            let prog = ctx.create_program_from_source(src);
+            prog.build(device).unwrap();
+
+            let k = prog.create_kernel("test");
+            let v = ctx.create_buffer_from(vec![1isize], CL_MEM_READ_WRITE);
+
+            k.set_arg(0, &v);
+            k.alloc_local::<isize>(1, 1);
+
+            queue.enqueue_async_kernel(&k, 1isize, None, ()).wait();
+
+            let v: Vec<isize> = queue.get(&v, ());
+
+            expect!(v[0], 2);
+        })
+
+    }
+
     #[test]
     fn event_get_times() {
         let src = "__kernel void test(__global int *i) { \