Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Clang Generating Invalid SPIRV #92135

Open
Calandracas606 opened this issue May 14, 2024 · 0 comments
Open

Clang Generating Invalid SPIRV #92135

Calandracas606 opened this issue May 14, 2024 · 0 comments
Labels
SPIR-V SPIR-V language support

Comments

@Calandracas606
Copy link
Contributor

Calandracas606 commented May 14, 2024

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

@github-actions github-actions bot added the clang Clang issues not falling into any other category label May 14, 2024
@EugeneZelenko EugeneZelenko added SPIR-V SPIR-V language support and removed clang Clang issues not falling into any other category labels May 14, 2024
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
SPIR-V SPIR-V language support
Projects
None yet
Development

No branches or pull requests

2 participants