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

[hughperkins]

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?

[hughperkins]

Per SPIR-V 1.1 reference pdf https://www.khronos.org/registry/spir-v/specs/1.1/SPIRV.pdf, OpCapability is a standard top-level entity:

So... some issue with the parser?

(edit: also, more explicitly:

)

[hughperkins]

ah, probably I should be linking with spirv-tools https://github.com/KhronosGroup/SPIRV-Tools#usage Closing for now

[hughperkins]

actually, unclear on this point. Opening, pending confirmation/clarification on what are the standard way(s) of reading spir-v code, from c++.

[pierremoreau]

SPIR-V modules are stored as a stream of bytes. So, if it has the same endianness as your CPU, you can do something like:

#include <fstream>
#include <string>

#include "spirv.hpp11"

int main()
{
  std::ifstream file("filename.spv");
  std::string string_binary((std::istreambuf_iterator<char>(file)), (std::istreambuf_iterator<char>()));
  unsigned int num_words = string_binary.size() / 4u;
  const unsigned int *binary = reinterpret_cast<const unsigned int *>(string_binary.c_str());

  unsigned int i = 0u;
  while (i < num_words) {
    unsigned int opcode_description = binary[i];
    spv::Op opcode = static_cast<spv::Op>(opcode_description && spv::OpCodeMask);
    unsigned int num_operands = opcode_description >> spv::WordCountShift;

    if (opcode == spv::Op::OpLoad) {
      spv::Id result_type = static_cast<spv::Id>(binary[i + 1u]);
      spv::Id result      = static_cast<spv::Id>(binary[i + 2u]);
      spv::Id pointer     = static_cast<spv::Id>(binary[i + 3u]);
      spv::MemoryAccessMask access = spv::MemoryAccessMask::None;
      if (num_operands == 5u)
        access = static_cast<spv::MemoryAccessMask>(binary[i + 4u]);
    }

    i += num_operands;
  }

  return 0;
}

NOTE: I haven’t test compiled/run it, but it should give you an idea on how to do it. If I understood correctly what you wanted to do.

[hughperkins]

Cool. :+1: Questions:

• how to parse it using standard llvm tools, like parseIRFile
• assuming we have to use the parser in SPIRV-Tools, what option(s) exist to then somehow convert the parsed SPIR-V into a format we can manipulate using standard llvm functions, such as Module iterators et al?

[hughperkins]

Figured it out: use ReadSPIRV, in header file llvm/Support/SPIRV.h => closing

[hughperkins]

Well.. question: what is the difference between the methods available in llvm/Support/SPIRV.h, adn the methods available in spirv-tools/spirv.hpp1/spirv-tools/libspirv.h? Is it like:

• the methods in llvm/Support/SPIRV.h allow one to bring the full power of llvm to bear on the problem?, however
• the methods in spirv-tools/ are lighter weight, avoid needing to link with llvm libraries, at runtime?