Re: [PATCH i-g-t 3/5] lib/intel_compute: Allow the user to provide a custom compute kernel

[Francois Dugast <francois.dugast@intel.com>]

Hi Zbigniew,

On Thu, Feb 13, 2025 at 08:00:06AM +0100, Zbigniew Kempczyński wrote:
> On Wed, Feb 05, 2025 at 11:17:04AM +0100, Francois Dugast wrote:
> > Allow the user to provide a custom compute kernel which will be used
> > instead of the default compute square one. This will be helpful to
> > try out corner cases which require a specific compute kernel.
> > 
> > Signed-off-by: Francois Dugast <francois.dugast@intel.com>
> > ---
> >  lib/intel_compute.c | 26 ++++++++++++++++++--------
> >  lib/intel_compute.h |  2 ++
> >  2 files changed, 20 insertions(+), 8 deletions(-)
> > 
> > diff --git a/lib/intel_compute.c b/lib/intel_compute.c
> > index e0776fb6d..a826d58c0 100644
> > --- a/lib/intel_compute.c
> > +++ b/lib/intel_compute.c
> > @@ -1770,6 +1770,8 @@ static bool __run_intel_compute_kernel(int fd,
> >  	unsigned int batch;
> >  	const struct intel_compute_kernels *kernels = intel_compute_square_kernels;
> >  	enum intel_driver driver = get_intel_driver(fd);
> > +	const unsigned char *kernel;
> > +	unsigned int kernel_size;
> >  
> >  	for (batch = 0; batch < ARRAY_SIZE(intel_compute_batches); batch++) {
> >  		if (ip_ver == intel_compute_batches[batch].ip_ver)
> > @@ -1787,16 +1789,24 @@ static bool __run_intel_compute_kernel(int fd,
> >  		return false;
> >  	}
> >  
> > -	while (kernels->kernel) {
> > -		if (ip_ver == kernels->ip_ver)
> > -			break;
> > -		kernels++;
> > +	/* If the user provides a kernel, use it */
> > +	if (user && user->kernel) {
> > +		kernel = user->kernel;
> > +		kernel_size = user->kernel_size;
> > +	} else {
> > +		while (kernels->kernel) {
> > +			if (ip_ver == kernels->ip_ver)
> > +				break;
> > +			kernels++;
> > +		}
> > +		if (!kernels->kernel)
> > +			return false;
> > +		kernel = kernels->kernel;
> > +		kernel_size = kernels->size;
> 
> According to how we build pipeline indirect data we're limited to
> three arguments - *input, *output, count. Allowing the user to provide
> a custom kernel won't work unless type constraint will be met.
> I don't like this change, because someone who doesn't know how
> this work will start providing its own kernels will be surprised
> it doesn't work.

The way we statically build the pipeline indirect data is a good
balance of simplicity and flexibility, as we can test a lot even
with the constraint *input, *output, count.

More complex KMD tests will require simple specific kernels which
still comply with this constraint. For example the one below can
trigger a page fault at 0x10000 from the compute kernel context
so that we run other checks in KMD, all from IGT:

    __kernel void square(__global float* input,
                         __global float* output,
                         const unsigned int count) {
        int i = get_global_id(0);
        const __global uint* addr = 0x10000;
        output[i] = *addr;
    }

This is the reason for allowing custom compute kernels. This way
we can leverage the existing lib/intel_compute infrastructure to
test corner cases. I believe having this possibility is far more
important than the risk of a user incorrectly expecting IGT to
provide a full "compute runtime" able to run any kernel.

Would you be fine if the doc of struct user_execenv::kernel would
make the constraint of *input, *output, count explicit?

If not, what about replacing "const unsigned char *kernel" with
an enum which would be used to obtain the right kernel from a
library of "approved" kernels in intel_compute_kernels.c?

Francois

> 
> --
> Zbigniew
> 
> >  	}
> > -	if (!kernels->kernel)
> > -		return false;
> >  
> > -	intel_compute_batches[batch].compute_exec(fd, kernels->kernel,
> > -						  kernels->size, eci, user);
> > +	intel_compute_batches[batch].compute_exec(fd, kernel,
> > +						  kernel_size, eci, user);
> >  
> >  	return true;
> >  }
> > diff --git a/lib/intel_compute.h b/lib/intel_compute.h
> > index c4b4ee5e1..6096bb83a 100644
> > --- a/lib/intel_compute.h
> > +++ b/lib/intel_compute.h
> > @@ -35,6 +35,8 @@ struct intel_compute_kernels {
> >  
> >  struct user_execenv {
> >  	uint32_t vm;
> > +	const unsigned char *kernel;
> > +	unsigned int kernel_size;
> >  };
> >  
> >  extern const struct intel_compute_kernels intel_compute_square_kernels[];
> > -- 
> > 2.43.0
> >