Merge branch 'sycl' into multiple-build

Author: ayylol

clang/test/Driver/sycl-device-traits-macros-amdgcn.cpp

@@ -50,6 +50,10 @@
 // RUN:   FileCheck %s --check-prefix=CHECK-SYCL-AMDGCN-AMD-AMDHSA-DEVICE-TRIPLE
 // RUN: %clangxx -fsycl -nogpulib -fsycl-targets=amd_gpu_gfx940 \
 // RUN:   -fsycl-libspirv-path=%S/Inputs/SYCL/libspirv.bc -### %s 2>&1 | \
+// RUN: %clangxx -fsycl -nogpulib -fsycl-targets=amd_gpu_gfx941 \
+// RUN:   -fsycl-libspirv-path=%S/Inputs/SYCL/libspirv.bc -### %s 2>&1 | \
+// RUN: %clangxx -fsycl -nogpulib -fsycl-targets=amd_gpu_gfx942 \
+// RUN:   -fsycl-libspirv-path=%S/Inputs/SYCL/libspirv.bc -### %s 2>&1 | \
 // RUN:   FileCheck %s --check-prefix=CHECK-SYCL-AMDGCN-AMD-AMDHSA-DEVICE-TRIPLE
 // RUN: %clangxx -fsycl -nogpulib -fsycl-targets=amd_gpu_gfx1010 \
 // RUN:   -fsycl-libspirv-path=%S/Inputs/SYCL/libspirv.bc -### %s 2>&1 | \
@@ -156,6 +160,12 @@
 // RUN: %clangxx -fsycl -nogpulib -fsycl-targets=amdgcn-amd-amdhsa -Xsycl-target-backend --offload-arch=gfx940 \
 // RUN:   -fsycl-libspirv-path=%S/Inputs/SYCL/libspirv.bc -### %s 2>&1 | \
 // RUN:   FileCheck %s --check-prefix=CHECK-SYCL-AMDGCN-AMD-AMDHSA-OFFLOAD-ARCH
+// RUN: %clangxx -fsycl -nogpulib -fsycl-targets=amdgcn-amd-amdhsa -Xsycl-target-backend --offload-arch=gfx941 \
+// RUN:   -fsycl-libspirv-path=%S/Inputs/SYCL/libspirv.bc -### %s 2>&1 | \
+// RUN:   FileCheck %s --check-prefix=CHECK-SYCL-AMDGCN-AMD-AMDHSA-OFFLOAD-ARCH
+// RUN: %clangxx -fsycl -nogpulib -fsycl-targets=amdgcn-amd-amdhsa -Xsycl-target-backend --offload-arch=gfx942 \
+// RUN:   -fsycl-libspirv-path=%S/Inputs/SYCL/libspirv.bc -### %s 2>&1 | \
+// RUN:   FileCheck %s --check-prefix=CHECK-SYCL-AMDGCN-AMD-AMDHSA-OFFLOAD-ARCH
 // RUN: %clangxx -fsycl -nogpulib -fsycl-targets=amdgcn-amd-amdhsa -Xsycl-target-backend --offload-arch=gfx1010 \
 // RUN:   -fsycl-libspirv-path=%S/Inputs/SYCL/libspirv.bc -### %s 2>&1 | \
 // RUN:   FileCheck %s --check-prefix=CHECK-SYCL-AMDGCN-AMD-AMDHSA-OFFLOAD-ARCH

libdevice/imf_rounding_op.hpp

@@ -860,7 +860,7 @@ template <typename Ty> Ty __fp_div(Ty x, Ty y, int rd) {
   }
 }
 
-unsigned get_grs_bits(uint64_t dbits, unsigned bit_num) {
+static unsigned get_grs_bits(uint64_t dbits, unsigned bit_num) {
   if (bit_num == 1)
     return (dbits & 0x1) << 2;
   else if (bit_num == 2)
@@ -873,7 +873,7 @@ unsigned get_grs_bits(uint64_t dbits, unsigned bit_num) {
   }
 }
 
-unsigned get_grs_bits(__iml_ui128 dbits, unsigned bit_num) {
+static unsigned get_grs_bits(__iml_ui128 dbits, unsigned bit_num) {
   if (bit_num == 1)
     return static_cast<uint32_t>(dbits & 0x1) << 2;
   else if (bit_num == 2)

llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td

@@ -275,7 +275,11 @@ class HipTargetInfo<string targetName, list<Aspect> aspectList, list<int> subGro
 defvar HipSubgroupSizesGCN2 = [16];     // gfx7
 defvar HipSubgroupSizesGCN3 = [16];     // gfx8, GCN 3rd gen and 4th gen have the same subgroup sizes
 defvar HipSubgroupSizesGCN5 = [64];     // gfx900-gfx906 GCN5.0 (known as "Vega"), gfx90c GCN5.1 (known as "Vega 7nm")
-defvar HipSubgroupSizesRDNA = [32, 64]; // gfxX10-gfx11 (encapsulates RDNA1..3), natively 32 (64-waves mode available)
+// According to the "Accelerator and GPU hardware specifications table" docs,
+// (see: https://rocm.docs.amd.com/en/latest/reference/gpu-arch-specs.html)
+// the ROCm driver selects wave32 mode for the gfx10 and gfx11 family of GPUs.
+// Also, see relevant ROCm issue: https://github.com/ROCm/hipamd/issues/59
+defvar HipSubgroupSizesRDNA = [32];     // gfxX10-gfx11 (encapsulates RDNA1..3), (wave64 mode available but not used).
 defvar HipSubgroupSizesCDNA = [64];     // gfx908, gfx90a (encapsulates CDNA1..2)
 
 defvar HipMinAspects = [AspectGpu, AspectFp64, AspectOnline_compiler, AspectOnline_linker, AspectQueue_profiling,
@@ -290,9 +294,18 @@ def : HipTargetInfo<"amd_gpu_gfx908", !listconcat(HipMinAspects, AllUSMAspects,
 def : HipTargetInfo<"amd_gpu_gfx90a", !listconcat(HipMinAspects, AllUSMAspects,
     [AspectExt_intel_device_info_uuid, AspectExt_oneapi_graph, AspectExt_oneapi_limited_graph, AspectExt_oneapi_native_assert]),
     HipSubgroupSizesCDNA>;
+// TODO: Need to verify whether device-side asserts (oneapi_native_assert) are
+// now working for the new CDNA3 gfx940, gfx941, gfx942 GPUs and fixed for the
+// other supported, gfx1030 and gfx1100, RDNA3 GPUs.
 def : HipTargetInfo<"amd_gpu_gfx940", !listconcat(HipMinAspects, AllUSMAspects,
     [AspectExt_intel_device_info_uuid, AspectExt_oneapi_graph, AspectExt_oneapi_limited_graph]),
     HipSubgroupSizesCDNA>;
+def : HipTargetInfo<"amd_gpu_gfx941", !listconcat(HipMinAspects, AllUSMAspects,
+    [AspectExt_intel_device_info_uuid, AspectExt_oneapi_graph, AspectExt_oneapi_limited_graph]),
+    HipSubgroupSizesCDNA>;
+def : HipTargetInfo<"amd_gpu_gfx942", !listconcat(HipMinAspects, AllUSMAspects,
+    [AspectExt_intel_device_info_uuid, AspectExt_oneapi_graph, AspectExt_oneapi_limited_graph]),
+    HipSubgroupSizesCDNA>;
 def : HipTargetInfo<"amd_gpu_gfx1030", !listconcat(HipMinAspects, AllUSMAspects,
     [AspectExt_intel_device_info_uuid, AspectExt_oneapi_graph, AspectExt_oneapi_limited_graph]),
     HipSubgroupSizesRDNA>;
@@ -332,7 +345,5 @@ def : HipTargetInfo<"amd_gpu_gfx1103", !listconcat(HipMinAspects, AllUSMAspects)
 def : HipTargetInfo<"amd_gpu_gfx1150", !listconcat(HipMinAspects, AllUSMAspects), HipSubgroupSizesRDNA>;
 def : HipTargetInfo<"amd_gpu_gfx1151", !listconcat(HipMinAspects, AllUSMAspects), HipSubgroupSizesRDNA>;
 // TBA
-def : HipTargetInfo<"amd_gpu_gfx941", [], []>;  // CDNA 3
-def : HipTargetInfo<"amd_gpu_gfx942", [], []>;  // CDNA 3
 def : HipTargetInfo<"amd_gpu_gfx1200", [], []>; // RDNA 4
 def : HipTargetInfo<"amd_gpu_gfx1201", [], []>; // RDNA 4

sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc renamed to sycl/doc/extensions/experimental/sycl_ext_intel_cache_controls.asciidoc

@@ -36,7 +36,7 @@ https://github.com/intel/llvm/issues
 
 == Dependencies
 
-This extension is written against the SYCL 2020 revision 7 specification.  All
+This extension is written against the SYCL 2020 revision 9 specification.  All
 references below to the "core SYCL specification" or to section numbers in the
 SYCL specification refer to that revision.
 
@@ -49,11 +49,12 @@ This extension depends on the following SYCL extensions:
 
 == Status
 
-This is a proposed extension specification, intended to gather community
-feedback.  Interfaces defined in this specification may not be implemented yet
-or may be in a preliminary state.  The specification itself may also change in
-incompatible ways before it is finalized.  *Shipping software products should
-not rely on APIs defined in this specification.*
+This is an experimental extension specification, intended to provide early
+access to features and gather community feedback.  Interfaces defined in this
+specification are implemented in {dpcpp}, but they are not finalized and may
+change incompatibly in future versions of {dpcpp} without prior notice.
+*Shipping software products should not rely on APIs defined in this
+specification.*
 
 
 == Overview

sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_kernels.asciidoc

@@ -705,12 +705,12 @@ static constexpr size_t WGSIZE = 16;
 SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
 void iota(float start, float *ptr) {
   // Get the ID of this kernel iteration.
-  size_t id = syclext::this_work_item::get_nd_item().get_global_linear_id();
+  size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
 
   ptr[id] = start + static_cast<float>(id);
 }
 
-void main() {
+int main() {
   sycl::queue q;
   sycl::context ctxt = q.get_context();
 
@@ -963,6 +963,13 @@ int main() {
 ```
 +
 Where `kfp` would have some nicer name.
++
+With either form above, it seems like we have enough type information for the
+header to check that the types of the actual kernel arguments are implicitly
+convertible to the types of the formal kernel parameters, and we can raise a
+compile-time error if they are not.
+In addition, the header can perform any necessary implicit conversions when
+setting the kernel argument values.
 
 * We are debating whether we should allow a free function kernel to be defined
   with an initial "iteration index" parameter such as:

sycl/doc/extensions/proposed/sycl_ext_oneapi_work_group_memory.asciidoc

@@ -457,7 +457,7 @@ constexpr size_t WGSIZE = 256;
 
 SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
 void mykernel(syclexp::work_group_memory<int[WGSIZE]> mem) {
-  size_t id = syclext::this_work_item::get_nd_item().get_local_linear_id();
+  size_t id = syclext::this_work_item::get_nd_item<1>().get_local_linear_id();
 
   // Each work-item has its own dedicated element of the device local memory
   // array.

sycl/test-e2e/AddressSanitizer/lit.local.cfg

@@ -12,3 +12,7 @@ config.unsupported_features += ['cuda', 'hip']
 
 # FIXME: Skip some of gpu devices, waiting for gfx driver uplifting
 config.unsupported_features += ['gpu-intel-gen9', 'gpu-intel-gen11', 'gpu-intel-gen12', 'gpu-intel-pvc']
+
+# GPU testing requires level_zero
+if 'opencl:gpu' in config.sycl_devices:
+   config.required_features += ['level_zero']