Open
Description
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