The number of parameters for buffer in spv file ranges from 2-4

[shixinlishixinli]

Describe the bug

i try to use spv file , which is generated by "icpx -save-temps". but i have found that the num of parameters for buffer is different from 2 to 4. why the num is different ? how to set arguments by in my level0 code? how to calculate the num of params in level0 through sycl code?

• index: 0 name: _arg_d_acc address_qualifier: __global access_qualifier: NONE type_name: 'float*;8' type_qualifiers: NONE
  • index: 1 name: _arg_d_acc2 address_qualifier: __private access_qualifier: NONE type_name: 'class.sycl::_V1::range;16' type_qualifiers: NONE
  • index: 2 name: _arg_d_acc3 address_qualifier: __private access_qualifier: NONE type_name: 'class.sycl::_V1::range;16' type_qualifiers: NONE

Best Lisa Shi

To reproduce

1. Include a code snippet that is as short as possible
2. Specify the command which should be used to compile the program
3. Specify the command which should be used to launch the program
4. Indicate what is wrong and what was expected

Environment

• OS: [e.g Windows/Linux]
• Target device and vendor: [e.g. Intel GPU]
• DPC++ version: [e.g. commit hash or output of clang++ --version]
• Dependencies version: [e.g. the output of sycl-ls --verbose]

Additional context

No response

[AlexeySachkov]

Hi @shixinlishixinli,

i try to use spv file , which is generated by "icpx -save-temps". but i have found that the num of parameters for buffer is different from 2 to 4.

I assume that by "buffer" you mean accessor. accessor has plenty of methods like querying its max range/size or offset and therefore it has to know some information about a buffer it points to, which turns a single accessor into 4 arguments, see https://github.com/intel/llvm/blob/9457ac27fd03906448b984fe4a17093815ecfb34/sycl/source/handler.cpp#L956-L989 for an example of how its handled.

why the num is different ?

We also have an optimization to eliminate unused kernel arguments, so if some of accessor fields aren't needed for a kernel - they will be optimized out. That explains why for some kernels we don't need all four. You can try -fno-sycl-dead-args-optimization (its ON by default) and I expect the result to stabilize at 4 kernel arguments per accessor.

how to set arguments by in my level0 code? how to calculate the num of params in level0 through sycl code?

How is your kernel written? Is it a lambda? Is it a functor?

I think that the only reliable way to submit a SYCL kernel directly through L0 is to use sycl_ext_oneapi_free_function_kernels extension which should guarantee a stable arguments interface, but even then you may have to stay away from complex data types like accessors.