search

KhronosGroup / SPIRV-LLVM

This project is no longer active. Please join us at
https://github.com/KhronosGroup/SPIRV-LLVM-Translator
Other
262 stars 60 forks source link

Invalid function parameter attributes generated #239

Closed aarongreig closed 5 years ago

aarongreig commented 5 years ago

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 5 years ago

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

aarongreig commented 5 years ago

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 5 years ago

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

AlexeySotkin commented 5 years ago

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 5 years ago

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