Skip to content

Commit fb96c4d

Browse files
tedstatedsta
committed
Support global_work_offset
1 parent aa6465a commit fb96c4d

File tree

4 files changed

+83
-18
lines changed

4 files changed

+83
-18
lines changed

‎examples/demo/main.rs

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -32,7 +32,7 @@ fn main()
3232
kernel.set_arg(1, &b);
3333
kernel.set_arg(2, &c);
3434

35-
let event = queue.enqueue_async_kernel(&kernel, vec_a.len(), None, ());
35+
let event = queue.enqueue_async_kernel(&kernel, None, vec_a.len(), None, ());
3636

3737
let vec_c: Vec<isize> = queue.get(&c, &event);
3838

‎src/array.rs

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -180,6 +180,14 @@ impl<T: Clone> Array2D<T> {
180180
}
181181
}
182182

183+
pub fn width(&self) -> usize {
184+
self.width
185+
}
186+
187+
pub fn height(&self) -> usize {
188+
self.height
189+
}
190+
183191
pub fn set(&mut self, x: usize, y: usize, val: T) {
184192
self.dat[self.width*y + x] = val;
185193
}

‎src/hl.rs

Lines changed: 16 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -413,8 +413,13 @@ unsafe impl Send for CommandQueue {}
413413
impl CommandQueue
414414
{
415415
//synchronous
416-
pub fn enqueue_kernel<I: KernelIndex, E: EventList>(&self, k: &Kernel, global: I, local: Option<I>, wait_on: E)
417-
-> Event
416+
pub fn enqueue_kernel<I: KernelIndex, E: EventList>(&self,
417+
k: &Kernel,
418+
global_offset: Option<I>,
419+
global: I,
420+
local: Option<I>,
421+
wait_on: E)
422+
-> Event
418423
{
419424
unsafe
420425
{
@@ -424,7 +429,7 @@ impl CommandQueue
424429
self.cqueue,
425430
k.kernel,
426431
KernelIndex::num_dimensions(None::<I>),
427-
ptr::null(),
432+
global_offset.map(|x| x.get_ptr()).unwrap_or(ptr::null()),
428433
global.get_ptr(),
429434
match local {
430435
Some(ref l) => l.get_ptr() as *const libc::size_t,
@@ -442,8 +447,13 @@ impl CommandQueue
442447
}
443448

444449
//asynchronous
445-
pub fn enqueue_async_kernel<I: KernelIndex, E: EventList>(&self, k: &Kernel, global: I, local: Option<I>, wait_on: E)
446-
-> Event
450+
pub fn enqueue_async_kernel<I: KernelIndex, E: EventList>(&self,
451+
k: &Kernel,
452+
global_offset: Option<I>,
453+
global: I,
454+
local: Option<I>,
455+
wait_on: E)
456+
-> Event
447457
{
448458
unsafe
449459
{
@@ -453,7 +463,7 @@ impl CommandQueue
453463
self.cqueue,
454464
k.kernel,
455465
KernelIndex::num_dimensions(None::<I>),
456-
ptr::null(),
466+
global_offset.map(|x| x.get_ptr()).unwrap_or(ptr::null()),
457467
global.get_ptr(),
458468
match local {
459469
Some(ref l) => l.get_ptr() as *const libc::size_t,

‎tests/test.rs

Lines changed: 58 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -160,7 +160,7 @@ mod hl {
160160

161161
k.set_arg(0, &v);
162162

163-
queue.enqueue_async_kernel(&k, 1isize, None, ()).wait();
163+
queue.enqueue_async_kernel(&k, None, 1isize, None, ()).wait();
164164

165165
let v: Vec<isize> = queue.get(&v, ());
166166

@@ -185,7 +185,7 @@ mod hl {
185185
k.set_arg(0, &v);
186186
k.set_arg(1, &42isize);
187187

188-
queue.enqueue_async_kernel(&k, 1isize, None, ()).wait();
188+
queue.enqueue_async_kernel(&k, None, 1isize, None, ()).wait();
189189

190190
let v: Vec<isize> = queue.get(&v, ());
191191

@@ -209,7 +209,7 @@ mod hl {
209209

210210
k.set_arg(0, &v);
211211

212-
queue.enqueue_async_kernel(&k, 1isize, None, ()).wait();
212+
queue.enqueue_async_kernel(&k, None, 1isize, None, ()).wait();
213213

214214
let v: Vec<isize> = queue.get(&v, ());
215215

@@ -234,7 +234,7 @@ mod hl {
234234

235235
let mut e : Option<Event> = None;
236236
for _ in 0isize .. 8 {
237-
e = Some(queue.enqueue_async_kernel(&k, 1isize, None, e));
237+
e = Some(queue.enqueue_async_kernel(&k, None, 1isize, None, e));
238238
}
239239
e.wait();
240240

@@ -269,15 +269,15 @@ mod hl {
269269
k_inc_b.set_arg(0, &b);
270270

271271
let event_list = [
272-
queue.enqueue_async_kernel(&k_inc_a, 1isize, None, ()),
273-
queue.enqueue_async_kernel(&k_inc_b, 1isize, None, ()),
272+
queue.enqueue_async_kernel(&k_inc_a, None, 1isize, None, ()),
273+
queue.enqueue_async_kernel(&k_inc_b, None, 1isize, None, ()),
274274
];
275275

276276
k_add.set_arg(0, &a);
277277
k_add.set_arg(1, &b);
278278
k_add.set_arg(2, &c);
279279

280-
let event = queue.enqueue_async_kernel(&k_add, 1isize, None, &event_list[..]);
280+
let event = queue.enqueue_async_kernel(&k_add, None, 1isize, None, &event_list[..]);
281281

282282
let v: Vec<isize> = queue.get(&c, event);
283283

@@ -312,7 +312,7 @@ mod hl {
312312

313313
k.set_arg(0, &v);
314314

315-
queue.enqueue_async_kernel(&k, (3isize, 3isize), None, ()).wait();
315+
queue.enqueue_async_kernel(&k, None, (3isize, 3isize), None, ()).wait();
316316

317317
let v: Vec<isize> = queue.get(&v, ());
318318

@@ -385,7 +385,7 @@ mod hl {
385385

386386
k.set_arg(0, &v);
387387

388-
let e = queue.enqueue_async_kernel(&k, 1isize, None, ());
388+
let e = queue.enqueue_async_kernel(&k, None, 1isize, None, ());
389389
e.wait();
390390

391391
// the that are returned are not useful for unit test, this test
@@ -469,7 +469,7 @@ mod array {
469469
let k = prog.create_kernel("test");
470470

471471
k.set_arg(0, &a_cl);
472-
let event = queue.enqueue_async_kernel(&k, (8isize, 8isize), None, ());
472+
let event = queue.enqueue_async_kernel(&k, None, (8isize, 8isize), None, ());
473473
queue.read(&a_cl, &mut a, &event);
474474

475475
for x in 0usize .. 8usize {
@@ -480,6 +480,53 @@ mod array {
480480
})
481481
}
482482

483+
#[test]
484+
fn kernel_2d_offset()
485+
{
486+
::test_all_platforms_devices(&mut |device, ctx, queue| {
487+
let mut a = Array2D::new(8, 8, |_, _| {(1) as i32});
488+
let b = Array2D::new(8, 8, |x, y| {(x*y) as i32});
489+
let a_cl = ctx.create_buffer_from(&a, CL_MEM_READ_WRITE);
490+
491+
let src = "__kernel void test(__global int *a, ulong size_x) { \
492+
int x_off = get_global_offset(0); \
493+
int x = get_global_id(0) + x_off; \
494+
int y_off = get_global_offset(1); \
495+
int y = get_global_id(1) + y_off; \
496+
a[size_x*y + x] = x*y; \
497+
}";
498+
let prog = ctx.create_program_from_source(src);
499+
match prog.build(device) {
500+
Ok(_) => (),
501+
Err(build_log) => {
502+
println!("Error building program:\n");
503+
println!("{}", build_log);
504+
panic!("");
505+
}
506+
}
507+
let k = prog.create_kernel("test");
508+
509+
k.set_arg(0, &a_cl);
510+
k.set_arg(1, &a.width());
511+
let event = queue.enqueue_async_kernel(&k, Some((3, 3)), (5isize, 5isize), None, ());
512+
queue.read(&a_cl, &mut a, &event);
513+
514+
println!("");
515+
for y in 0usize .. 8usize {
516+
for x in 0usize .. 8usize {
517+
let _a = a.get(x, y);
518+
print!("{:?}\t", _a);
519+
if x < 3 || y < 3 {
520+
expect!(a.get(x, y), 1);
521+
} else {
522+
expect!(a.get(x, y), b.get(x, y));
523+
}
524+
}
525+
println!("");
526+
}
527+
})
528+
}
529+
483530
#[test]
484531
fn put_get_3d()
485532
{
@@ -552,7 +599,7 @@ mod array {
552599
let k = prog.create_kernel("test");
553600

554601
k.set_arg(0, &a_cl);
555-
let event = queue.enqueue_async_kernel(&k, (8isize, 8isize, 8isize), None, ());
602+
let event = queue.enqueue_async_kernel(&k, None, (8isize, 8isize, 8isize), None, ());
556603
queue.read(&a_cl, &mut a, &event);
557604

558605
for x in 0usize .. 8usize {

0 commit comments

Comments
 (0)