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

llvm/sycl/source/handler.cpp

Lines 956 to 989 in 9457ac2

	static void addArgsForGlobalAccessor(detail::Requirement *AccImpl, size_t Index,
	size_t &IndexShift, int Size,
	bool IsKernelCreatedFromSource,
	size_t GlobalSize,
	std::vector<detail::ArgDesc> &Args,
	bool isESIMD) {
	using detail::kernel_param_kind_t;
	if (AccImpl->PerWI)
	AccImpl->resize(GlobalSize);
	Args.emplace_back(kernel_param_kind_t::kind_accessor, AccImpl, Size,
	Index + IndexShift);
	// TODO ESIMD currently does not suport offset, memory and access ranges -
	// accessor::init for ESIMD-mode accessor has a single field, translated
	// to a single kernel argument set above.
	if (!isESIMD && !IsKernelCreatedFromSource) {
	// Dimensionality of the buffer is 1 when dimensionality of the
	// accessor is 0.
	const size_t SizeAccField =
	sizeof(size_t) * (AccImpl->MDims == 0 ? 1 : AccImpl->MDims);
	++IndexShift;
	Args.emplace_back(kernel_param_kind_t::kind_std_layout,
	&AccImpl->MAccessRange[0], SizeAccField,
	Index + IndexShift);
	++IndexShift;
	Args.emplace_back(kernel_param_kind_t::kind_std_layout,
	&AccImpl->MMemoryRange[0], SizeAccField,
	Index + IndexShift);
	++IndexShift;
	Args.emplace_back(kernel_param_kind_t::kind_std_layout,
	&AccImpl->MOffset[0], SizeAccField, Index + IndexShift);
	}
	}

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.