KhronosGroup/SPIRV-LLVM

Invalid function parameter attributes generated

Closed this issue · 5 comments

aarongreig commented

It is possible to generate SPIR-V that applies function parameter attributes to invalid types. Repro:

__kernel void foo(__global float* dst, const int bar) {
  int globalId = get_global_id(0);

  if (globalId < bar) {
    dst[globalId] = globalId * 3.0f;
  } else {
    dst[globalId] = globalId * 4.0f;
  }
}

Since bar is const the spir-v generated from this kernel applies the NoWrite decoration to it, but NoWrite is only valid for pointer types.

AlexeySotkin commented

@aarongreig, thank you for reporting the issue. Could you attach LLVM IR and SPIR-V files for the repro.

aarongreig commented

Yes: paramAttr.ll.txt produces paramAttr.spvasm.txt (I had to add the .txt extension for github to accept the file). I realise now that I'm actually using the llvm 6.0 version of the tool, currently building 3.6.1 to see if this will repro with that version as well.

aarongreig commented

the same IR passed into llvm-spirv built from the 3.6.1 branch produces this paramAttr.spvasm.txt

AlexeySotkin commented

Sorry, I mixed this repository with https://github.com/KhronosGroup/SPIRV-LLVM-Translator.
There is no active development in this repository anymore.
So I suggest to use translator from https://github.com/KhronosGroup/SPIRV-LLVM-Translator. It should work with latest versions of LLVM/Clang.

aarongreig commented

The new version doesn't exhibit this bug, thanks for pointing me in the right direction.