[SYCL] CPU JIT ICE in VPlan Vectorization Driver

[agozillon]

I've attached an example SYCL program it's a pretty contrived and nonsensical example and will segfault if it actually runs properly. It's a problematic excerpt I found from a larger project that runs fine when using the Intel GPU OpenCL runtime but dies with the CPU implementation due to at minimum this segment I've found. Happy to redirect to the much larger project if that's of use.

However, for the moment at least on my machine and in my environment the attached example will result in what I think is a CPU JIT Compiler ICE:

   Stack dump:
      0.	Running pass 'Function Pass Manager' on module 'main'.
      1.	Running pass 'VPlan Vectorization Driver' on function ''
    Segmentation fault (core dumped)

I've unfortunately not had much time to dig into the IR after somewhat isolating the problem yet but the main problem area appears to be:

     for (; t < tfar; t += 1.0f) {
       f_tt = get(v);
       if (f_tt < 0)
         break;
     }
     t = t + 1.0f * f_tt / (1.0f - f_tt);

But there's a lot of very small things you can tweak in the example and it will compile properly, so it's been quite problematic to shrink it to a more concise code example sadly.

Currently using the most up-to-date Intel OpenCL CPU implementation and a recent (but not top of the tree, a few days or so old) version of the compiler.

VPlanFailure.zip

[AlexeySachkov]

Hi @agozillon, thanks for the report!

Just to be sure, am I right that

Currently using the most up-to-date Intel OpenCL CPU implementation

is oclcpuexp-2019.9.11.0.1106_rel?

[agozillon]

yeah that's the one, the one that fixed the previous JIT issue for me! Perhaps there is a new release with the October notes I am missing though?

[AlexeySachkov]

Hi @agozillon,

Perhaps there is a new release with the October notes I am missing though?

No, AFAIK, above mentioned release is the latest available one.

Good news: the issue seems to be not reproducible on the latest internal builds. Bad news: you have to wait for the next release

However, I believe that this issue is not blocking you since you already have workaround and it seems that attached reproducer is not fully correct:

template <typename T>
inline float get(T v) {
  return v->get()[0].x();
}
  
template<typename T>
float4 do_something(T v) {
	const float tnear = cl::sycl::fmax(0.1f, 1.0f);
	const float tfar  = cl::sycl::fmin(2.0f, 3.0f);

  if (tnear < tfar) {
     float t        = tnear;
     float f_tt     = 0;
    
     for (; t < tfar; t += 1.0f) {
       f_tt = get(v);
       if (f_tt < 0)
         break;
     }
     t = t + 1.0f * f_tt / (1.0f - f_tt);
     return float4{0.1f,0.1f,0.1f,t};
  }

  return float4{0,0,0,0};
}

int main() {
   queue q;    
   buffer<float4> ob{range<1>{1}};
    
   q.submit([&](handler &cgh) {      
      auto wb = ob.get_access<access::mode::write>(cgh);      
      cgh.single_task<codegen_or_jit_bug>([=]() {
        cl::sycl::global_ptr<short2> v[2];        
        wb[0] = do_something(v);
      });
  });
     
  return 0;
}

Please note the following:

cl::sycl::global_ptr<short2> v[2]; // array of pointers, initialized with default value (nullptr, I guess)
wb[0] = do_something(v); // passed to a function
// ...
f_tt = get(v); // then again passed to a function
// ...
return v->get()[0].x(); // and here we have dereference of nullptr

So, we have stored null into m_Pointer field of multi_ptr, in constructor:

; Function Attrs: nounwind
define linkonce_odr dso_local spir_func void @_ZN2cl4sycl9multi_ptrINS0_3vecIsLi2EEELNS0_6access13address_spaceE1EEC2Ev(%"class._ZTSN2cl4sycl9multi_ptrINS0_3vecIsLi2EEELNS0_6access13address_spaceE1EEE.cl::sycl::multi_ptr" addrspace(4)* %this) unnamed_addr #4 comdat align 2 {
entry:
  %this.addr = alloca %"class._ZTSN2cl4sycl9multi_ptrINS0_3vecIsLi2EEELNS0_6access13address_spaceE1EEE.cl::sycl::multi_ptr" addrspace(4)*, align 8
  store %"class._ZTSN2cl4sycl9multi_ptrINS0_3vecIsLi2EEELNS0_6access13address_spaceE1EEE.cl::sycl::multi_ptr" addrspace(4)* %this, %"class._ZTSN2cl4sycl9multi_ptrINS0_3vecIsLi2EEELNS0_6access13address_spaceE1EEE.cl::sycl::multi_ptr" addrspace(4)** %this.addr, align 8, !tbaa !8
  %this1 = load %"class._ZTSN2cl4sycl9multi_ptrINS0_3vecIsLi2EEELNS0_6access13address_spaceE1EEE.cl::sycl::multi_ptr" addrspace(4)*, %"class._ZTSN2cl4sycl9multi_ptrINS0_3vecIsLi2EEELNS0_6access13address_spaceE1EEE.cl::sycl::multi_ptr" addrspace(4)** %this.addr, align 8
  %m_Pointer = getelementptr inbounds %"class._ZTSN2cl4sycl9multi_ptrINS0_3vecIsLi2EEELNS0_6access13address_spaceE1EEE.cl::sycl::multi_ptr", %"class._ZTSN2cl4sycl9multi_ptrINS0_3vecIsLi2EEELNS0_6access13address_spaceE1EEE.cl::sycl::multi_ptr" addrspace(4)* %this1, i32 0, i32 0
  store %"class._ZTSN2cl4sycl3vecIsLi2EEE.cl::sycl::vec" addrspace(1)* null, %"class._ZTSN2cl4sycl3vecIsLi2EEE.cl::sycl::vec" addrspace(1)* addrspace(4)* %m_Pointer, align 8, !tbaa !23
  ret void
}

If I understand correctly, this confuses some of optimization within our OpenCL runtime.

Summarizing: reproducer contains a bug which causes crash. Issue is fixed on the latest internal build, and fix will be available in the next release.

Please let me know if you have any questions or spot another bugs

[agozillon]

Hi @AlexeySachkov thank you very much for the reply. You are correct it's not initialized, that was more me being lazy at the time I created the test and trying to shrink the example! I didn't realize it would actually impact the IR in this case. My apologies for being misleading! I should have been more careful.

A reproducer more true to the source would be something like below:

template <typename T>
inline float get(T v) {
  return v[0].x();
}
  
template<typename T>
float4 do_something(T v) {
  const float tnear = cl::sycl::fmax(0.1f, 1.0f);
  const float tfar  = cl::sycl::fmin(2.0f, 3.0f);

  if (tnear < tfar) {
     float t        = tnear;
     float f_tt     = 0;
    
     for (; t < tfar; t += 1.0f) {
       f_tt = get(v);
       if (f_tt < 0)
         break;
     }
     t = t + 1.0f * f_tt / (1.0f - f_tt);
     return float4{0.1f,0.1f,0.1f,t};
  }

  return float4{0,0,0,0};
}

int main() {
   queue q;    
   buffer<float4> ob{range<1>{1}};
   buffer<short2,1> *ib = new buffer<short2>(range<1>{2});

   q.submit([&](handler &cgh) {      
      auto wb = ob.get_access<access::mode::write>(cgh);
      auto rwb = ib->get_access<access::mode::read_write>(cgh);

      cgh.single_task<codegen_or_jit_bug>([=]() {
        cl::sycl::global_ptr<short2> v = rwb;
        wb[0] = do_something(v);
      });
  });
     
  return 0;
}

Where I believe the data is initialized via new in a buffer external to the kernel and then passed in via accessor. In this case there appears to be no null when I search through the LLVM-IR. It still seems to be a very related issue and is also perhaps fixed in your internal build however (or perhaps just an environment issue on my end)!

Hopefully in this slightly updated version of the reproducer I've not done something silly again, it's an attempt at extracting the problem from this much larger kernel: https://github.com/agozillon/syclslambench/blob/master/kfusion/src/sycl/kernels.cpp#L963

[AlexeySachkov]

Hi @agozillon, thanks for the updated reproducer.

I can confirm that new reproducer crashes on the latest oclcpuexp-2019.9.11.0.1106_rel:

Stack dump:
0.      Running pass 'Function Pass Manager' on module 'main'.
1.      Running pass 'VPlan Vectorization Driver' on function ''
Segmentation fault (core dumped)

Also, as I can see, it works fine on the latest internal build. So, stay tuned for the next release!
Feel free to leave this issue open until your problem is actually fixed

Just to confirm: am I right that you have some work-around (via slight changes in code)?

[agozillon]

Ah that's good news! Thanks again @AlexeySachkov for the help!

At the moment I don't have a work-around unfortunately, however it's not a blocking issue as the benchmark works for Intel GPU at the moment, so I was happy to wait on the updated runtime rather than seek a work-around.

[AlexeySachkov]

Hi @agozillon,

We have recently released newer version of Experimental Intel(R) CPU Runtime for OpenCL(TM) Applications with SYCL support: oclcpuexp-2019.9.12.0.12_rel: I've checked the updated reproducer and it seems to work fine with it.

Could you please double-check on your side and close the issue if problem disappeared?

Moreover, using the latest SYCL compiler, it works even with previous release (oclcpuexp-2019.9.11.0.1106_rel)

[agozillon]

Thanks for the update @AlexeySachkov it LGTM, I've tested a ToT version of the compiler and the new runtime and the issue is fixed for me as well. No other errors to report either, it now compiles and runs the benchmark as expected for Intel CPU and GPU, thanks for the awesome work guys!

As an off-topic aside I'm not sure if the documentation in the getting started guide for the Intel OpenCL CPU runtime setup might need a small update, I had to create these two other symbolic links with the new 2020 TBB release:

ln -s /opt/intel/tbb/lib/intel64/gcc4.8/libtbb.so.2 /opt/intel/oclcpuexp/x64/libtbb.so.2
ln -s /opt/intel/tbb/lib/intel64/gcc4.8/libtbbmalloc.so.2 /opt/intel/oclcpuexp/x64/libtbbmalloc.so.2

To get the runtime to function correctly, but perhaps I got one of the prior steps incorrect and this was just a weird fix/workaround I've made.

[AlexeySachkov]

Tagging @vladimirlaz to comment about TBB symlinks