CPU OpenCL runtime: Apparent miscompilation of conditional #2038
Description
Running this snippet:
import numpy as np
import pyopencl as cl
import pyopencl.array
src = """
#define lid(N) ((int) get_local_id(N))
#define gid(N) ((int) get_group_id(N))
__kernel void __attribute__ ((reqd_work_group_size(16, 1, 1))) actx_special_cos(__global double const *__restrict__ inp0, int const inp0_offset, int const n0, int const n1, __global double *__restrict__ out, int const out_offset)
{
if (-1 * lid(0) + n1 >= 0) // CRASHES/WRONG OUTPUT ON INTEL
//if (n1 >= lid(0)) // NO FAILURE ON INTEL
{
out[(1 + n1) * gid(0) + lid(0)] = cos(inp0[(1 + n1) * gid(0) + lid(0)]);
}
}
"""
def main():
ctx = cl._csc()
queue = cl.CommandQueue(ctx)
inp = cl.array.zeros(queue, (3, 4), dtype=float)
out = cl.array.zeros_like(inp)
prg = cl.Program(ctx, src).build()
knl = prg.actx_special_cos
knl(
queue, (16*3,), (16,), inp.data, np.int32(0),
np.int32(inp.shape[0]-1), np.int32(inp.shape[1]-1),
out.data, np.int32(0))
c = out.get()
print(c)
good = (c == 1)
assert good.all(), np.where(good)
main()on the CPU runtime version oclcpuexp-2020.10.6.0.4_rel distributed from this release page gives me
[[0. 0. 0. 1.]
[1. 1. 1. 1.]
[1. 1. 1. 1.]]
Traceback (most recent call last):
File "testac2.py", line 38, in <module>
main()
File "testac2.py", line 36, in main
assert good.all(), np.where(good)
AssertionError: (array([0, 1, 1, 1, 1, 2, 2, 2, 2]), array([3, 0, 1, 2, 3, 0, 1, 2, 3]))
double free or corruption (out)
- Oclgrind has no complaints about this kernel.
- Every other OpenCL implementation with which I try this kernel gives me what I'd view as the correct output:
[[1. 1. 1. 1.] [1. 1. 1. 1.] [1. 1. 1. 1.]] - Uncommenting the (IMO) equivalent rewrite of the conditional at the top and commenting the other version makes the Intel CPU runtime behave the same as all other CL implementations.
x-ref: https://gitlab.tiker.net/inducer/meshmode/-/merge_requests/80#note_49286