Missing function in hierarchical when targetting PTX

[steffenlarsen]

Compiling

#include <CL/sycl.hpp>

using namespace cl::sycl;

int main() {
  queue q(default_selector().select_device());
  q.submit([&](handler &cgh) {
    cgh.parallel_for_work_group(range<1>(1), range<1>(1), [=](group<1> g) {
      g.parallel_for_work_item(
          [=](h_item<1> i) {});
    });
  });
  return 0;
}

using the nvptx64-nvidia-cuda-sycldevice triple causes the following error:

Call parameter type does not match function signature!
%"class._ZTSN2cl4sycl5groupILi1EEE.cl::sycl::group" addrspace(3)* @ArgShadow
 %"class._ZTSN2cl4sycl5groupILi1EEE.cl::sycl::group"*  call void @"_ZNK2cl4sycl5groupILi1EE22parallel_for_work_itemIZZZ4mainENK3$_0clERNS0_7handlerEENKUlS2_E_clES2_EUlNS0_6h_itemILi1EEEE_EEvT_"(%"class._ZTSN2cl4sycl5groupILi1EEE.cl::sycl::group" addrspace(3)* @ArgShadow, %"class._ZTSZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlNS1_5groupILi1EEEE_clES5_EUlNS1_6h_itemILi1EEEE_.anon"* byval(%"class._ZTSZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlNS1_5groupILi1EEEE_clES5_EUlNS1_6h_itemILi1EEEE_.anon") align 1 %agg.tmp)
Call parameter type does not match function signature!
%"class._ZTSN2cl4sycl5groupILi1EEE.cl::sycl::group" addrspace(3)* @ArgShadow
 %"class._ZTSN2cl4sycl5groupILi1EEE.cl::sycl::group"*  call void @"_ZNK2cl4sycl5groupILi1EE22parallel_for_work_itemIZZZ4mainENK3$_0clERNS0_7handlerEENKUlS2_E_clES2_EUlNS0_6h_itemILi1EEEE_EEvT_"(%"class._ZTSN2cl4sycl5groupILi1EEE.cl::sycl::group" addrspace(3)* @ArgShadow, %"class._ZTSZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlNS1_5groupILi1EEEE_clES5_EUlNS1_6h_itemILi1EEEE_.anon"* byval(%"class._ZTSZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlNS1_5groupILi1EEEE_clES5_EUlNS1_6h_itemILi1EEEE_.anon") align 1 %agg.tmp)
in function _ZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlNS1_5groupILi1EEEE_clES5_
fatal error: error in backend: Broken function found, compilation aborted!
clang-9: error: clang frontend command failed with exit code 70 (use -v to see invocation)

[pvchupin]

@againull, can you take a look please?

[againull]

@againull, can you take a look please?

Yes, I will work on it.

[againull]

Hi @steffenlarsen
I get the following error on your test case:

test.cpp:8:59: error: kernel name class and its template argument classes' declarations can only nest in a namespace: '(lambda at test.cpp:7:12)'
    cgh.parallel_for_work_group(range<1>(1), range<1>(1), [=](group<1> g) {
                                                          ^

And problem is that kernel name class is not specified.
After the following change test is compiled successfully:

diff --git a/test.cpp b/test.cpp
index 0b5f385..857890f 100644
--- a/test.cpp
+++ b/test.cpp
@@ -5,7 +5,7 @@ using namespace cl::sycl;
 int main() {
   queue q(default_selector().select_device());
   q.submit([&](handler &cgh) {
-    cgh.parallel_for_work_group(range<1>(1), range<1>(1), [=](group<1> g) {
+    cgh.parallel_for_work_group<class example_kernel>(range<1>(1), range<1>(1), [=](group<1> g) {
       g.parallel_for_work_item(
           [=](h_item<1> i) {});
     });

Please close the issue if you agree.

[steffenlarsen]

I don't know why that error didn't occur before, but the error I reported is still there with the changes you suggested. It doesn't happen if sycl-targets is not nvptx64-nvidia-cuda-sycldevice.

As a note, there is currently another error (#1342) which can be circumvented by setting the include path explicitly.

[againull]

I am aware of #1342.
I found out how to reproduce the error. You haven't mentioned in the description of the issue that build must be with assertions enabled to reproduce the error. If assertions are not enabled then "-disable-llvm-verifier" option is added by clang and there is no error.