On 2016/3/25 上午 02:40, Martin Jambor wrote:
> On the whole, I am fine with the patch but there are two issues:
>
> First, and generally, when you change the return type of a function,
> you must document what return values mean in the comment of the
> function. Most importantly, it must be immediately apparent whether a
> function returns true or false on failure from its comment. So please
> fix that.
Thanks, I'll update on that.
>> > /* Callback of dispatch queues to report errors. */
>> > @@ -454,7 +471,7 @@ queue_callback (hsa_status_t status,
>> > hsa_queue_t *queue __attribute__ ((unused)),
>> > void *data __attribute__ ((unused)))
>> > {
>> > - hsa_fatal ("Asynchronous queue error", status);
>> > + hsa_error ("Asynchronous queue error", status);
>> > }
> ...I believe this hunk is wrong. Errors reported in this way mean
> that something is very wrong and generally happen during execution of
> code on HSA GPU, i.e. within GOMP_OFFLOAD_run. And since you left
> calls in create_single_kernel_dispatch, which is called as a part of
> GOMP_OFFLOAD_run, intact, I believe you actually want to leave
> hsa_fatel here too.
Yes, a fatal exit is okay within the 'run' hook, since we're not holding
the device lock there. I was only trying to audit the GOMP_OFFLOAD_init_device()
function, where the queues are created.
I'm not familiar with the HSA runtime API; will the callback only be triggered
during GPU kernel execution (inside the 'run' hook), and not for example,
within hsa_queue_create()? If so, then yes as you advised, the above change to
queue_callback() should be reverted.
Chung-Lin
