Merge remote-tracking branch 'upstream/sycl' into konrad/AddSyncToCBRelease

Author: Konrad Kusiak

clang/include/clang/Basic/Attr.td

@@ -1637,12 +1637,14 @@ def SYCLType: InheritableAttr {
   let Subjects = SubjectList<[CXXRecord, Enum], ErrorDiag>;
   let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
   let Args = [EnumArgument<"Type", "SYCLType", /*is_string=*/true,
-                           ["accessor", "local_accessor", "work_group_memory", "dynamic_work_group_memory",
+                           ["accessor", "local_accessor", "dynamic_local_accessor",
+                           "work_group_memory", "dynamic_work_group_memory",
                             "specialization_id", "kernel_handler", "buffer_location",
                             "no_alias", "accessor_property_list", "group",
                             "private_memory", "aspect", "annotated_ptr", "annotated_arg",
                             "stream", "sampler", "host_pipe", "multi_ptr"],
-                           ["accessor", "local_accessor", "work_group_memory", "dynamic_work_group_memory",
+                           ["accessor", "local_accessor", "dynamic_local_accessor",
+                           "work_group_memory", "dynamic_work_group_memory",
                             "specialization_id", "kernel_handler", "buffer_location",
                             "no_alias", "accessor_property_list", "group",
                             "private_memory", "aspect", "annotated_ptr", "annotated_arg",

clang/include/clang/Sema/SemaSYCL.h

@@ -64,7 +64,8 @@ class SYCLIntegrationHeader {
     kind_stream,
     kind_work_group_memory,
     kind_dynamic_work_group_memory,
-    kind_last = kind_dynamic_work_group_memory
+    kind_dynamic_accessor,
+    kind_last = kind_dynamic_accessor
   };
 
 public:

clang/lib/Sema/SemaSYCL.cpp

@@ -94,7 +94,8 @@ bool SemaSYCL::isSyclType(QualType Ty, SYCLTypeAttr::SYCLType TypeName) {
 
 static bool isSyclAccessorType(QualType Ty) {
   return SemaSYCL::isSyclType(Ty, SYCLTypeAttr::accessor) ||
-         SemaSYCL::isSyclType(Ty, SYCLTypeAttr::local_accessor);
+         SemaSYCL::isSyclType(Ty, SYCLTypeAttr::local_accessor) ||
+         SemaSYCL::isSyclType(Ty, SYCLTypeAttr::dynamic_local_accessor);
 }
 
 // FIXME: Accessor property lists should be modified to use compile-time
@@ -1151,7 +1152,8 @@ static QualType GetSYCLKernelObjectType(const FunctionDecl *KernelCaller) {
 /// \return the target of given SYCL accessor type
 static target getAccessTarget(QualType FieldTy,
                               const ClassTemplateSpecializationDecl *AccTy) {
-  if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::local_accessor))
+  if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::local_accessor) ||
+      SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::dynamic_local_accessor))
     return local;
 
   return static_cast<target>(
@@ -4815,7 +4817,13 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
       int Dims = static_cast<int>(
           AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue());
       int Info = getAccessTarget(FieldTy, AccTy) | (Dims << 11);
-      Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info,
+
+      SYCLIntegrationHeader::kernel_param_kind_t ParamKind =
+          SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::dynamic_local_accessor)
+              ? SYCLIntegrationHeader::kind_dynamic_accessor
+              : SYCLIntegrationHeader::kind_accessor;
+
+      Header.addParamDesc(ParamKind, Info,
                           CurOffset +
                               offsetOf(RD, BC.getType()->getAsCXXRecordDecl()));
     } else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::work_group_memory)) {
@@ -4841,8 +4849,12 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
           AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue());
       int Info = getAccessTarget(FieldTy, AccTy) | (Dims << 11);
 
-      Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info,
-                          CurOffset + offsetOf(FD, FieldTy));
+      SYCLIntegrationHeader::kernel_param_kind_t ParamKind =
+          SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::dynamic_local_accessor)
+              ? SYCLIntegrationHeader::kind_dynamic_accessor
+              : SYCLIntegrationHeader::kind_accessor;
+
+      Header.addParamDesc(ParamKind, Info, CurOffset + offsetOf(FD, FieldTy));
     } else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::stream)) {
       addParam(FD, FieldTy, SYCLIntegrationHeader::kind_stream);
     } else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::work_group_memory)) {
@@ -6056,6 +6068,7 @@ static const char *paramKind2Str(KernelParamKind K) {
     CASE(pointer);
     CASE(work_group_memory);
     CASE(dynamic_work_group_memory);
+    CASE(dynamic_accessor);
   }
   return "<ERROR>";
 

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -440,6 +440,9 @@ local_accessor: public accessor<dataT,
 #ifdef __SYCL_DEVICE_ONLY__
   void __init(__attribute__((opencl_local)) dataT *Ptr, range<dimensions> AccessRange,
               range<dimensions> MemRange, id<dimensions> Offset) {}
+
+template <typename, int>
+  friend class dynamic_local_accessor;
 #endif
 };
 
@@ -693,6 +696,23 @@ __SYCL_TYPE(dynamic_work_group_memory) dynamic_work_group_memory {
   work_group_memory<DataT> LocalMem;
 };
 
+template <typename DataT, int Dimensions>
+class __attribute__((sycl_special_class))
+__SYCL_TYPE(dynamic_local_accessor) dynamic_local_accessor {
+public:
+  dynamic_local_accessor() = default;
+
+void __init(__attribute__((opencl_local)) DataT *Ptr,
+            range<Dimensions> AccessRange, range<Dimensions> range,
+            id<Dimensions> id) {
+  this->LocalMem.__init(Ptr, AccessRange, range, id);
+}
+  local_accessor<DataT, Dimensions> get() const { return LocalMem; }
+
+private:
+  local_accessor<DataT, Dimensions> LocalMem;
+};
+
 template <typename T, int dimensions = 1,
           typename AllocatorT = int /*fake type as AllocatorT is not used*/>
 class buffer {

clang/test/CodeGenSYCL/dynamic_local_accessor.cpp

@@ -0,0 +1,40 @@
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o %t.ll
+// RUN: FileCheck < %t.ll %s --check-prefix CHECK-IR
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -fsycl-int-header=%t.h %s
+// RUN: FileCheck < %t.h %s --check-prefix CHECK-INT-HEADER
+//
+// Tests for dynamic_local_accessor kernel parameter using the dummy implementation in Inputs/sycl.hpp.
+// The first two RUN commands verify that the init call is generated with the correct arguments in LLVM IR
+// and the second two RUN commands verify the contents of the integration header produced by the frontend.
+//
+// CHECK-IR: define dso_local spir_kernel void @
+// CHECK-IR-SAME: ptr addrspace(3) noundef align 4 [[PTR:%[a-zA-Z0-9_]+]]
+//
+// CHECK-IR: [[PTR]].addr = alloca ptr addrspace(3), align 8
+// CHECK-IR: [[PTR]].addr.ascast = addrspacecast ptr [[PTR]].addr to ptr addrspace(4)
+// CHECK-IR: store ptr addrspace(3) [[PTR]], ptr addrspace(4) [[PTR]].addr.ascast, align 8
+// CHECK-IR: [[PTR_LOAD:%[a-zA-Z0-9_]+]] = load ptr addrspace(3), ptr addrspace(4) [[PTR]].addr.ascast, align 8
+//
+// CHECK-IR: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) noundef align 4 dereferenceable_or_null(24) %{{[a-zA-Z0-9_]+}}, ptr addrspace(3) noundef %{{[a-zA-Z0-9_]+}}, ptr noundef byval(%"struct.sycl::_V1::range") align 4 %agg.{{[a-zA-Z0-9_]+}}.ascast.ascast, ptr noundef byval(%"struct.sycl::_V1::range") align 4 %agg.{{[a-zA-Z0-9_]+}}.ascast.ascast, ptr noundef byval(%"struct.sycl::_V1::id") align 4 %agg.{{[a-zA-Z0-9_]+}}.ascast.ascast) #{{[0-9_]+}}
+//
+// CHECK-INT-HEADER: const kernel_param_desc_t kernel_signatures[] = {
+// CHECK-INT-HEADER-NEXT: //--- _ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlNS0_4itemILi1EEEE_
+// CHECK-INT-HEADER-NEXT: { kernel_param_kind_t::kind_dynamic_accessor, 4064, 0 },
+// CHECK-INT-HEADER-EMPTY:
+// CHECK-INT-HEADER-NEXT: { kernel_param_kind_t::kind_invalid, -987654321, -987654321 },
+// CHECK-INT-HEADER-NEXT: };
+
+#include "Inputs/sycl.hpp"
+
+int main() {
+  sycl::queue Q;
+  sycl::dynamic_local_accessor<int, 1> DynLocalAcc;
+  Q.submit([&](sycl::handler &CGH) {
+    sycl::range<1> ndr;
+    CGH.parallel_for(ndr, [=](sycl::item<1> it) { 
+      auto LocalAcc = DynLocalAcc.get();
+      auto* Ptr = &LocalAcc;
+    });
+  });
+  return 0;
+}

devops/scripts/benchmarks/html/scripts.js

@@ -681,6 +681,18 @@ function getLayerTags(metadata) {
 
 function processLayerComparisonsData(benchmarkRuns) {
     const groupedResults = {};
+    const labelsByGroup = {};
+
+    benchmarkRuns.forEach(run => {
+        run.results.forEach(result => {
+            if (result.explicit_group) {
+                if (!labelsByGroup[result.explicit_group]) {
+                    labelsByGroup[result.explicit_group] = new Set();
+                }
+                labelsByGroup[result.explicit_group].add(result.label);
+            }
+        });
+    });
 
     benchmarkRuns.forEach(run => {
         run.results.forEach(result => {
@@ -691,13 +703,7 @@ function processLayerComparisonsData(benchmarkRuns) {
             if (!metadata) return;
 
             // Get all benchmark labels in this group
-            const labelsInGroup = new Set(
-                benchmarkRuns.flatMap(r =>
-                    r.results
-                        .filter(res => res.explicit_group === result.explicit_group)
-                        .map(res => res.label)
-                )
-            );
+            const labelsInGroup = labelsByGroup[result.explicit_group];
 
             // Check if this group compares different layers
             const uniqueLayers = new Set();

libclc/CMakeLists.txt

@@ -413,23 +413,27 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
     DIRS ${opencl_dirs}
   )
 
-  set( libspirv_lib_files )
-  set( libspirv_gen_files )
-
-  if( NOT ARCH STREQUAL spirv AND NOT ARCH STREQUAL spirv64 )
-    if( ARCH STREQUAL clspv OR ARCH STREQUAL clspv64 )
-      list( APPEND libspirv_gen_files clspv-convert.cl )
-    elseif ( NOT ENABLE_RUNTIME_SUBNORMAL )
-      list( APPEND libspirv_gen_files spirv-convert.cl )
-      list( APPEND libspirv_lib_files libspirv/lib/generic/subnormal_use_default.ll )
+  set( BUILD_LIBSPIRV_${t} FALSE )
+  if ( t STREQUAL amdgcn--amdhsa OR t STREQUAL nvptx64--nvidiacl OR t STREQUAL native_cpu )
+    set( libspirv_lib_files )
+    set( libspirv_gen_files )
+    set( BUILD_LIBSPIRV_${t} TRUE )
+
+    if( NOT ARCH STREQUAL spirv AND NOT ARCH STREQUAL spirv64 )
+      if( ARCH STREQUAL clspv OR ARCH STREQUAL clspv64 )
+        list( APPEND libspirv_gen_files clspv-convert.cl )
+      elseif ( NOT ENABLE_RUNTIME_SUBNORMAL )
+        list( APPEND libspirv_gen_files spirv-convert.cl )
+        list( APPEND libspirv_lib_files libspirv/lib/generic/subnormal_use_default.ll )
+      endif()
     endif()
-  endif()
 
-  libclc_configure_lib_source(
-    libspirv_lib_files
-    LIB_ROOT_DIR libspirv
-    DIRS ${libspirv_dirs} ${DARCH} ${DARCH}-${OS} ${DARCH}-${VENDOR}-${OS}
-  )
+    libclc_configure_lib_source(
+      libspirv_lib_files
+      LIB_ROOT_DIR libspirv
+      DIRS ${libspirv_dirs} ${DARCH} ${DARCH}-${OS} ${DARCH}-${VENDOR}-${OS}
+    )
+  endif()
 
   foreach( d ${${t}_devices} )
     get_libclc_device_info(
@@ -568,28 +572,30 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
       GEN_FILES ${clc_gen_files}
     )
 
-    set( spirv_build_flags ${build_flags} )
-    list( APPEND spirv_build_flags
-      -I${CMAKE_CURRENT_SOURCE_DIR}/generic/include
-      -I${CMAKE_CURRENT_SOURCE_DIR}/libspirv/include/
-      # FIXME: Fix libspirv to not require disabling this noisy warning
-      -Wno-bitwise-conditional-parentheses
-    )
-
-    add_libclc_builtin_set(
-      ARCH ${ARCH}
-      ARCH_SUFFIX libspirv-${arch_suffix}
-      TRIPLE ${clang_triple}
-      TARGET_ENV libspirv-
-      COMPILE_FLAGS ${spirv_build_flags}
-      OPT_FLAGS ${opt_flags}
-      LIB_FILES ${libspirv_lib_files}
-      GEN_FILES ${libspirv_gen_files}
-      ALIASES ${${d}_aliases}
-      PARENT_TARGET libspirv-builtins
-      # Link in the CLC builtins and internalize their symbols
-      INTERNAL_LINK_DEPENDENCIES builtins.link.clc-${arch_suffix}
-    )
+    if( BUILD_LIBSPIRV_${t} )
+      set( spirv_build_flags ${build_flags} )
+      list( APPEND spirv_build_flags
+        -I${CMAKE_CURRENT_SOURCE_DIR}/generic/include
+        -I${CMAKE_CURRENT_SOURCE_DIR}/libspirv/include/
+        # FIXME: Fix libspirv to not require disabling this noisy warning
+        -Wno-bitwise-conditional-parentheses
+      )
+
+      add_libclc_builtin_set(
+        ARCH ${ARCH}
+        ARCH_SUFFIX libspirv-${arch_suffix}
+        TRIPLE ${clang_triple}
+        TARGET_ENV libspirv-
+        COMPILE_FLAGS ${spirv_build_flags}
+        OPT_FLAGS ${opt_flags}
+        LIB_FILES ${libspirv_lib_files}
+        GEN_FILES ${libspirv_gen_files}
+        ALIASES ${${d}_aliases}
+        PARENT_TARGET libspirv-builtins
+        # Link in the CLC builtins and internalize their symbols
+        INTERNAL_LINK_DEPENDENCIES builtins.link.clc-${arch_suffix}
+      )
+    endif()
 
     set( opencl_build_flags ${build_flags} )
     list( APPEND opencl_build_flags

libclc/libspirv/include/libspirv/image/image.h

@@ -6,6 +6,8 @@
 //
 //===----------------------------------------------------------------------===//
 
+#if defined(__opencl_c_images)
+
 _CLC_OVERLOAD _CLC_DECL float __spirv_ImageRead__Rfloat(image2d_t image,
                                                         int2 coord);
 _CLC_OVERLOAD _CLC_DECL float __spirv_ImageRead__Rfloat(image2d_t image,
@@ -113,3 +115,5 @@ _CLC_OVERLOAD _CLC_DECL void __spirv_ImageWrite(image2d_t image, int4 coord,
 _CLC_OVERLOAD _CLC_DECL void __spirv_ImageWrite(image3d_t image, int4 coord,
                                                 half4 texel);
 #endif
+
+#endif

libclc/libspirv/include/libspirv/spirv_types.h

@@ -50,12 +50,16 @@ typedef struct {
   float real, imag;
 } complex_float;
 
+#ifdef cl_khr_fp64
 typedef struct {
   double real, imag;
 } complex_double;
+#endif
 
+#ifdef cl_khr_fp16
 typedef struct {
   half real, imag;
 } complex_half;
+#endif
 
 #endif // CLC_SPIRV_TYPES