Append cache controls instead of overwriting.

Author: maarquitos14

llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp

@@ -950,7 +950,10 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation(
     if (!FPGAProp && llvm::isa<llvm::Instruction>(IntrInst->getArgOperand(0))) {
       // If there are no annotations other than cache controls we can apply the
       // controls to the pointer and remove the intrinsic.
-      auto PtrInstr = cast<Instruction>(IntrInst->getArgOperand(0));
+      Instruction *PtrInstr = cast<Instruction>(IntrInst->getArgOperand(0));
+      if (MDNode *CurrentMD = PtrInstr->getMetadata(MDKindID))
+        for (Metadata *Op : CurrentMD->operands())
+          MDOpsCacheProp.push_back(Op);
       PtrInstr->setMetadata(MDKindID, MDTuple::get(Ctx, MDOpsCacheProp));
       // Replace all uses of IntrInst with first operand
       IntrInst->replaceAllUsesWith(PtrInstr);

sycl/test/check_device_code/extensions/properties/properties_cache_control.cpp

@@ -7,6 +7,15 @@ using namespace sycl;
 using namespace ext::oneapi::experimental;
 using namespace ext::intel::experimental;
 
+using ST_L1 = cache_control<cache_mode::streaming, cache_level::L1>;
+using WB_L1 = cache_control<cache_mode::write_back, cache_level::L1>;
+using UC_L1 = cache_control<cache_mode::uncached, cache_level::L1>;
+using CA_L1 = cache_control<cache_mode::cached, cache_level::L1>;
+using UC_L2 = cache_control<cache_mode::uncached, cache_level::L2>;
+using CA_L2 = cache_control<cache_mode::cached, cache_level::L2>;
+using UC_L3 = cache_control<cache_mode::uncached, cache_level::L3>;
+using CA_L3 = cache_control<cache_mode::cached, cache_level::L3>;
+
 using load_hint = annotated_ptr<
     float, decltype(properties(
                read_hint<cache_control<cache_mode::cached, cache_level::L1>,
@@ -33,6 +42,22 @@ using load_store_hint = annotated_ptr<
         write_hint<
             cache_control<cache_mode::write_through, cache_level::L4>>))>;
 
+template <typename t>
+using ap_load_ca_uc_uc =
+    annotated_ptr<t, decltype(properties(read_hint<CA_L1, UC_L2, UC_L3>))>;
+
+template <typename t>
+using ap_load_st_ca_uc =
+    annotated_ptr<t, decltype(properties(read_hint<ST_L1, CA_L2, CA_L3>))>;
+
+template <typename T>
+using ap_store_uc_uc_uc =
+    annotated_ptr<T, decltype(properties(write_hint<UC_L1, UC_L2, UC_L3>))>;
+
+template <typename T>
+using ap_store_wb_uc_uc =
+    annotated_ptr<T, decltype(properties(write_hint<WB_L1, UC_L2, UC_L3>))>;
+
 void cache_control_read_hint_func() {
   queue q;
   constexpr int N = 10;
@@ -81,6 +106,65 @@ void cache_control_read_write_func() {
   });
 }
 
+void cache_control_load_store_func() {
+  queue q(gpu_selector_v);
+
+  constexpr int N = 512;
+  double *x_buf = malloc_device<double>(N, q);
+  double *y_buf = malloc_device<double>(N, q);
+  double *d_buf = malloc_device<double>(1, q);
+  double *d_buf_h = malloc_host<double>(1, q);
+
+  q.fill<double>(d_buf, 0.0, 1).wait();
+
+  constexpr int SG_SIZE = 16;
+
+  q.submit([&](handler &cgh) {
+     const int nwg = N / SG_SIZE;
+     auto x = x_buf;
+     auto y = y_buf;
+     auto d = d_buf;
+     auto d_h = d_buf_h;
+
+     auto kernel =
+         [=](nd_item<2> item) [[intel::reqd_sub_group_size(SG_SIZE)]] {
+           const int global_tid = item.get_global_id(0);
+           const int row_st = global_tid * SG_SIZE;
+
+           if (row_st > N)
+             return;
+
+           const sub_group sgr = item.get_sub_group();
+           const int sgr_tid = sgr.get_local_id();
+
+           ap_store_uc_uc_uc<double> x_s(x);
+           ap_store_wb_uc_uc<double> y_s(y);
+
+           x_s[row_st + sgr_tid] = 1.0;
+           y_s[row_st + sgr_tid] = 1.0;
+
+           group_barrier(sgr);
+
+           ap_load_ca_uc_uc<double> x_l(x);
+           ap_load_st_ca_uc<double> y_l(y);
+
+           const double xVal = x_l[row_st + sgr_tid];
+           const double yVal = y_l[row_st + sgr_tid];
+           double T = xVal * yVal;
+           T = reduce_over_group(sgr, T, 0.0, std::plus<>());
+
+           if (sgr.leader()) {
+             atomic_ref<double, memory_order::relaxed, memory_scope::device,
+                        access::address_space::global_space>
+                 d_atomic(d[0]);
+             d_atomic.fetch_add(T);
+           }
+         };
+     cgh.parallel_for<class write_ker>(
+         nd_range<2>(range<2>(nwg, SG_SIZE), range<2>(1, SG_SIZE)), kernel);
+   }).wait();
+}
+
 // Test that annotated pointer parameter functions don't crash.
 SYCL_EXTERNAL void annotated_ptr_func_param_test(float *p) {
   *(store_hint{p}) = 42.0f;
@@ -106,6 +190,11 @@ SYCL_EXTERNAL void annotated_ptr_func_param_test(float *p) {
 // CHECK-IR: {{.*}}addrspacecast ptr addrspace(1){{.*}}!spirv.Decorations [[RWHINT:.*]]
 // CHECK-IR: ret void
 
+// CHECK-IR: spir_kernel{{.*}}cache_control_load_store_func
+// CHECK-IR: {{.*}}getelementptr{{.*}}addrspace(4){{.*}}!spirv.Decorations [[LDSTHINT_A:.*]]
+// CHECK-IR: {{.*}}getelementptr{{.*}}addrspace(4){{.*}}!spirv.Decorations [[LDSTHINT_B:.*]]
+// CHECK-IR: ret void
+
 // CHECK-IR: [[WHINT]] = !{[[WHINT1:.*]], [[WHINT2:.*]], [[WHINT3:.*]], [[WHINT4:.*]]}
 // CHECK-IR: [[WHINT1]] = !{i32 6443, i32 3, i32 3}
 // CHECK-IR: [[WHINT2]] = !{i32 6443, i32 0, i32 1}
@@ -126,3 +215,13 @@ SYCL_EXTERNAL void annotated_ptr_func_param_test(float *p) {
 // CHECK-IR: [[RWHINT1]] = !{i32 6442, i32 2, i32 1}
 // CHECK-IR: [[RWHINT2]] = !{i32 6442, i32 3, i32 4}
 // CHECK-IR: [[RWHINT3]] = !{i32 6443, i32 3, i32 1}
+
+// CHECK-IR: [[LDSTHINT_A]] = !{[[RHINT1]], [[RHINT2]], [[RHINT3]], [[LDSTHINT_A1:.*]], [[LDSTHINT_A2:.*]], [[LDSTHINT_A3:.*]]}
+// CHECK-IR: [[LDSTHINT_A1]] = !{i32 6443, i32 0, i32 0}
+// CHECK-IR: [[LDSTHINT_A2]] = !{i32 6443, i32 1, i32 0}
+// CHECK-IR: [[LDSTHINT_A3]] = !{i32 6443, i32 2, i32 0}
+
+// CHECK-IR: [[LDSTHINT_B]] = !{[[LDSTHINT_B1:.*]], [[RWHINT1]], [[LDSTHINT_B2:.*]], [[LDSTHINT_A2]], [[LDSTHINT_A3]], [[LDSTHINT_B3:.*]]}
+// CHECK-IR: [[LDSTHINT_B1]] = !{i32 6442, i32 1, i32 1}
+// CHECK-IR: [[LDSTHINT_B2]] = !{i32 6442, i32 0, i32 2}
+// CHECK-IR: [[LDSTHINT_B3]] = !{i32 6443, i32 0, i32 2}