Merge branch 'sycl' into review/yang/fix_kernel_filter

Author: AllanZyne

clang/lib/Driver/ToolChains/HIPUtility.cpp

@@ -324,7 +324,10 @@ void HIP::constructHIPFatbinCommand(Compilation &C, const JobAction &JA,
       Args.MakeArgString(std::string("-output=").append(Output));
   BundlerArgs.push_back(BundlerOutputArg);
 
-  addOffloadCompressArgs(Args, BundlerArgs);
+  // For SYCL, the compression is occurring during the wrapping step, so we do
+  // not want to do additional compression here.
+  if (!JA.isDeviceOffloading(Action::OFK_SYCL))
+    addOffloadCompressArgs(Args, BundlerArgs);
 
   const char *Bundler = Args.MakeArgString(
       T.getToolChain().GetProgramPath("clang-offload-bundler"));

clang/test/Driver/sycl-offload-wrapper-compression.cpp

@@ -10,5 +10,11 @@
 // RUN: %clangxx -### -fsycl %s 2>&1 | FileCheck %s --check-prefix=CHECK-NO-COMPRESS
 // RUN: %clangxx -### -fsycl --offload-compression-level=3 %s 2>&1 | FileCheck %s --check-prefix=CHECK-NO-COMPRESS
 
+// For SYCL offloading to HIP, make sure we don't pass '--compress' to offload-bundler.
+// RUN: %clangxx -### -fsycl --offload-compress -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx1031 -fsycl-targets=amdgcn-amd-amdhsa -fno-sycl-libspirv -nogpulib %s &> %t.driver
+// RUN: FileCheck %s --check-prefix=CHECK-NO-COMPRESS-BUNDLER --input-file=%t.driver
+
 // CHECK-NO-COMPRESS-NOT: {{.*}}clang-offload-wrapper{{.*}}"-offload-compress"{{.*}}
 // CHECK-NO-COMPRESS-NOT: {{.*}}clang-offload-wrapper{{.*}}"-offload-compression-level=3"{{.*}}
+
+// CHECK-NO-COMPRESS-BUNDLER-NOT: {{.*}}clang-offload-bundler{{.*}}"-compress"{{.*}}

libdevice/cmake/modules/SYCLLibdevice.cmake

@@ -112,13 +112,29 @@ function(compile_lib filename)
     "FILETYPE"
     "SRC;EXTRA_OPTS;DEPENDENCIES"
     ${ARGN})
+    set(compile_opt_list ${compile_opts}
+                         ${${ARG_FILETYPE}_device_compile_opts}
+                         ${ARG_EXTRA_OPTS})
+    compile_lib_ext(${filename}
+      FILETYPE ${ARG_FILETYPE}
+      SRC ${ARG_SRC}
+      DEPENDENCIES ${ARG_DEPENDENCIES}
+      OPTS ${compile_opt_list})
+endfunction()
+
+function(compile_lib_ext filename)
+  cmake_parse_arguments(ARG
+    ""
+    "FILETYPE"
+    "SRC;OPTS;DEPENDENCIES"
+    ${ARGN})
 
   set(devicelib-file
     ${${ARG_FILETYPE}_binary_dir}/${filename}.${${ARG_FILETYPE}-suffix})
 
   add_custom_command(
     OUTPUT ${devicelib-file}
-    COMMAND ${clang} ${compile_opts} ${ARG_EXTRA_OPTS}
+    COMMAND ${clang} ${ARG_OPTS}
             ${CMAKE_CURRENT_SOURCE_DIR}/${ARG_SRC} -o ${devicelib-file}
     MAIN_DEPENDENCY ${ARG_SRC}
     DEPENDS ${ARG_DEPENDENCIES}

sycl/include/sycl/ext/oneapi/experimental/prefetch.hpp

@@ -20,7 +20,8 @@ enum class cache_level { L1 = 0, L2 = 1, L3 = 2, L4 = 3 };
 
 struct nontemporal;
 
-struct prefetch_hint_key {
+struct prefetch_hint_key
+    : detail::compile_time_property_key<detail::PropKind::Prefetch> {
   template <cache_level Level, typename Hint>
   using value_t =
       property_value<prefetch_hint_key,
@@ -51,8 +52,6 @@ inline constexpr prefetch_hint_key::value_t<cache_level::L4, nontemporal>
 namespace detail {
 using namespace sycl::detail;
 
-template <> struct IsCompileTimeProperty<prefetch_hint_key> : std::true_type {};
-
 template <cache_level Level, typename Hint>
 struct PropertyMetaInfo<prefetch_hint_key::value_t<Level, Hint>> {
   static constexpr const char *name = std::is_same_v<Hint, nontemporal>

sycl/include/sycl/ext/oneapi/properties/property.hpp

@@ -214,8 +214,9 @@ enum PropKind : uint32_t {
   ResponseCapacity = 73,
   MaxWorkGroupSize = 74,
   MaxLinearWorkGroupSize = 75,
+  Prefetch = 76,
   // PropKindSize must always be the last value.
-  PropKindSize = 76,
+  PropKindSize = 77,
 };
 
 struct property_key_base_tag {};

sycl/test-e2e/Compression/compression.cpp

@@ -1,9 +1,6 @@
 // End-to-End test for testing device image compression.
 // REQUIRES: zstd
 
-// XFAIL: hip_amd
-// XFAIL-TRACKER: https://github.com/intel/llvm/issues/15829
-
 // RUN: %{build} -O0 -g %S/Inputs/single_kernel.cpp -o %t_not_compress.out
 // RUN: %{build} -O0 -g --offload-compress --offload-compression-level=3 %S/Inputs/single_kernel.cpp -o %t_compress.out
 // RUN: %{run} %t_not_compress.out

sycl/test-e2e/Compression/compression_multiple_tu.cpp

@@ -2,9 +2,6 @@
 // translation units, one compressed and one not compressed.
 // REQUIRES: zstd, linux
 
-// XFAIL: hip_amd
-// XFAIL-TRACKER: https://github.com/intel/llvm/issues/15829
-
 // RUN: %{build} --offload-compress -DENABLE_KERNEL1 -shared -fPIC -o %T/kernel1.so
 // RUN: %{build} -DENABLE_KERNEL2 -shared -fPIC -o %T/kernel2.so