Address more comments

Author: Alexander Batashev

sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp

@@ -61,10 +61,8 @@ class accessor_property_list : protected sycl::detail::PropertyListBase {
   template <typename PropT> struct ContainsProperty<PropT> : std::false_type {};
   template <typename PropT, typename Head, typename... Tail>
   struct ContainsProperty<PropT, Head, Tail...>
-      : std::conditional<
-            AreSameTemplate<PropT, Head>::value ||
-                !sycl::detail::IsCompileTimePropertyInstance<PropT>::value,
-            std::true_type, ContainsProperty<PropT, Tail...>>::type {};
+      : std::conditional<AreSameTemplate<PropT, Head>::value, std::true_type,
+                         ContainsProperty<PropT, Tail...>>::type {};
 
   // PropertyContainer is a helper structure, that holds list of properties.
   // It is used to avoid multiple parameter packs in templates.
@@ -90,10 +88,8 @@ class accessor_property_list : protected sycl::detail::PropertyListBase {
   struct ContainsPropertyInstance
       : std::conditional_t<
             !std::is_same_v<typename ContainerT::Head, void> &&
-                (AreSameTemplate<PropT<Args...>,
-                                 typename ContainerT::Head>::value ||
-                 !sycl::detail::IsCompileTimePropertyInstance<
-                     typename ContainerT::Head>::value),
+                AreSameTemplate<PropT<Args...>,
+                                typename ContainerT::Head>::value,
             std::true_type,
             ContainsPropertyInstance<typename ContainerT::Rest, PropT,
                                      Args...>> {};
@@ -103,11 +99,14 @@ class accessor_property_list : protected sycl::detail::PropertyListBase {
 #endif
 
   // This template checks if two lists of properties contain the same set of
-  // compile-time-constant properties in any order.
-  template <typename ContainerT, typename... OtherProps>
+  // compile-time-constant properties in any order. Run time properties are
+  // skipped.
   struct ContainsSameProperties
       : std::conditional<
-            ContainsProperty<typename ContainerT::Head, OtherProps...>::value,
+            !detail::IsCompileTimePropertyInstance<
+                typename ContainerT::Head>::value ||
+                ContainsProperty<typename ContainerT::Head,
+                                 OtherProps...>::value,
             ContainsSameProperties<typename ContainerT::Rest, OtherProps...>,
             std::false_type>::type {};
   template <typename... OtherProps>

sycl/include/CL/sycl/properties/accessor_properties.hpp

@@ -39,19 +39,13 @@ namespace INTEL {
 namespace property {
 struct buffer_location {
   template <int A> struct instance {
-    constexpr bool operator==(const buffer_location::instance<A> &) const {
-      return true;
-    }
-    constexpr bool operator!=(const buffer_location::instance<A> &) const {
-      return false;
-    }
     template <int B>
     constexpr bool operator==(const buffer_location::instance<B> &) const {
-      return false;
+      return A == B;
     }
     template <int B>
     constexpr bool operator!=(const buffer_location::instance<B> &) const {
-      return true;
+      return A != B;
     }
   };
 };

sycl/test/basic_tests/accessor/accessor_property_list_rt.cpp

@@ -11,6 +11,7 @@ int main() {
   {
     // Single RT property
     accessor_property_list PL{sycl::noinit};
+    static_assert(!PL.has_property<property::no_offset>(), "Property is found");
     assert(PL.has_property<sycl::property::noinit>() && "Property not found");
   }
 

sycl/test/fpga_tests/buffer_location.cpp

@@ -0,0 +1,22 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUNx: %ACC_RUN_PLACEHOLDER %t.out
+
+#include <CL/sycl.hpp>
+
+int main() {
+  sycl::queue Queue;
+  sycl::buffer<int, 1> Buf{sycl::range{1}};
+
+  Queue.submit([&](sycl::handler &CGH) {
+    sycl::ONEAPI::accessor_property_list PL{sycl::INTEL::buffer_location<1>};
+    sycl::accessor Acc(Buf, CGH, sycl::write_only, PL);
+    CGH.single_task<class Test>([=]() { Acc[0] = 42; });
+  });
+
+  Queue.wait();
+
+  auto Acc = Buf.template get_access<sycl::access::mode::read_write>();
+  assert(Acc[0] == 42 && "Value mismatch");
+
+  return 0;
+}

sycl/test/fpga_tests/buffer_location_codegen.cpp

@@ -0,0 +1,47 @@
+// RUN: %clangxx -fsycl -c -fsycl-device-only -S -emit-llvm %s -o - | FileCheck %s
+
+// CHECK: define {{.*}}spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE15kernel_function"{{.*}} !kernel_arg_buffer_location ![[MDBL:[0-9]+]]
+// CHECK: ![[MDBL]] = !{i32 3, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 2, i32 -1, i32 -1, i32 -1, i32 2, i32 -1, i32 -1, i32 -1, i32 -1}
+
+#include <CL/sycl.hpp>
+
+struct Base {
+  int A, B;
+  cl::sycl::accessor<
+      char, 1, cl::sycl::access::mode::read,
+      cl::sycl::access::target::global_buffer,
+      cl::sycl::access::placeholder::false_t,
+      cl::sycl::ONEAPI::accessor_property_list<
+          cl::sycl::INTEL::property::buffer_location::instance<2>>>
+      AccField;
+};
+
+struct Captured
+    : Base,
+      cl::sycl::accessor<
+          char, 1, cl::sycl::access::mode::read,
+          cl::sycl::access::target::global_buffer,
+          cl::sycl::access::placeholder::false_t,
+          cl::sycl::ONEAPI::accessor_property_list<
+              cl::sycl::INTEL::property::buffer_location::instance<2>>> {
+  int C;
+};
+
+int main() {
+  Captured Obj;
+  cl::sycl::accessor<
+      int, 1, cl::sycl::access::mode::read_write,
+      cl::sycl::access::target::global_buffer,
+      cl::sycl::access::placeholder::false_t,
+      cl::sycl::ONEAPI::accessor_property_list<
+          cl::sycl::INTEL::property::buffer_location::instance<3>>>
+      accessorA;
+  sycl::queue Queue;
+  Queue.submit([&](sycl::handler &CGH) {
+    CGH.single_task<class kernel_function>([=]() {
+      (int)accessorA[0];
+      (int)Obj[0];
+    });
+  });
+  return 0;
+}