`expected top-level entity`, when using `parseIRFile`, on output of `spirv_dis`

The error message:

The c++ code:
https://github.com/hughperkins/pub-prototyping/blob/3310247911a6b26a6768baf165f435faaa25e24e/spirv/test_spir_cpp.cpp#L27-L34

    llvm::LLVMContext context;
    llvm::SMDiagnostic smDiagnostic;
    std::string llFilename = "cl_kernel1.ll";
    std::unique_ptr<llvm::Module> M = parseIRFile(llFilename, smDiagnostic, context);
    if(!M) {
        smDiagnostic.print("irtoopencl", llvm::errs());
        throw std::runtime_error("failed to parse IR");
    }

The input SPIR-V code:

~/git-local/pub-prototyping/spirv/build (master|…5) $ cat cl_kernel1.ll
; SPIR-V
; Version: 1.0
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 22
; Schema: 0
               OpCapability Addresses
               OpCapability Linkage
               OpCapability Kernel
               OpCapability Int8
          %1 = OpExtInstImport "OpenCL.std"
               OpMemoryModel Physical32 OpenCL
               OpEntryPoint Kernel %10 "mykernel"
               OpSource OpenCL_C 102000
               OpName %5 "__spirv_BuiltInGlobalInvocationId"
               OpName %11 "cmem0"
               OpName %12 "offset"
               OpName %13 "entry"
               OpName %14 "add.ptr"
               OpName %18 "call"
               OpName %20 "add"
               OpName %21 "arrayidx"
               OpDecorate %5 BuiltIn GlobalInvocationId
               OpDecorate %5 Constant
               OpDecorate %5 LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
          %2 = OpTypeInt 32 0
          %7 = OpTypeInt 8 0
         %19 = OpConstant %2 123
          %3 = OpTypeVector %2 3
          %4 = OpTypePointer UniformConstant %3
          %6 = OpTypeVoid
          %8 = OpTypePointer CrossWorkgroup %7
          %9 = OpTypeFunction %6 %8 %2
         %15 = OpTypePointer CrossWorkgroup %2
          %5 = OpVariable %4 UniformConstant
         %10 = OpFunction %6 None %9
         %11 = OpFunctionParameter %8
         %12 = OpFunctionParameter %2
         %13 = OpLabel
         %14 = OpInBoundsPtrAccessChain %8 %11 %12
         %16 = OpBitcast %15 %14
         %17 = OpLoad %3 %5
         %18 = OpCompositeExtract %2 %17 0
         %20 = OpIAdd %2 %18 %19
         %21 = OpInBoundsPtrAccessChain %15 %16 %18
               OpStore %21 %20 Aligned 4
               OpReturn
               OpFunctionEnd

Generated by running this script https://github.com/hughperkins/pub-prototyping/blob/3310247911a6b26a6768baf165f435faaa25e24e/spirv/cl-to-spirv.sh

clang -cc1 -emit-spirv -triple spir-unknown-unknown -cl-std=CL1.2 -include opencl.h -x cl -o cl_kernel1.spv ../cl_kernel1.cl
spirv-dis cl_kernel1.spv -o cl_kernel1.ll

against this minimal opencl kernel: https://github.com/hughperkins/pub-prototyping/blob/3310247911a6b26a6768baf165f435faaa25e24e/spirv/cl_kernel1.cl

kernel void mykernel(global char *cmem0, unsigned int offset) {
    global int *data0 = (global int *)(cmem0 + offset);
    int tid = get_global_id(0);
    data0[tid] = tid + 123;
}

Versoins:

• spir-v 1.1
• khronos clang: 3.6.1

Thoughts?

Reproduce the failure with test_spir_cpp.cpp and the cl-to-spirv.sh commands, using the generated cl_kernel1.ll from spirv-dis. Start by inspecting parseIRFile's diagnostic and the displayed SPIR-V input, then determine whether the parser should accept this output or whether the generated file needs a different handling path. Done means the incompatibility is confirmed and covered by an appropriate regression test or clearly documented limitation.