Merge from 'sycl' to 'sycl-web'

Author: iclsrc

.github/workflows/scorecard.yml

@@ -57,6 +57,6 @@ jobs:
 
       # Upload the results to GitHub's code scanning dashboard.
       - name: "Upload to code-scanning"
-        uses: github/codeql-action/upload-sarif@1b1aada464948af03b950897e5eb522f92603cc2 # v3.24.9
+        uses: github/codeql-action/upload-sarif@d39d31e687223d841ef683f52467bd88e9b21c14 # v3.25.3
         with:
           sarif_file: results.sarif

.github/workflows/sycl-issues-ping-assignee.yml

@@ -0,0 +1,64 @@
+name: Ping issue assignees
+# We have some specific pool of the issues we would like to handle. Sometimes
+# the issue from this pool has an assignee, but doesn't get any updates for a
+# long time. In this case it'd be useful to periodically ping the assignee.
+
+# Note: may be we could use "actions/stale@v*", but I'm not sure if it's
+# possible to not set the "stale" label at all. Even so, this action will not
+# ping the assignee of the "stale" issue more than onсe.
+
+# Note2: probably it'd be useful to have a small doc describing this "specific
+# pool" to refer to.
+
+on:
+  schedule:
+    - cron: '0 0 * * *'
+
+jobs:
+  run:
+    runs-on: ubuntu-20.04
+    env:
+      GH_TOKEN: ${{ secrets.GITHUB_TOKEN }}
+      REPO: ${{ github.repository }}
+    steps:
+    # List specific issues with an assignee but without recent updates.
+    # Note: for some reason gh returns 0 results if the "assignee:*" filter is
+    # added, so we have to manually filter the results.
+    - name: Get the specicifc list of issues
+      run: |
+          gh issue list --search '-label:"help wanted" -label:cuda \
+          -label:confirmed -label:hip -label:sycl-mlir -label:upstream is:open \
+          -label:genx -label:sycl-bindless-images -label:sycl-graph \
+          -label:native-cpu' --limit 200 --json assignees --json number \
+          --json updatedAt \
+          -R https://github.com/${{ env.REPO }}.git > issues.json
+
+    - name: Filter issues and ping
+      run: |
+          days_to_stale=60
+          current_time=$(date +%s)
+
+          cat issues.json | jq -c '.[]' | while read -r issue; do
+            assignees=$(echo "$issue" | jq '.assignees | length')
+            [ "$assignees" -eq 0 ] && continue
+
+            updated_at=$(echo "$issue" | jq -r '.updatedAt')
+            updated_at_seconds=$(date -d "$updated_at" +%s)
+            difference_days=$(( (current_time - updated_at_seconds) / 86400 ))
+            [ "$difference_days" -lt $days_to_stale ] && continue
+
+            issue_number=$(echo "$issue" | jq '.number')
+            assignee_logins=$(echo "$issue" | jq -r '.assignees[].login' | sed 's/^/@/' | paste -s -d ' ' -)
+            comment_body="Hi! There have been no updates for at least the last $days_to_stale days, though the issue has assignee(s).
+
+          $assignee_logins, could you please take one of the following actions:
+          - provide an update if you have any
+          - unassign yourself if you're not looking / going to look into this issue
+          - mark this issue with the 'confirmed' label if you have confirmed the problem/request and our team should work on it
+          - close the issue if it has been resolved
+          - take any other suitable action.
+
+          Thanks!"
+
+            gh issue comment $issue_number -R https://github.com/${{ env.REPO }}.git -b "$comment_body" >> $GITHUB_STEP_SUMMARY
+          done

clang/lib/CodeGen/CGVTT.cpp

@@ -90,6 +90,10 @@ CodeGenVTables::EmitVTTDefinition(llvm::GlobalVariable *VTT,
      llvm::Constant *Init = llvm::ConstantExpr::getGetElementPtr(
          VTable->getValueType(), VTable, Idxs, /*InBounds=*/true, InRange);
 
+     if (CGM.getTriple().isSPIR())
+       Init = llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast(
+           Init, CGM.DefaultInt8PtrTy);
+
      VTTComponents.push_back(Init);
   }
 

clang/lib/CodeGen/CGVTables.cpp

@@ -770,6 +770,19 @@ void CodeGenVTables::addVTableComponent(ConstantArrayBuilder &builder,
       // Method is acceptable, continue processing as usual.
     }
 
+    if (CGM.getLangOpts().SYCLIsDevice) {
+      // Some virtual function may be only valid in host code. Don't emit
+      // virtual functions which were not specifically marked as device code.
+      const CXXMethodDecl *MD = cast<CXXMethodDecl>(GD.getDecl());
+      // FIXME: Update check to use compile-time property instead of
+      // SYCLDeviceIndirectlyCallableAttr once Virtual function extension
+      // specification is accepted.
+      if (!MD->hasAttr<SYCLDeviceIndirectlyCallableAttr>() &&
+          !MD->hasAttr<SYCLDeviceAttr>())
+        return builder.add(
+            llvm::ConstantExpr::getNullValue(CGM.DefaultInt8PtrTy));
+    }
+
     auto getSpecialVirtualFn = [&](StringRef name) -> llvm::Constant * {
       // FIXME(PR43094): When merging comdat groups, lld can select a local
       // symbol as the signature symbol even though it cannot be accessed

clang/lib/CodeGen/CodeGenModule.cpp

@@ -393,6 +393,17 @@ CodeGenModule::CodeGenModule(ASTContext &C,
       LLVMContext, C.getTargetAddressSpace(GetGlobalConstantAddressSpace()));
   ASTAllocaAddressSpace = getTargetCodeGenInfo().getASTAllocaAddressSpace();
 
+  if (getTriple().isSPIR()) {
+    // Currently code uses GlobalsInt8PtrTy for virtual table elements but for
+    // SPIR-V targets default address space pointers are needed.
+    GlobalsInt8PtrTy = DefaultInt8PtrTy;
+    // Pointer to runtime globals such as virtual tables.
+    RuntimeGlobalsInt8PtrTy = Int8Ty->getPointerTo(
+        getContext().getTargetAddressSpace(LangAS::opencl_global));
+  } else {
+    RuntimeGlobalsInt8PtrTy = GlobalsInt8PtrTy;
+  }
+
   // Build C++20 Module initializers.
   // TODO: Add Microsoft here once we know the mangling required for the
   // initializers.
@@ -5537,8 +5548,10 @@ llvm::GlobalVariable *CodeGenModule::CreateOrReplaceCXXRuntimeVariable(
   }
 
   // Create a new variable.
-  GV = new llvm::GlobalVariable(getModule(), Ty, /*isConstant=*/true,
-                                Linkage, nullptr, Name);
+  GV = new llvm::GlobalVariable(getModule(), Ty, /*isConstant=*/true, Linkage,
+                                nullptr, Name, nullptr,
+                                llvm::GlobalValue::NotThreadLocal,
+                                RuntimeGlobalsInt8PtrTy->getAddressSpace());
 
   if (OldGV) {
     // Replace occurrences of the old variable if needed.

clang/lib/CodeGen/CodeGenTypeCache.h

@@ -69,6 +69,9 @@ struct CodeGenTypeCache {
   /// void* in target address space
   llvm::PointerType *DefaultInt8PtrTy;
 
+  /// void* used for virtual table globals.
+  llvm::PointerType *RuntimeGlobalsInt8PtrTy;
+
   /// void* in default globals address space
   union {
     llvm::PointerType *GlobalsVoidPtrTy;

clang/test/CodeGenSYCL/simple-sycl-virtual-function.cpp

@@ -1,21 +1,27 @@
-// Test verifying that RTTI information is not emitted during SYCL device compilation.
+// Test verifies the following
+//  1. RTTI information is not emitted during SYCL device compilation.
+//  2. Virtual table elements are generated in AS4.
+//  3. Runtime Global Variables are generated in AS1.
 
 // RUN: %clang_cc1 -triple spir64 -fsycl-allow-virtual-functions -fsycl-is-device -emit-llvm %s -o - | FileCheck %s --implicit-check-not _ZTI4Base --implicit-check-not _ZTI8Derived1 -check-prefix VTABLE
-// RUN: %clang_cc1 -triple spir64 -fsycl-allow-virtual-functions -fsycl-is-device -fexperimental-relative-c++-abi-vtables -emit-llvm %s -o - | FileCheck %s --implicit-check-not _ZTI4Base --implicit-check-not _ZTI8Derived1
+// RUNx: %clang_cc1 -triple spir64 -fsycl-allow-virtual-functions -fsycl-is-device -fexperimental-relative-c++-abi-vtables -emit-llvm %s -o - | FileCheck %s --implicit-check-not _ZTI4Base --implicit-check-not _ZTI8Derived1
 
-// VTABLE: @_ZTV8Derived1 = linkonce_odr unnamed_addr constant { [3 x ptr] } { [3 x ptr] [ptr null, ptr null, ptr @_ZN8Derived17displayEv] }, comdat, align 8
-// VTABLE: @_ZTV4Base = linkonce_odr unnamed_addr constant { [3 x ptr] } { [3 x ptr] [ptr null, ptr null, ptr @_ZN4Base7displayEv] }, comdat, align 8
+// Since experimental-relative-c++-abi-vtables is some experimental option, temporary disabling the check for now
+// until we emit proper address spaces (and casts) everywhere.
+
+// VTABLE: @_ZTV8Derived1 = linkonce_odr unnamed_addr addrspace(1) constant { [3 x ptr addrspace(4)] } { [3 x ptr addrspace(4)] [ptr addrspace(4) null, ptr addrspace(4) null, ptr addrspace(4) addrspacecast (ptr @_ZN8Derived17displayEv to ptr addrspace(4))] }, comdat, align 8
+// VTABLE: @_ZTV4Base = linkonce_odr unnamed_addr addrspace(1) constant { [3 x ptr addrspace(4)] } { [3 x ptr addrspace(4)] [ptr addrspace(4) null, ptr addrspace(4) null, ptr addrspace(4) addrspacecast (ptr @_ZN4Base7displayEv to ptr addrspace(4))] }, comdat, align 8
 
 SYCL_EXTERNAL bool rand();
 
 class Base {
    public:
-    virtual void display() {}
+    [[intel::device_indirectly_callable]] virtual void display() {}
 };
 
 class Derived1 : public Base {
    public:
-    void display() {}
+    [[intel::device_indirectly_callable]] void display() override {}
 };
 
 SYCL_EXTERNAL void test() {

clang/test/CodeGenSYCL/sycl-device-only-virtual-functions.cpp

@@ -0,0 +1,145 @@
+// Purpose of this test is to ensure that in SYCL device code mode, we are only
+// emitting those virtual functions into vtable, which are marked with
+// [[intel::device_indirectly_callable]] attribute or SYCL_EXTERNAL macro.
+//
+// RUN: %clang_cc1 -emit-llvm -o - -fsycl-is-device \
+// RUN:     -fsycl-allow-virtual-functions -internal-isystem %S/Inputs \
+// RUN:     -triple spir64 %s -o %t.ll
+// RUN: FileCheck %s --input-file %t.ll --implicit-check-not host \
+// RUN:     --implicit-check-not _ZN8Derived416maybe_device_barEv
+//
+// Check that vtable contains null pointers for host-only virtual functions:
+//
+// CHECK: @_ZTV4Base =
+// CHECK-SAME:, {{.*}} @_ZN4Base10device_fooEv
+// CHECK-SAME:, {{.*}} @_ZN4Base10device_barEv
+// CHECK-SAME:, ptr addrspace(4) null
+// CHECK-SAME:, ptr addrspace(4) null
+//
+// CHECK: @_ZTV8Derived1 =
+// CHECK-SAME:, {{.*}} @_ZN8Derived110device_fooEv
+// CHECK-SAME:, {{.*}} @_ZN4Base10device_barEv
+// CHECK-SAME:, ptr addrspace(4) null
+// CHECK-SAME:, ptr addrspace(4) null
+// CHECK-SAME:, {{.*}} @_ZN8Derived110device_bazEv
+// CHECK-SAME:, ptr addrspace(4) null
+//
+// CHECK: @_ZTV8Derived2 =
+// CHECK-SAME:, {{.*}} @_ZN4Base10device_fooEv
+// CHECK-SAME:, {{.*}} @_ZN8Derived210device_barEv
+// CHECK-SAME:, ptr addrspace(4) null
+// CHECK-SAME:, ptr addrspace(4) null
+//
+// CHECK: @_ZTV10SubDerived =
+// CHECK-SAME:, {{.*}} @_ZN8Derived110device_fooEv
+// CHECK-SAME:, {{.*}} @_ZN4Base10device_barEv
+// CHECK-SAME:, ptr addrspace(4) null
+// CHECK-SAME:, ptr addrspace(4) null
+// CHECK-SAME:, {{.*}} @_ZN10SubDerived10device_bazEv
+// CHECK-SAME:, ptr addrspace(4) null
+//
+// CHECK: @_ZTV8Derived3 =
+// CHECK-SAME:, {{.*}} @_ZN8Derived310device_fooEv
+// CHECK-SAME:, {{.*}} @_ZN8Derived316maybe_device_barEv
+// CHECK-SAME:, ptr addrspace(4) null
+//
+// CHECK: @_ZTV12AbstractBase =
+// CHECK-SAME: zeroinitializer
+//
+// CHECK: @_ZTV8Derived4 =
+// CHECK-SAME:, {{.*}} @_ZN8Derived410device_fooEv
+// CHECK-SAME:, ptr addrspace(4) null
+// CHECK-SAME:, ptr addrspace(4) null
+//
+// Check that bodies of device virtual functions are present:
+//
+// CHECK-DAG: define linkonce_odr spir_func void @_ZN4Base10device_fooEv
+// CHECK-DAG: define linkonce_odr spir_func void @_ZN4Base10device_barEv
+//
+// CHECK-DAG: define linkonce_odr spir_func void @_ZN8Derived110device_fooEv
+// CHECK-DAG: define linkonce_odr spir_func void @_ZN8Derived110device_bazEv
+//
+// CHECK-DAG: define linkonce_odr spir_func void @_ZN10SubDerived10device_bazEv
+//
+// CHECK-DAG: define linkonce_odr spir_func void @_ZN8Derived210device_barEv
+//
+// CHECK-DAG: define linkonce_odr spir_func void @_ZN8Derived310device_fooEv
+// CHECK-DAG: define linkonce_odr spir_func void @_ZN8Derived316maybe_device_barEv
+//
+// CHECK-DAG: define linkonce_odr spir_func void @_ZN8Derived410device_fooEv
+
+#include "sycl.hpp"
+
+struct Base {
+  [[intel::device_indirectly_callable]] virtual void device_foo() {}
+  [[intel::device_indirectly_callable]] virtual void device_bar() {}
+
+  virtual void host_foo() {}
+  virtual void host_bar() {}
+};
+
+struct Derived1 : public Base {
+  [[intel::device_indirectly_callable]] void device_foo() override {}
+
+  void host_bar() override {}
+
+  [[intel::device_indirectly_callable]] virtual void device_baz() {}
+  virtual void host_baz() {}
+};
+
+struct SubDerived : public Derived1 {
+  [[intel::device_indirectly_callable]] void device_baz() override {}
+  void host_baz() override {}
+};
+
+struct Derived2 : public Base {
+  [[intel::device_indirectly_callable]] void device_bar() override {}
+
+  void host_foo() override {}
+};
+
+class AbstractBase {
+  virtual void device_foo() = 0;
+  virtual void maybe_device_bar() = 0;
+  virtual void host_foo() = 0;
+};
+
+class Derived3 : public AbstractBase {
+  SYCL_EXTERNAL void device_foo() override {}
+  SYCL_EXTERNAL void maybe_device_bar() override {}
+  void host_foo() override {}
+};
+
+class Derived4 : public AbstractBase {
+  SYCL_EXTERNAL void device_foo() override {}
+  void host_foo() override {}
+  void maybe_device_bar() override {}
+};
+
+int main(int argc, char *argv[]) {
+  sycl::kernel_single_task<class kernel_function>([=]() {
+    Base b;
+    Derived1 d1;
+    Derived2 d2;
+    SubDerived sd;
+    Base *ptr;
+    if (argc > 5) {
+      ptr = &d1;
+    } else if (argc > 42) {
+      ptr = &d2;
+    } else {
+      ptr = &sd;
+    }
+
+    Derived3 d3;
+    Derived4 d4;
+    AbstractBase *aptr;
+    if (argc > 5) {
+      aptr = &d3;
+    } else {
+      aptr = &d4;
+    }
+  });
+
+  return 0;
+}