[SYCLomatic] Remove -fgpu-exclude-wrong-side-overloads to avoid other parsing errors

clang/lib/Sema/SemaOverload.cpp

@@ -10683,8 +10683,25 @@ bool clang::isBetterOverloadCandidate(
   // to determine which is better.
   if (S.getLangOpts().CUDA && Cand1.Function && Cand2.Function) {
     FunctionDecl *Caller = S.getCurFunctionDecl(/*AllowLambda=*/true);
+#ifdef SYCLomatic_CUSTOMIZATION
+    auto Preference1 = S.CUDA().IdentifyPreference(Caller, Cand1.Function);
+    auto Preference2 = S.CUDA().IdentifyPreference(Caller, Cand2.Function);
+    if (Preference1 != Preference2)
+      return Preference1 > Preference2;
+    // This is a workaround to align to the behavior of nvcc
+    const CXXMethodDecl *MD1 = dyn_cast<CXXMethodDecl>(Cand1.Function);
+    const CXXMethodDecl *MD2 = dyn_cast<CXXMethodDecl>(Cand2.Function);
+    if (MD1 && MD2 && Caller && !Caller->hasAttr<CUDAHostAttr>()) {
+      if ((MD1->isMoveAssignmentOperator() ||
+           MD1->isCopyAssignmentOperator()) &&
+          !(MD2->isMoveAssignmentOperator() || MD2->isCopyAssignmentOperator()))
+        return true;
+    }
+    return false;
+#else
     return S.CUDA().IdentifyPreference(Caller, Cand1.Function) >
            S.CUDA().IdentifyPreference(Caller, Cand2.Function);
+#endif // SYCLomatic_CUSTOMIZATION
   }
 
   // General member function overloading is handled above, so this only handles

clang/lib/Tooling/CommonOptionsParser.cpp

@@ -438,10 +438,6 @@ llvm::Error CommonOptionsParser::init(
     Adjuster = combineAdjusters(
         std::move(Adjuster),
         getInsertArgumentAdjuster("-D__NVCC__", ArgumentInsertPosition::BEGIN));
-    Adjuster = combineAdjusters(
-        std::move(Adjuster),
-        getInsertArgumentAdjuster("-fgpu-exclude-wrong-side-overloads",
-                                  ArgumentInsertPosition::BEGIN));
   }
 #endif // SYCLomatic_CUSTOMIZATION
   AdjustingCompilations->appendArgumentsAdjuster(Adjuster);

clang/lib/Tooling/Tooling.cpp

@@ -901,10 +901,6 @@ int ClangTool::processFiles(llvm::StringRef File, bool &ProcessingFailed,
             std::move(CudaArgsAdjuster),
             getInsertArgumentAdjuster(CUDAVerMinor.c_str(),
                                       ArgumentInsertPosition::BEGIN));
-        CudaArgsAdjuster = combineAdjusters(
-            std::move(CudaArgsAdjuster),
-            getInsertArgumentAdjuster("-fgpu-exclude-wrong-side-overloads",
-                                      ArgumentInsertPosition::BEGIN));
         CudaArgsAdjuster =
             combineAdjusters(std::move(CudaArgsAdjuster),
                              getInsertArgumentAdjuster(

clang/test/dpct/compat_with_clang.cu

@@ -3,16 +3,8 @@
 // RUN: FileCheck %s --match-full-lines --input-file %T/compat_with_clang/compat_with_clang.dp.cpp
 // RUN: %if build_lit %{icpx -c -fsycl %T/compat_with_clang/compat_with_clang.dp.cpp -o %T/compat_with_clang/compat_with_clang.dp.o %}
 
-#include "cuda_fp16.h"
 #include <cstdint>
 
-// CHECK: inline void foo1(sycl::half2 *array, sycl::half a) {
-// CHECK-NEXT:   array[dpct::reverse_bits<unsigned int>(123)] = {a, sycl::vec<float, 1>(2.3f).convert<sycl::half, sycl::rounding_mode::automatic>()[0]};
-// CHECK-NEXT: }
-__device__ inline void foo1(__half2 *array, __half a) {
-  array[__brev(123)] = {a, __float2half(2.3f)};
-}
-
 // CHECK: void foo2(int a, int b) {
 // CHECK-NEXT:   dpct::dim3 block{dpct::min(512, uint32_t(a * b))};
 // CHECK-NEXT: }

clang/test/dpct/compat_with_clang_4.cu

@@ -0,0 +1,14 @@
+// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1
+// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1
+// RUN: dpct --format-range=none -out-root %T/compat_with_clang_4 %s --cuda-include-path="%cuda-path/include" --stop-on-parse-err
+// RUN: FileCheck %s --match-full-lines --input-file %T/compat_with_clang_4/compat_with_clang_4.dp.cpp
+// RUN: %if build_lit %{icpx -c -fsycl %T/compat_with_clang_4/compat_with_clang_4.dp.cpp -o %T/compat_with_clang_4/compat_with_clang_4.dp.o %}
+
+#include "cuda_fp16.h"
+
+// CHECK: inline void foo1(sycl::half2 *array, sycl::half a) {
+// CHECK-NEXT:   array[dpct::reverse_bits<unsigned int>(123)] = {a, sycl::vec<float, 1>(2.3f).convert<sycl::half, sycl::rounding_mode::automatic>()[0]};
+// CHECK-NEXT: }
+__device__ inline void foo1(__half2 *array, __half a) {
+  array[__brev(123)] = {a, __float2half(2.3f)};
+}