Skip to content

Clang Generating Invalid SPIRV #92135

Open
@Calandracas606

Description

@Calandracas606

I am trying to compile C++ for OpenCL kernels to spirv, however clang is creating invalid spirv (verified with spirv-val)

clang 18.1.4 built from source on x86_64 linux

I am using the -O flag, which appears to be experimental:
https://github.com/KhronosGroup/OpenCL-Guide/blob/main/chapters/os_tooling.md

simple reproducer kernel.clcpp:

void sample(read_only image2d_t in_img, write_only image2d_t out_img) {

  const int col = get_global_id(0);
  const int row = get_global_id(1);

  const sampler_t samplerA = CLK_NORMALIZED_COORDS_FALSE |
                             CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

  float4 px = {0.0F, 0.0F, 0.0F, 1.0F};
  for (int i = 0; i < 3; ++i) {
    for (int j = 0; j < 3; ++j) {
      float4 val = read_imagef(in_img, samplerA, int2{col + i, row + j});
      px += val.xxxx;
    }
  }

  write_imagef(out_img, int2{col, row}, px);
};

I try to compile with:

clang -c -target spirv64 -O -o sample.spv kernel.clcpp
spirv-val sample.spv

i get the output:

error: line 51: Block '17[%17]' is already a merge block for another header
  %_Z6sample14ocl_image2d_ro14ocl_image2d_wo = OpFunction %void None %9

when I remove the -O flag, spirv-val no longer reports an error

when i manually invoke llvm-spirv and spirv-opt there is no error:

clang -c -target spirv64 -emit-llvm -o sample2.bc kernel.clcpp
llvm-spirv sample2.bc -o sample2.spv
spirv-opt sample2.spv -o sample2_opt.spv
spirv-val sample2_opt.spv

please let me know if this is not the correct area to report this bug, so I can report it elsewhere

Metadata

Metadata

Assignees

No one assigned

    Labels

    SPIR-VSPIR-V language support

    Type

    No type

    Projects

    No projects

    Milestone

    No milestone

    Relationships

    None yet

    Development

    No branches or pull requests

    Issue actions