Merge branch 'sycl' of https://github.com/intel/llvm into op-forward-ann-arg

Author: Chen, Brox

.github/workflows/sycl_linux_precommit.yml

@@ -8,6 +8,7 @@ on:
     branches:
     - sycl
     - sycl-devops-pr/**
+    - sycl-rel-**
     # Do not run builds if changes are only in the following locations
     paths-ignore:
     - '.github/ISSUE_TEMPLATE/**'

.github/workflows/sycl_post_commit.yml

@@ -5,6 +5,7 @@ on:
     branches:
     - sycl
     - sycl-devops-pr/**
+    - sycl-rel-**
 
   pull_request:
     branches:

.github/workflows/sycl_windows_precommit.yml

@@ -6,6 +6,7 @@ on:
     - sycl
     - sycl-devops-pr/**
     - llvmspirv_pulldown
+    - sycl-rel-**
     # Do not run builds if changes are only in the following locations
     paths-ignore:
     - '.github/ISSUE_TEMPLATE/**'

clang-tools-extra/clangd/Compiler.cpp

@@ -89,6 +89,12 @@ void disableUnsupportedOptions(CompilerInvocation &CI) {
   CI.getLangOpts().ProfileListFiles.clear();
   CI.getLangOpts().XRayAlwaysInstrumentFiles.clear();
   CI.getLangOpts().XRayNeverInstrumentFiles.clear();
+  if (CI.getLangOpts().SYCLIsDevice) {
+    CI.getLangOpts().SYCLIsDevice = false;
+    CI.getLangOpts().SYCLUnnamedLambda = false;
+    CI.getLangOpts().DeclareSPIRVBuiltins = false;
+    CI.getTargetOpts().Triple = CI.getTargetOpts().HostTriple;
+  }
 }
 
 std::unique_ptr<CompilerInvocation>

clang-tools-extra/clangd/test/sycl.test

@@ -0,0 +1,40 @@
+# RUN: rm -rf %t.dir/* && mkdir -p %t.dir
+# RUN: echo '[{"directory": "%/t.dir", "command": "clang++ -fsycl main.cpp", "file": "main.cpp"}]' > %t.dir/compile_commands.json
+# RUN: sed -e "s|INPUT_DIR|%/t.dir|g" %s > %t.test.1
+
+# On Windows, we need the URI in didOpen to look like "uri":"file:///C:/..."
+# (with the extra slash in the front), so we add it here.
+# RUN: sed -E -e 's|"file://([A-Z]):/|"file:///\1:/|g' %t.test.1 > %t.test
+
+# RUN: clangd -lit-test < %t.test | FileCheck -strict-whitespace %t.test
+
+{"jsonrpc":"2.0","id":0,"method":"initialize","params":{"initializationOptions":{"compilationDatabasePath":"INPUT_DIR"}}}
+---
+{"jsonrpc":"2.0","method":"textDocument/didOpen","params":{"textDocument":{"uri":"test:///main.cpp","languageId":"cpp","version":1,"text":"#include <sycl/sycl.hpp>\nsycl::queue q{};"}}}
+---
+{"jsonrpc":"2.0","id":1,"method":"textDocument/symbolInfo","params":{
+  "textDocument":{"uri":"test:///main.cpp"},
+  "position":{"line":1,"character":13}
+}}
+#      CHECK:  "id": 1
+# CHECK-NEXT:  "jsonrpc": "2.0",
+# CHECK-NEXT:  "result": [
+# CHECK-NEXT:    {
+# CHECK-NEXT:      "containerName": "sycl::queue::",
+# CHECK-NEXT:      "declarationRange": {
+# CHECK-NEXT:        "range": {
+# CHECK-NEXT:          "end": {
+# CHECK-NEXT:            "character": 16,
+# CHECK-NEXT:            "line": 124
+# CHECK-NEXT:          },
+# CHECK-NEXT:          "start": {
+# CHECK-NEXT:            "character": 11,
+# CHECK-NEXT:            "line": 124
+# CHECK-NEXT:          }
+# CHECK-NEXT:        },
+# CHECK-NEXT:        "uri": "file://{{.*}}/include/sycl/queue.hpp"
+# CHECK-NEXT:      },
+---
+{"jsonrpc":"2.0","id":3,"method":"shutdown"}
+---
+{"jsonrpc":"2.0","method":"exit"}

clang/lib/Sema/SemaDeclAttr.cpp

@@ -7520,9 +7520,7 @@ static void handleSYCLIntelDoublePumpAttr(Sema &S, Decl *D,
 
 /// Handle the [[intel::fpga_memory]] attribute.
 /// This is incompatible with the [[intel::fpga_register]] attribute.
-static void handleSYCLIntelMemoryAttr(Sema &S, Decl *D,
-                                  const ParsedAttr &AL) {
-  checkForDuplicateAttribute<SYCLIntelMemoryAttr>(S, D, AL);
+static void handleSYCLIntelMemoryAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   if (checkAttrMutualExclusion<SYCLIntelRegisterAttr>(S, D, AL))
     return;
 
@@ -7543,10 +7541,20 @@ static void handleSYCLIntelMemoryAttr(Sema &S, Decl *D,
     }
   }
 
-  // We are adding a user memory attribute, drop any implicit default.
-  if (auto *MA = D->getAttr<SYCLIntelMemoryAttr>())
-    if (MA->isImplicit())
-      D->dropAttr<SYCLIntelMemoryAttr>();
+  if (auto *MA = D->getAttr<SYCLIntelMemoryAttr>()) {
+    // Check to see if there's a duplicate memory attribute with different
+    // values already applied to the declaration.
+    if (!MA->isImplicit()) {
+      if (MA->getKind() != Kind) {
+        S.Diag(AL.getLoc(), diag::warn_duplicate_attribute) << &AL;
+        S.Diag(MA->getLocation(), diag::note_previous_attribute);
+      }
+      // Drop the duplicate attribute.
+      return;
+    }
+    // We are adding a user memory attribute, drop any implicit default.
+    D->dropAttr<SYCLIntelMemoryAttr>();
+  }
 
   D->addAttr(::new (S.Context) SYCLIntelMemoryAttr(S.Context, AL, Kind));
 }

clang/test/SemaSYCL/intel-fpga-local-ast.cpp

@@ -265,6 +265,46 @@ void check_ast()
   [[intel::merge("mrg6", "depth")]]
   [[intel::merge("mrg6", "depth")]] unsigned int mrg_mrg6[4];
 
+  // Check to see if there's a duplicate attribute with different values
+  // already applied to the declaration. Drop the duplicate attribute.
+  // CHECK: VarDecl{{.*}}mem_block_ram
+  // CHECK: SYCLIntelMemoryAttr{{.*}}MLAB
+  // CHECK-NOT: SYCLIntelMemoryAttr
+  [[intel::fpga_memory("MLAB")]]
+  [[intel::fpga_memory("BLOCK_RAM")]] unsigned int mem_block_ram[32];
+
+  // Check to see if there's a duplicate attribute with same Default values
+  // already applied to the declaration. Drop the duplicate attribute.
+  // CHECK: VarDecl{{.*}}mem_memory
+  // CHECK: SYCLIntelMemoryAttr{{.*}}Default
+  // CHECK-NOT: SYCLIntelMemoryAttr
+  [[intel::fpga_memory]]
+  [[intel::fpga_memory]] unsigned int mem_memory[64];
+
+  // Check to see if there's a duplicate attribute with same values
+  // already applied to the declaration. Drop the duplicate attribute.
+  // CHECK: VarDecl{{.*}}mem_memory_block
+  // CHECK: SYCLIntelMemoryAttr{{.*}}BlockRAM
+  // CHECK-NOT: SYCLIntelMemoryAttr
+  [[intel::fpga_memory("BLOCK_RAM")]]
+  [[intel::fpga_memory("BLOCK_RAM")]] unsigned int mem_memory_block[64];
+
+  // Check to see if there's a duplicate attribute with different values
+  // already applied to the declaration. Drop the duplicate attribute.
+  // CHECK: VarDecl{{.*}}mem_mlabs
+  // CHECK: SYCLIntelMemoryAttr{{.*}}Default
+  // CHECK-NOT: SYCLIntelMemoryAttr
+  [[intel::fpga_memory]]
+  [[intel::fpga_memory("MLAB")]] unsigned int mem_mlabs[64];
+
+  // Check to see if there's a duplicate attribute with different values
+  // already applied to the declaration. Drop the duplicate attribute.
+  // CHECK: VarDecl{{.*}}mem_mlabs_block_ram
+  // CHECK: SYCLIntelMemoryAttr{{.*}}BlockRAM
+  // CHECK-NOT: SYCLIntelMemoryAttr
+  [[intel::fpga_memory("BLOCK_RAM")]]
+  [[intel::fpga_memory]] unsigned int mem_mlabs_block_ram[64];
+
   // FIXME: Duplicate attribute should be ignored.
   // Both are applied at the moment.
   //CHECK: VarDecl{{.*}}bb_bb

clang/test/SemaSYCL/intel-fpga-local.cpp

@@ -143,12 +143,45 @@ void diagnostics()
   //expected-note@-2 {{conflicting attribute is here}}
   unsigned int mem_reg[64];
 
-  //expected-warning@+1{{attribute 'fpga_memory' is already applied}}
-  [[intel::fpga_memory]] [[intel::fpga_memory]] unsigned int mem_mem[64];
+  // Check to see if there's a duplicate attribute with same Default values
+  // already applied to the declaration.
+  // No diagnostic is emitted because the arguments match.
+  [[intel::fpga_memory]]
+  [[intel::fpga_memory]] unsigned int mem_mem[64]; // OK
 
   //expected-warning@+1 {{unknown attribute 'memory' ignored}}
   [[intelfpga::memory]] unsigned int memory_var[64];
 
+  // Check to see if there's a duplicate attribute with different values
+  // already applied to the declaration.
+  // Diagnostic is emitted because the arguments mismatch.
+  //expected-warning@+2{{attribute 'fpga_memory' is already applied with different arguments}}
+  [[intel::fpga_memory("MLAB")]] // expected-note {{previous attribute is here}}
+  [[intel::fpga_memory("BLOCK_RAM")]] unsigned int mem_block_ram[32];
+
+  // Check to see if there's a duplicate attribute with same values
+  // already applied to the declaration.
+  // No diagnostic is emitted because the arguments match.
+  [[intel::fpga_memory("MLAB")]]
+  [[intel::fpga_memory("MLAB")]] unsigned int mem_mlab[32]; // OK
+
+  [[intel::fpga_memory("BLOCK_RAM")]]
+  [[intel::fpga_memory("BLOCK_RAM")]] unsigned int mem_block[32]; // OK
+
+  // Check to see if there's a duplicate attribute with different values
+  // already applied to the declaration.
+  // Diagnostic is emitted because the arguments mismatch.
+  //expected-warning@+2{{attribute 'fpga_memory' is already applied with different arguments}}
+  [[intel::fpga_memory]] // expected-note {{previous attribute is here}}
+  [[intel::fpga_memory("MLAB")]] unsigned int mem_mlabs[64];
+
+  // Check to see if there's a duplicate attribute with different values
+  // already applied to the declaration.
+  // Diagnostic is emitted because the arguments mismatch.
+  //expected-warning@+2{{attribute 'fpga_memory' is already applied with different arguments}}
+  [[intel::fpga_memory("BLOCK_RAM")]] // expected-note {{previous attribute is here}}
+  [[intel::fpga_memory]] unsigned int mem_mlabs_block_ram[64];
+
   // **bankwidth
   //expected-warning@+1 {{unknown attribute 'bankwidth' ignored}}
   [[intelfpga::bankwidth(4)]] unsigned int bankwidth_var[32];

llvm-spirv/test/FPFastMathModeNotNaNFast.spvasm

@@ -13,9 +13,10 @@
                OpName %r3 "r3"
                OpName %r4 "r4"
                OpName %r5 "r5"
+               OpName %r6 "r6"
                OpDecorate %dec FPFastMathMode NotNaN|Fast
         %dec = OpDecorationGroup
-               OpGroupDecorate %dec %r1 %r2 %r3 %r4 %r5
+               OpGroupDecorate %dec %r1 %r2 %r3 %r4 %r5 %r6
        %void = OpTypeVoid
       %float = OpTypeFloat 32
           %5 = OpTypeFunction %void %float %float
@@ -28,6 +29,7 @@
          %r3 = OpFMul %float %a %b
          %r4 = OpFDiv %float %a %b
          %r5 = OpFRem %float %a %b
+         %r6 = OpFNegate %float %a
                OpReturn
                OpFunctionEnd
 
@@ -36,3 +38,4 @@
 ; CHECK: %r3 = fmul fast float %a, %b
 ; CHECK: %r4 = fdiv fast float %a, %b
 ; CHECK: %r5 = frem fast float %a, %b
+; CHECK: %r6 = fneg fast float %a