New sdk vulkan seems needs what is now above last tagged release.

It would be good to see better coordination about releasing stuff between spirv-* and vulkan-* repos.

With installed pkgconfig file other projects build processes can detest availability of the spirv-headers and require some minimum version of the spirv-headers to be present in build environment.

There i many problems when some other packages are failing only because installed spirv-headers version is to low. Adding that file will allow to avoid that kind of situations.

For code that needs to do minor inspection/validation of SPIR-V it is often useful to be able to build an id->typeid map, but there's currently no good way to know which opcodes have a result/type and which don't, and this makes it clumsy to try to build the map.

My original motivation for this is at: https://github.com/KhronosGroup/Vulkan-ValidationLayers/pull/742/files#diff-02aa55cdf8efb86e40256744eb74b841R1949. I needed to do some validation of cooperative matrix types, but the cooperative matrix operands could be a result of a bunch of different instructions. There were few enough that it was reasonable to whitelist the instructions in this case.

Baldur described a case where he needed to know what type of integer an OpAccessChain operand is. This would be easy to know by building an id->typeid map and only inspecting OpTypeInt and OpAccessChain instructions.

So I think this is generally useful, and the JSON has this information encoded in it. The core SPIR-V header seems like a natural place to put this, though I can imagine somebody objecting to having utility code creep into the header.

If the idea behind OpSelectionMerge and OpLoopMerge is that threads should know when they should diverge and converge, why have 2 separate instructions for conditional statements and loops? The fact that there are 2 instructions gives rise to many corner case questions, questions such as "should a OpLoopMerge block have OpSelectionMerge blocks within itself or not" (SPIR-V documentation example says no, while I sometimes see SPIR-V code in the wild that says yes). I see driver programmers from a big IHV hit these corner cases and complain about it on the Internet in the year of 2019, 5 years after the release of the spec, this is not normal, there are Android phones out there with broken SPIR-V compilers because of that, and many more will come if someone won't clarify all of this in detail.

Note: cross-posted from https://github.com/KhronosGroup/SPIRV-LLVM/issues/151

given the following OpenCL C code:

typedef struct {
    int val;
} test_struct;

kernel void struct_test(global int* buf, test_struct param) {
    buf[get_global_id(0)] = param.val;
}

kernel void int_test(global int* buf, int param) {
    buf[get_global_id(0)] = param;
}

resulting in the following IR (shortened for brevity):

%struct.test_struct = type { i32 }

define spir_kernel void @struct_test(i32 addrspace(1)* %buf, %struct.test_struct* %param) nounwind {
  %1 = getelementptr inbounds %struct.test_struct* %param, i64 0, i32 0
  %2 = load i32* %1, align 4, !tbaa !12
  %3 = tail call spir_func i64 @_Z13get_global_idj(i32 0) nounwind readnone
  %4 = getelementptr inbounds i32 addrspace(1)* %buf, i64 %3
  store i32 %2, i32 addrspace(1)* %4, align 4, !tbaa !12
  ret void
}

define spir_kernel void @int_test(i32 addrspace(1)* %buf, i32 %param) nounwind {
  %1 = tail call spir_func i64 @_Z13get_global_idj(i32 0) nounwind readnone
  %2 = getelementptr inbounds i32 addrspace(1)* %buf, i64 %1
  store i32 %param, i32 addrspace(1)* %2, align 4, !tbaa !12
  ret void
}

resulting in the following SPIR-V (shortened for brevity):

               OpEntryPoint Kernel %12 "struct_test"
               OpEntryPoint Kernel %25 "int_test"
               OpDecorate %5 LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
          %2 = OpTypeInt 64 0
          %7 = OpTypeInt 32 0
         %16 = OpConstant %2 0
         %17 = OpConstant %7 0
          %3 = OpTypeVector %2 3
          %4 = OpTypePointer UniformConstant %3
          %6 = OpTypeVoid
          %8 = OpTypePointer CrossWorkgroup %7
          %9 = OpTypeStruct %7
         %10 = OpTypePointer Function %9
         %11 = OpTypeFunction %6 %8 %10
         %18 = OpTypePointer Function %7
         %24 = OpTypeFunction %6 %8 %7
          %5 = OpVariable %4 UniformConstant
         %12 = OpFunction %6 None %11
         %13 = OpFunctionParameter %8
         %14 = OpFunctionParameter %10
         %15 = OpLabel
         %19 = OpInBoundsPtrAccessChain %18 %14 %16 %17
         %20 = OpLoad %7 %19 Aligned 4
         %21 = OpLoad %3 %5 Aligned 0
         %22 = OpCompositeExtract %2 %21 0
         %23 = OpInBoundsPtrAccessChain %8 %13 %22
               OpStore %23 %20 Aligned 4
               OpReturn
               OpFunctionEnd
         %25 = OpFunction %6 None %24
         %26 = OpFunctionParameter %8
         %27 = OpFunctionParameter %7
         %28 = OpLabel
         %29 = OpLoad %3 %5 Aligned 0
         %30 = OpCompositeExtract %2 %29 0
         %31 = OpInBoundsPtrAccessChain %8 %26 %30
               OpStore %31 %27 Aligned 4
               OpReturn
               OpFunctionEnd

Is the way kernel struct parameters are handled really the correct/intended behavior? Considering that scalar types are directly used in OpFunctionParameter/OpTypeFunction, shouldn't structs be handled the same way instead of going through an "OpTypePointer Function" indirection? Even more, doesn't this indirection say that only a pointer argument will be set/used (4 or 8 bytes), not so much a struct object (which could be any size)? I know that the issue here is that LLVM/SPIR can only handle struct parameters as pointers, but something like that isn't specified for SPIR-V.

How to solve this?

Option 1 (preferable): Keep it the way it is right now, but explicitly specify that kernel pointer parameters to Function/private memory actually perform some kind of allocation of the element/pointee type on the device side, and are set as this element/pointee type on the host side (not as the pointer type). Note that private address space pointer kernel arguments are otherwise invalid.

Option 2 (impossible?): Directly use OpTypeStruct in OpFunctionParameter/OpTypeFunction. This will however require IR/SPIR-V translator changes, since OpTypeStruct is no longer a pointer type (making all GEPs/Op*AccessChain instructions using it invalid). This might be impossible to do though, since there is no way of getting a pointer to this struct then in SPIR-V (afaik).

edit: Option 3: Require a OpVariable in OpFunctionParameter/OpTypeFunction for struct types. This way it should be clear what is actually happening + it is still a pointer.

…PIRV-LLVM-Translator

List of extensions: SPV_INTEL_fpga_memory_attributes https://github.com/KhronosGroup/SPIRV-Registry/blob/master/extensions/INTEL/SPV_INTEL_fpga_memory_attributes.asciidoc

SPV_INTEL_kernel_attributes https://github.com/KhronosGroup/SPIRV-Registry/blob/master/extensions/INTEL/SPV_INTEL_kernel_attributes.asciidoc

SPV_INTEL_fpga_reg https://github.com/KhronosGroup/SPIRV-Registry/blob/master/extensions/INTEL/SPV_INTEL_fpga_reg.asciidoc

SPV_INTEL_blocking_pipes https://github.com/KhronosGroup/SPIRV-Registry/blob/master/extensions/INTEL/SPV_INTEL_blocking_pipes.asciidoc

SPV_INTEL_fpga_loop_controls https://github.com/KhronosGroup/SPIRV-Registry/blob/master/extensions/INTEL/SPV_INTEL_fpga_loop_controls.asciidoc

SPV_INTEL_unstructured_loop_controls https://github.com/KhronosGroup/SPIRV-Registry/blob/master/extensions/INTEL/SPV_INTEL_unstructured_loop_controls.asciidoc

Signed-off-by: Dmitry Sidorov [email protected]

switchcase_goto_loop_merge.spv.txt

The switch-case is inside the loop. It branches from switch-case to loop merge block. I think it's invalid, but I'd like someone to confirm it. Is there some validator that can check it?

    %808 = OpLabel
           OpLoopMerge **%809** %810 None
           OpBranch %811
    ...
           OpSelectionMerge %813 None
           OpSwitch %849 %814 0 %815 1 %816 2 %817
    %814 = OpLabel
           OpBranch **%809**
    %815 = OpLabel
    %818 = OpLoad %211 %213
    %819 = OpLoad %215 %217
    %820 = OpSampledImage %219 %818 %819
    %822 = OpImageSampleImplicitLod %27 %820 %670
           OpBranch **%809**
    ...
    %813 = OpLabel
           OpUnreachable
    %810 = OpLabel
           OpBranch %808
    %809 = OpLabel
    ...

The current definition of a back edge in SPIR-V is not quite correct:

Back Edge: If a depth-first traversal is done on a function’s CFG, starting from the first block of the function, a back edge is a branch to a previously visited block. A back-edge block is the block containing such a branch.

This was a late addition that doesn't quite work with the intent of the structured control flow rules.

Suggestions from @dneto0, in issue https://github.com/KhronosGroup/SPIRV-Tools/issues/270 :

I think the spec's definition of back-edge incorrectly allows cross edges to be classified as back-edges. I think it should instead say the equivalent of "A back-edge is an edge in the CFG from a block B to a block C that dominates B. B may be the same as C."

But I think glslang's code generation is sensible, and if anything the spec should be amended to clearly allow unreachable continue-constructs.

For B-> C to be a back edge the "previously visited" phrasing to me means "I am currently visiting B, and the edge makes me want to visit C, but I have previously visited C". It's the "I am currently visiting B" part that is not satisfied if doing a DFS starting a the entry block of the function.

Fix grammar error for Google decorate_string & hlsl_functionality1.

We don't need to put "LiteralString" as a parameter to OpDecorateString itself; they should be specified by the Decoration instead.

Note that for the backport, I also included OpDecorateId, but guarded it with an "extension" field.

Let me know if this is not the preferred way to go. I'd happy to adjust.

According to the description of OpLoopMerge, the loop header block can contain a conditional branch (OpBranchConditional). Doing so would however require the block to contain an OpSelectionMerge in addition to the OpLoopMerge. This is currently not allowed by the spec, because both descriptions state that (OpLoopMerge|OpSelectionMerge) "must be the second-to-last instruction in its block".

E.g. this code is currently invalid:

   OpSelectionMerge %selmergebb None
        OpLoopMerge %loopmergebb %continuebb None
OpBranchConditional %condition %true %false

Solution: either remove OpBranchConditional from the OpLoopMerge description or refine the "second-to-last" text so that both merge operations can occur in a block.

If doing the latter: shouldn't OpLoopMerge allow OpSwitch as well then?

As specified in https://www.khronos.org/registry/spir-v/extensions/KHR/SPV_KHR_shader_ballot.html:

 (Add the SubgroupBallotKHR capability to SubgroupSize.)
 (Add the SubgroupBallotKHR capability to SubgroupLocalInvocationId.)

According to the spec, OpGroupNonUniformPartitionNV is currently reserved, but not defined (see https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpGroupNonUniformPartitionNV). However, as @greg-lunarg pointed out:

"Looking at the spec for the GLSL extension GL_NV_shader_subgroup_partitioned, subgroupPartitionNV maps directly to OpGroupNonUniformPartitionNV, so it seems OpGroupNonUniformPartitionNV is well defined."

(See https://github.com/KhronosGroup/glslang/issues/3085 for reference.)

The registry home page indicated that feedback should be posted to this repository.

This SPIRV extension update will be introducing the new execution mode RegisterMapInterfaceINTEL. This PR adds support for this new mode.

NOTE

This execution mode is being added to an existing extension SPV_INTEL_kernel_attributes and is guarded by a new capability FPGAKernelAttributesINTELv2 which implicitly defines the existing capability FPGAKernelAttributesINTEL.

The Vulkan 1.3.222 spec added a reference to three new SPIR-V capabilities to vk.xml:

• TextureSampleWeightedQCOM
• TextureBoxFilterQCOM
• TextureBlockMatchQCOM

These capabilities are currently missing from the spirv.core.grammar.json file.

from: https://github.com/KhronosGroup/Vulkan-ValidationLayers/pull/4634#discussion_r992749426

Would be nice to have a constant value in the headers to use for code who needs to skip to first instruction without just having a magic 5 in their code

Created an entry in https://github.com/KhronosGroup/Khronos-Schemas/pull/30 that would get pulled in first

Using the grammar file currently to generate things for SPIR-V tools has an issue of not even knowing what is valid to be in the grammar file. This PR is an attempt to create a schema for the grammar files.

I took inspiration from the https://github.com/KhronosGroup/Vulkan-Profiles usage of a schema file.

Has been tested on spirv.core.grammar.json and extinst.*.grammar.json

Happy to work out the details on how to make this logistically as simple as to not add a burden to future header releases