CLOO (or probably OpenCL): Anyone used non-blocking ComputeCommandQueue.Read?

Nythrix, maybe you can shed some light on this, but I don't see how the current queue.Read(....) command in Cloo can be asynchronous in it's current form even if the underlying graphics card and driver supports it. For example, say we wanted async reads like:

float[] resultsA = queue.Read(a, false, 0, resultsLength, events);
float[] resultsB = queue.Read(b, false, 0, resultsLength, events);
float[] resultsC = queue.Read(c, false, 0, resultsLength, events);
queue.Finish();

Each line in the code would have to be blocking for the array variables to store valid data since they have to wait for queue.Read(....) to read and generate the results into .Net form. And I don't believe you can have 3 separate return calls running at the same time in .Net like this without a different syntax.

I think I don't understand the problem... It looks to me like the implementation of read right now creates the return array and returns it empty, having passed a pointer to the memory to the underlying call to populate at leisure.

Is the problem that the "gcHandle.Free" really shouldn't be called until the async call is finished, but it instead is called right away (thereby allowing the GC to move or blow away your memory before the underlying call gets around to writing to it)?

Well, if I can't make async calls, do you have any other ideas how I might achieve what I want? I have a large array (millions) of data I need to load, process, and return, but I was hoping to queue it up as a series of smaller chunks (thousands) to try to get the GPU working on the first chunk while the second one was being transferred.

Right now (if I understand everything correctly) the GPU sits idle until the entire array is loaded into its memory, then it processes the whole thing, then it sits idle again while the results are sent back.

The other thread you linked to implies that they couldn't get a data transfer and processing to happen at the same time no matter what they tried... which would be unfortunate.

Ok. One bit at a time.
1. Read/Write methods
You are right, they don't work properly in case the implementation supports asynchronous calls. Letting the user manage the resources is a simple solution but I'm not a fan of it because it is quite error prone. Besides I believe there's this better approach I've been playing with for some time now. The main idea is to bind the GCHandles to the ComputeEvents that the R/W methods generate thus delaying the call to GCHandle.Free() until the corresponding command is complete. This is fully transparent to the user and doesn't require any changes to the API.
2.

Have you hit this issue? I haven't been able to enforce this exact bug, even though I was sure it would happen. I think this proves that async calls are working in your environment. I would very much appreciate some info on how to reproduce this!
3. Kernel args
There is
ComputeKernel.SetValueArgument<T>( int index, T data ) for sending in any blittable structure, or
ComputeKernel.SetArgument( int index, IntPtr dataSize, IntPtr dataAddr ) which, apart from the error check, just calls the underlying clSetKernelArg so you can use it to send in virtually anything.
You cannot read back an argument unless you wrap it up in a ComputeBuffer. The kernel treats pointers as buffers anyway so the hack you're using is pretty much the only way.

Does the second chunk of your data depend on the first chunk of data to return a final result or are they totally independent? If they are totally independent, you could try creating host threads in .Net with each one having a command queue as suggested in the link or, you could split up your data into separate files and run multiple copies of your program with each pointed to a separate file (I can confirm this second method works). If your data chunks depend on one another to return a final result, I think you should wait and see if I have any success getting the async read to work on my computer.

Attached is a test case that shows that async reads execute on my computer but it does not properly read the data from the kernel when using my ClooHelper hack for structs. Example output:

---Async Read Demo---
Executing.....
Blocking Output:
testInt: 12345
testFloat: 1.2345

Async Output:
testInt: -1
testFloat: -1

The correct result is from the blocking output. I ruled out that GCAlloc is causing problems by manually allocating the pointer to the struct myself and freeing it manually at the end of the program.

Ref parameters in OpenTK bindings do bypass the marshaler (otherwise half of these overloads wouldn't even work):

public static unsafe int EnqueueReadBuffer<T5>(IntPtr command_queue, IntPtr buffer, bool blocking_read, IntPtr offset, IntPtr cb, [InAttribute, OutAttribute] ref T5 ptr, Int32 num_events_in_wait_list, IntPtr* event_wait_list, IntPtr* @event)
            where T5 : struct
        {
            GCHandle ptr_ptr = GCHandle.Alloc(ptr, GCHandleType.Pinned);
            try
            {
                int retval = Delegates.clEnqueueReadBuffer((IntPtr)command_queue, (IntPtr)buffer, (bool)blocking_read, (IntPtr)offset, (IntPtr)cb, (IntPtr)ptr_ptr.AddrOfPinnedObject(), (uint)num_events_in_wait_list, (IntPtr*)event_wait_list, (IntPtr*)@event);
                ptr = (T5)ptr_ptr.Target;
                return retval;
            }
            finally
            {
                ptr_ptr.Free();
            }
        }

That said, this doesn't really help here as the parameter is pinned only for the duration of the call. Then again, pinning the structure for the duration of the program should rule that issue out.

Some piece of the puzzle is missing here, I'll check again tomorrow with a clear brain.

Does anyone actually have asynchronous reads working in a C/C++ program with OpenCL?

This works:

string source = @"
kernel void Copy(
    global read_only float* bufferIn,
    global write_only float* bufferOut )
{
    int index = get_global_id(0);
    bufferOut[index] = bufferIn[index];
}
";
ComputeContextPropertyList cpl = new ComputeContextPropertyList( platform );
ComputeContext context = new ComputeContext( ComputeDeviceTypes.Default, cpl, null, IntPtr.Zero );
 
int size = 10;
float[] arrayIn = new float[ size ];
for( int i = 0; i < size; i++ ) arrayIn[ i ] = i;
 
ComputeBuffer<float> bufferIn = new ComputeBuffer<float>( context, ComputeMemoryFlags.ReadOnly | ComputeMemoryFlags.CopyHostPointer, arrayIn );
ComputeBuffer<float> bufferOut = new ComputeBuffer<float>( context, ComputeMemoryFlags.WriteOnly, size );
 
ComputeProgram program = new ComputeProgram( context, source );
program.Build( null, null, null, IntPtr.Zero );
ComputeKernel kernel = program.CreateKernel( "Copy" );
kernel.SetMemoryArgument( 0, bufferIn );
kernel.SetMemoryArgument( 1, bufferOut );
 
ComputeCommandQueue commands = new ComputeCommandQueue( context, context.Devices[ 0 ], ComputeCommandQueueFlags.None );
commands.Execute( kernel, null, new long[]{ size }, null, null );
 
float[] arrayOut = new float[ size ];
GCHandle arrayOutHandle = GCHandle.Alloc( arrayOut, GCHandleType.Pinned );
unsafe
{
    CL.EnqueueReadBuffer(
        commands.Handle, bufferOut.Handle, false, IntPtr.Zero, new IntPtr( bufferOut.Size ),
        arrayOutHandle.AddrOfPinnedObject(), 0, ( IntPtr* )null, ( IntPtr* )null );
}
commands.Finish();
arrayOutHandle.Free();

@Nythrix, yes that asynchronous read code you post works on my system, too! However, I can't seem to make it work for asynchronous read of single structs. Could you take a look at the code I posted before and see if you could make that work? I tried following your code but I still don't get any data read out for a single struct in non-blocking mode.

Nope, that ReadKernelArgSync did not work. In fact, using dataPtr.AddrOfPinnedObject() will cause no data to be read for a single struct no matter if it is blocking or not. You must use ref data in order for a single struct to be read successfully into .Net. I've tried out other variants of EnqueueReadBuffer, such reading a Vector4 array and a TestStruct array asynchronously and those work. I don't know exactly why this is so, but maybe I should just covert all the single structs in my .Net code to struct arrays and be done with the async read issue. I don't think I would have to change anything on the kernel side since the pointer notation would still work for structs array of length 1.

Just tried a pack size of 1, 4, and 8, but it didn't change anything. OpenCL expects a pack size that aligns vectors (8 for Vector2, 16 for Vector4, ie: 4 bytes per each float in a Vector). However, since I'm using a simple TestStruct with just a plain int and a float, I don't expect any alignment issues here. It's quite puzzling why the pinned memory doesn't work at all.

@Nythrix, that fixed it, thanks. The problem seems to be that "ref data" was not working in the async read command and instead required "new IntPtr(&data)". Looks like this wraps it up.

CL.EnqueueReadBuffer( 
    queue.Handle, buffer, block_read, new IntPtr(0), new IntPtr(Marshal.SizeOf(data)), 
    ref data,  // <- Don't do this! Use new IntPtr(&data) instead!
    0, (IntPtr*)null, (IntPtr*)null );

Yeah, I didn't get why you didn't like "ref data" before, but now I see. With "new IntPtr(&data)", there is now a minor problem where the computer says that it can't take the address of "data" if it is a generic type T since it's managed, but that's a minor problem I can work around by just overloading the ClooHelpers. But, I think you will run into this say compiler error when you modify public T[] Read( ......) in Cloo since you will have to take out the pinned memory address with a variant of "new IntPtr(&data)".

Nythrix, just to clarify, your improvements will fix asynchronous operations in concept, I.E. against any drivers that support it, but not for NVidia because their driver still does asynchronous operations synchronously.

So the only way (maybe, haven't tried it yet) to do asynchronous operations against NVidia is still multiple command queues.

Speaking of which, when is Intel going to do anything with OpenCL? They are really late to the game with their own drivers. Seems like they don't want to be a part of it.