[SYCL] Device code generation for "free functions", a new way to define kernels (#13207)

rdeodhar · web-flow · commit 0d5addcc80fd · 2024-05-16T15:00:52.000-07:00
This PR is intended for review of the device-side code generation of "free functions". The free function spec is here: https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_kernels.asciidoc This PR does not implement the full specification. It does however allow free function markup and execution using a manual method. Later PRs will add the remaining parts of the spec. This is the current status: - Free functions are supported at file scope only. - The SYCL_EXTERNAL markup is needed for free functions in addition to the SYCL_EXT_ONEAPI_FUNCTION_PROPERTY property. - The traits described in the specification section "New traits for kernel functions" are not yet implemented. - The functions described in the specification section "New kernel bundle member functions" are not yet implemented. As a result, a free-function kernel can only be launched by finding the kernel_id that matches the function's name, and this requires knowing how the compiler mangles the function's name when creating the kernel. This is a temporary solution until the "New kernel bundle member functions" are implemented. - The compiler does not yet diagnose an error if the application violates any of the restrictions listed in the specification under the section "Defining a free function kernel". - Device code generation is supported for scalars and USM pointers only. It is not supported for complex kernel argument types requiring decomposition like accessor, local_accessor, or stream. - The implementation has not been tested to handle the case when a kernel argument is optimized away. The switch -fno-sycl-dead-args-optimization could be used to disable this optimization, if needed - The kernel information descriptor info::kernel::num_args cannot yet be used to query the number of arguments in a free function kernel.
diff --git a/clang/include/clang/Sema/SemaSYCL.h b/clang/include/clang/Sema/SemaSYCL.h
@@ -318,6 +318,7 @@ class SemaSYCL : public SemaBase {
   void ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, MangleContext &MC);
   void SetSYCLKernelNames();
   void MarkDevices();
+  void ProcessFreeFunction(FunctionDecl *FD);
 
   /// Get the number of fields or captures within the parsed type.
   ExprResult ActOnSYCLBuiltinNumFieldsExpr(ParsedType PT);
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
@@ -16651,6 +16651,13 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body,
   if (FD && !FD->isDeleted())
     checkTypeSupport(FD->getType(), FD->getLocation(), FD);
 
+  // Handle free functions.
+  if (LangOpts.SYCLIsDevice && FD->hasAttr<SYCLDeviceAttr>() && Body &&
+      (FD->getTemplatedKind() == FunctionDecl::TK_NonTemplate ||
+       FD->getTemplatedKind() ==
+           FunctionDecl::TK_FunctionTemplateSpecialization))
+    SYCL().ProcessFreeFunction(FD);
+
   return dcl;
 }
 
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
diff --git a/clang/test/CodeGenSYCL/free_function_int_header.cpp b/clang/test/CodeGenSYCL/free_function_int_header.cpp
@@ -0,0 +1,85 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -sycl-std=2020 -fsycl-int-header=%t.h %s
+// RUN: FileCheck -input-file=%t.h %s
+// 
+// This test checks integration header contents for free functions with scalar
+// and pointer parameters.
+
+#include "mock_properties.hpp"
+#include "sycl.hpp"
+
+// First overload of function ff_2.
+__attribute__((sycl_device))
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel",
+                                              2)]] void
+ff_2(int *ptr, int start, int end) {
+  for (int i = start; i <= end; i++)
+    ptr[i] = start + 66;
+}
+
+// Second overload of function ff_2.
+__attribute__((sycl_device))
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel",
+  2)]] void
+  ff_2(int* ptr, int start, int end, int value) {
+  for (int i = start; i <= end; i++)
+    ptr[i] = start + value;
+}
+
+// Templated definition of function ff_3.
+template <typename T>
+__attribute__((sycl_device))
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]] void
+ff_3(T *ptr, T start, T end) {
+  for (int i = start; i <= end; i++)
+    ptr[i] = start;
+}
+
+// Explicit instantiation of ff_3 with int type.
+template void ff_3(int *ptr, int start, int end);
+
+// Explicit instantiation of ff_3 with float type.
+template void ff_3(float* ptr, float start, float end);
+
+// Specialization of ff_3 with double type.
+template <> void ff_3<double>(double *ptr, double start, double end) {
+  for (int i = start; i <= end; i++)
+    ptr[i] = end;
+}
+
+// CHECK:      const char* const kernel_names[] = {
+// CHECK-NEXT:   {{.*}}__sycl_kernel_ff_2Piii
+// CHECK-NEXT:   {{.*}}__sycl_kernel_ff_2Piiii
+// CHECK-NEXT:   {{.*}}__sycl_kernel_ff_3IiEvPT_S0_S0_
+// CHECK-NEXT:   {{.*}}__sycl_kernel_ff_3IfEvPT_S0_S0_
+// CHECK-NEXT:   {{.*}}__sycl_kernel_ff_3IdEvPT_S0_S0_
+// CHECK-NEXT: };
+
+// CHECK:      const kernel_param_desc_t kernel_signatures[] = {
+// CHECK-NEXT:   {{.*}}__sycl_kernel_ff_2Piii
+// CHECK-NEXT:   { kernel_param_kind_t::kind_pointer, 8, 0 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 8 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 12 },
+
+// CHECK:        {{.*}}__sycl_kernel_ff_2Piiii
+// CHECK-NEXT:   { kernel_param_kind_t::kind_pointer, 8, 0 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 8 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 12 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 16 },
+
+// CHECK:        {{.*}}__sycl_kernel_ff_3IiEvPT_S0_S0_
+// CHECK-NEXT:   { kernel_param_kind_t::kind_pointer, 8, 0 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 8 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 12 },
+
+// CHECK:        {{.*}}__sycl_kernel_ff_3IfEvPT_S0_S0_
+// CHECK-NEXT:   { kernel_param_kind_t::kind_pointer, 8, 0 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 8 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 12 },
+
+// CHECK:        {{.*}}__sycl_kernel_ff_3IdEvPT_S0_S0_
+// CHECK-NEXT:   { kernel_param_kind_t::kind_pointer, 8, 0 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 8, 8 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 8, 16 },
+
+// CHECK:        { kernel_param_kind_t::kind_invalid, -987654321, -987654321 },
+// CHECK-NEXT: };
diff --git a/clang/test/SemaSYCL/free_function_kernel_params.cpp b/clang/test/SemaSYCL/free_function_kernel_params.cpp
@@ -0,0 +1,57 @@
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -ast-dump \
+// RUN: %s -o - | FileCheck %s
+// This test checks parameter rewriting for free functions with parameters
+// of type scalar and pointer.
+
+#include "sycl.hpp"
+
+__attribute__((sycl_device))
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
+void ff_2(int *ptr, int start, int end) {
+  for (int i = start; i <= end; i++)
+    ptr[i] = start;
+}
+// CHECK: FunctionDecl {{.*}}__sycl_kernel_{{.*}} 'void (__global int *, int, int)'
+// CHECK-NEXT: ParmVarDecl {{.*}} __arg_ptr '__global int *'
+// CHECK-NEXT: ParmVarDecl {{.*}} __arg_start 'int'
+// CHECK-NEXT: ParmVarDecl {{.*}} __arg_end 'int'
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: CallExpr {{.*}} 'void'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'void (*)(int *, int, int)' <FunctionToPointerDecay>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'void (int *, int, int)' lvalue Function {{.*}} 'ff_2' 'void (int *, int, int)'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '__arg_ptr' '__global int *'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '__arg_start' 'int'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '__arg_end' 'int'
+
+
+// Templated free function definition.
+template <typename T>
+__attribute__((sycl_device))
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
+  void ff_3(T* ptr, T start, int end) {
+    for (int i = start; i <= end; i++)
+      ptr[i] = start;
+}
+
+// Explicit instantiation with "int*"
+template void ff_3(int* ptr, int start, int end);
+
+// CHECK: FunctionDecl {{.*}}__sycl_kernel_{{.*}} 'void (__global int *, int, int)'
+// CHECK-NEXT: ParmVarDecl {{.*}} __arg_ptr '__global int *'
+// CHECK-NEXT: ParmVarDecl {{.*}} __arg_start 'int'
+// CHECK-NEXT: ParmVarDecl {{.*}} __arg_end 'int'
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: CallExpr {{.*}} 'void'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'void (*)(int *, int, int)' <FunctionToPointerDecay>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'void (int *, int, int)' lvalue Function {{.*}} 'ff_3' 'void (int *, int, int)'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '__arg_ptr' '__global int *'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '__arg_start' 'int'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '__arg_end' 'int'
diff --git a/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp b/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp
@@ -61,6 +61,16 @@ struct device_has_key
                                  std::integral_constant<aspect, Aspects>...>;
 };
 
+struct nd_range_kernel_key {
+  template <int Dims>
+  using value_t =
+      property_value<nd_range_kernel_key, std::integral_constant<int, Dims>>;
+};
+
+struct single_task_kernel_key {
+  using value_t = property_value<single_task_kernel_key>;
+};
+
 template <size_t Dim0, size_t... Dims>
 struct property_value<work_group_size_key, std::integral_constant<size_t, Dim0>,
                       std::integral_constant<size_t, Dims>...> {
@@ -113,6 +123,21 @@ struct property_value<device_has_key,
   static constexpr std::array<aspect, sizeof...(Aspects)> value{Aspects...};
 };
 
+template <int Dims>
+struct property_value<nd_range_kernel_key, std::integral_constant<int, Dims>> {
+  static_assert(
+      Dims >= 1 && Dims <= 3,
+      "nd_range_kernel_key property must use dimension of 1, 2 or 3.");
+
+  using key_t = nd_range_kernel_key;
+  using value_t = int;
+  static constexpr int dimensions = Dims;
+};
+
+template <> struct property_value<single_task_kernel_key> {
+  using key_t = single_task_kernel_key;
+};
+
 template <size_t Dim0, size_t... Dims>
 inline constexpr work_group_size_key::value_t<Dim0, Dims...> work_group_size;
 
@@ -126,6 +151,11 @@ inline constexpr sub_group_size_key::value_t<Size> sub_group_size;
 template <aspect... Aspects>
 inline constexpr device_has_key::value_t<Aspects...> device_has;
 
+template <int Dims>
+inline constexpr nd_range_kernel_key::value_t<Dims> nd_range_kernel;
+
+inline constexpr single_task_kernel_key::value_t single_task_kernel;
+
 struct work_group_progress_key
     : detail::compile_time_property_key<detail::PropKind::WorkGroupProgress> {
   template <forward_progress_guarantee Guarantee,
@@ -209,6 +239,7 @@ template <> struct is_property_key<sub_group_progress_key> : std::true_type {};
 template <> struct is_property_key<work_item_progress_key> : std::true_type {};
 
 namespace detail {
+
 template <size_t Dim0, size_t... Dims>
 struct PropertyMetaInfo<work_group_size_key::value_t<Dim0, Dims...>> {
   static constexpr const char *name = "sycl-work-group-size";
@@ -230,6 +261,15 @@ struct PropertyMetaInfo<device_has_key::value_t<Aspects...>> {
   static constexpr const char *value =
       SizeListToStr<static_cast<size_t>(Aspects)...>::value;
 };
+template <int Dims>
+struct PropertyMetaInfo<nd_range_kernel_key::value_t<Dims>> {
+  static constexpr const char *name = "sycl-nd-range-kernel";
+  static constexpr int value = Dims;
+};
+template <> struct PropertyMetaInfo<single_task_kernel_key::value_t> {
+  static constexpr const char *name = "sycl-single-task-kernel";
+  static constexpr int value = 0;
+};
 
 template <typename T, typename = void>
 struct HasKernelPropertiesGetMethod : std::false_type {};
@@ -251,7 +291,6 @@ struct HasKernelPropertiesGetMethod<T,
 #ifdef __SYCL_DEVICE_ONLY__
 #define SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(PROP)                                \
   [[__sycl_detail__::add_ir_attributes_function(                               \
-      {"sycl-device-has"},                                                     \
       sycl::ext::oneapi::experimental::detail::PropertyMetaInfo<               \
           std::remove_cv_t<std::remove_reference_t<decltype(PROP)>>>::name,    \
       sycl::ext::oneapi::experimental::detail::PropertyMetaInfo<               \
diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp
@@ -203,8 +203,10 @@ enum PropKind : uint32_t {
   WorkGroupProgress = 62,
   SubGroupProgress = 63,
   WorkItemProgress = 64,
+  NDRangeKernel = 65,
+  SingleTaskKernel = 66,
   // PropKindSize must always be the last value.
-  PropKindSize = 65,
+  PropKindSize = 67,
 };
 
 struct property_key_base_tag {};
diff --git a/sycl/source/feature_test.hpp.in b/sycl/source/feature_test.hpp.in
@@ -105,6 +105,7 @@ inline namespace _V1 {
 #define SYCL_EXT_INTEL_FPGA_TASK_SEQUENCE 1
 #define SYCL_EXT_ONEAPI_PRIVATE_ALLOCA 1
 #define SYCL_EXT_ONEAPI_FORWARD_PROGRESS 1
+#define SYCL_EXT_ONEAPI_FREE_FUNCTION_KERNELS 1
 
 #ifndef __has_include
 #define __has_include(x) 0
diff --git a/sycl/test-e2e/KernelAndProgram/free_function_kernels.cpp b/sycl/test-e2e/KernelAndProgram/free_function_kernels.cpp
