Merge remote-tracking branch 'remote/sycl' into stream-class

Author: zahiraam

clang/lib/Driver/ToolChains/Clang.cpp

@@ -4208,6 +4208,17 @@ static bool ContainsWrapperAction(const Action *A) {
   return false;
 }
 
+/// Check whether the given input tree contains any append footer actions
+static bool ContainsAppendFooterAction(const Action *A) {
+  if (isa<AppendFooterJobAction>(A))
+    return true;
+  for (const auto &AI : A->inputs())
+    if (ContainsAppendFooterAction(AI))
+      return true;
+
+  return false;
+}
+
 // Put together an external compiler compilation call which is used instead
 // of the clang invocation for the host compile of an offload compilation.
 // Enabling command line:  clang++ -fsycl -fsycl-host-compiler=<HostExe>
@@ -4641,8 +4652,11 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
     if (!IsSYCLOffloadDevice) {
       // Add the -include option to add the integration header
       StringRef Header = D.getIntegrationHeader(Input.getBaseInput());
+      // Do not add the integration header if we are compiling after the
+      // integration footer has been applied.  Check for the append job
+      // action to determine this.
       if (types::getPreprocessedType(Input.getType()) != types::TY_INVALID &&
-          !Header.empty()) {
+          !Header.empty() && !ContainsAppendFooterAction(&JA)) {
         CmdArgs.push_back("-include");
         CmdArgs.push_back(Args.MakeArgString(Header));
         // When creating dependency information, filter out the generated

clang/test/Driver/sycl-int-footer.cpp

@@ -0,0 +1,8 @@
+/// Check compilation tool steps when using the integrated footer
+// RUN:  %clangxx -fsycl -fsycl-use-footer %s -### 2>&1 \
+// RUN:   | FileCheck -check-prefix FOOTER %s
+// FOOTER: clang{{.*}} "-fsycl-is-device"{{.*}} "-fsycl-int-header=[[INTHEADER:.+\.h]]" "-fsycl-int-footer=[[INTFOOTER:.+\h]]" "-sycl-std={{.*}}"
+// FOOTER: clang{{.*}} "-include" "[[INTHEADER]]"{{.*}} "-fsycl-is-host"{{.*}} "-E"{{.*}} "-o" "[[PREPROC:.+\.ii]]"
+// FOOTER: append-file{{.*}} "[[PREPROC]]" "--append=[[INTFOOTER]]" "--output=[[APPENDEDSRC:.+\.cpp]]"
+// FOOTER: clang{{.*}} "-fsycl-is-host"{{.*}} "[[APPENDEDSRC]]"
+// FOOTER-NOT: "-include" "[[INTHEADER]]"

libdevice/device_itt.h

@@ -15,6 +15,7 @@
 #include "spirv_vars.h"
 
 #define ITT_STUB_ATTRIBUTES __attribute__((noinline, optnone))
+#define ITT_WRAPPER_ATTRIBUTES __attribute__((always_inline))
 
 /// Atomic operation type
 enum __itt_atomic_mem_op_t {
@@ -36,7 +37,7 @@ DEVICE_EXTERN_C char __spirv_SpecConstant(int, char);
 
 #define ITT_SPEC_CONSTANT 0xFF747469
 
-static inline bool isITTEnabled() {
+static ITT_WRAPPER_ATTRIBUTES bool isITTEnabled() {
   return __spirv_SpecConstant(ITT_SPEC_CONSTANT, 0) != 0;
 }
 
@@ -50,14 +51,15 @@ static inline bool isITTEnabled() {
 // be computed in the wrapper itself and has to be passed from outside.
 // If a compiler needs to invoke such an API, it has to use the user
 // visible API directly (i.e. __itt_offload_sync_acquired).
-DEVICE_EXTERN_C
-void __itt_offload_wi_start_wrapper();
-DEVICE_EXTERN_C
-void __itt_offload_wi_finish_wrapper();
-DEVICE_EXTERN_C
-void __itt_offload_wg_barrier_wrapper();
-DEVICE_EXTERN_C
-void __itt_offload_wi_resume_wrapper();
+//
+// FIXME: we need to add always_inline compiler wrappers
+//        for atomic_op_start/finish. Compiler calls user
+//        wrappers right now, and they may interfere with
+//        debugging user code in non-ITT mode.
+DEVICE_EXTERN_C ITT_WRAPPER_ATTRIBUTES void __itt_offload_wi_start_wrapper();
+DEVICE_EXTERN_C ITT_WRAPPER_ATTRIBUTES void __itt_offload_wi_finish_wrapper();
+DEVICE_EXTERN_C ITT_WRAPPER_ATTRIBUTES void __itt_offload_wg_barrier_wrapper();
+DEVICE_EXTERN_C ITT_WRAPPER_ATTRIBUTES void __itt_offload_wi_resume_wrapper();
 
 // Non-inlinable and non-optimizable APIs that are recognized
 // by profiling tools.

sycl/plugins/cuda/CMakeLists.txt

@@ -43,7 +43,3 @@ install(TARGETS pi_cuda
   LIBRARY DESTINATION "lib${LLVM_LIBDIR_SUFFIX}" COMPONENT pi_cuda
   RUNTIME DESTINATION "bin" COMPONENT pi_cuda
 )
-
-# `sycl/source/CMakeLists.txt` adapted when SYCL_BUILD_PI_CUDA is defined:
-# target_link_libraries(sycl PUBLIC pi_cuda)
-# target_compile_definitions(sycl PUBLIC USE_PI_CUDA)

sycl/source/CMakeLists.txt

@@ -82,10 +82,6 @@ function(add_sycl_rt_library LIB_NAME)
       $<$<BOOL:${SYCL_BUILD_PI_CUDA}>:pi_cuda>
   )
 
-  target_compile_definitions(${LIB_OBJ_NAME}
-    PUBLIC
-      $<$<BOOL:${SYCL_BUILD_PI_CUDA}>:USE_PI_CUDA>)
-
   add_common_options(${LIB_NAME} ${LIB_OBJ_NAME})
 
   set_target_properties(${LIB_NAME} PROPERTIES

sycl/source/detail/context_impl.cpp

@@ -45,7 +45,6 @@ context_impl::context_impl(const vector_class<cl::sycl::device> Devices,
 
   const auto Backend = getPlugin().getBackend();
   if (Backend == backend::cuda) {
-#if USE_PI_CUDA
     const bool UseCUDAPrimaryContext =
         MPropList.has_property<property::context::cuda::use_primary_context>();
     const pi_context_properties Props[] = {
@@ -55,10 +54,6 @@ context_impl::context_impl(const vector_class<cl::sycl::device> Devices,
 
     getPlugin().call<PiApiKind::piContextCreate>(
         Props, DeviceIds.size(), DeviceIds.data(), nullptr, nullptr, &MContext);
-#else
-    cl::sycl::detail::pi::die(
-        "CUDA support was not enabled at compilation time");
-#endif
   } else {
     getPlugin().call<PiApiKind::piContextCreate>(nullptr, DeviceIds.size(),
                                                  DeviceIds.data(), nullptr,