fix work_group_scrach_memory

Signed-off-by: Victor Lomuller <[email protected]>

Author: Victor Lomuller

llvm/lib/SYCLLowerIR/LowerWGLocalMemory.cpp

@@ -205,6 +205,7 @@ static bool dynamicWGLocalMemory(Module &M) {
         GlobalVariable::NotThreadLocal, // ThreadLocalMode
         LocalAS                         // AddressSpace
     );
+    LocalMemArrayGV->setUnnamedAddr(GlobalVariable::UnnamedAddr::Local);
     constexpr int DefaultMaxAlignment = 128;
     if (!TT.isSPIROrSPIRV())
       LocalMemArrayGV->setAlignment(Align{DefaultMaxAlignment});

sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_local_accessor.cpp

@@ -2,7 +2,7 @@
 // RUN: %{run} %t.out
 //
 
-// UNSUPPORTED: gpu-intel-gen12, cpu
+// UNSUPPORTED: gpu-intel-gen12
 // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/16072
 
 // Test work_group_dynamic extension with allocation size specified at runtime
@@ -35,8 +35,8 @@ int main() {
     sycl_ext::properties properties{static_size};
     auto LocalAccessor =
         sycl::local_accessor<int>(WgSize * RepeatWG * sizeof(int), Cgh);
-    Cgh.parallel_for(nd_range<1>(range<1>(Size), range<1>(WgSize)), properties,
-                     [=](nd_item<1> Item) {
+    Cgh.parallel_for(nd_range<1>(range<1>(WgSize * WgCount), range<1>(WgSize)),
+                     properties, [=](nd_item<1> Item) {
                        int *Ptr = reinterpret_cast<int *>(
                            sycl_ext::get_work_group_scratch_memory());
                        size_t GroupOffset =

sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_ptr_alias.cpp

@@ -2,7 +2,7 @@
 // RUN: %{run} %t.out
 //
 
-// UNSUPPORTED: gpu-intel-gen12, cpu
+// UNSUPPORTED: gpu-intel-gen12
 // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/16072
 
 // Test work_group_dynamic extension with allocation size specified at runtime
@@ -33,8 +33,8 @@ int main() {
     sycl_ext::work_group_scratch_size static_size(WgSize * RepeatWG *
                                                   sizeof(int));
     sycl_ext::properties properties{static_size};
-    Cgh.parallel_for(nd_range<1>(range<1>(Size), range<1>(WgSize)), properties,
-                     [=](nd_item<1> Item) {
+    Cgh.parallel_for(nd_range<1>(range<1>(WgSize * WgCount), range<1>(WgSize)),
+                     properties, [=](nd_item<1> Item) {
                        int *Ptr = reinterpret_cast<int *>(
                            sycl_ext::get_work_group_scratch_memory());
                        size_t GroupOffset =

sycl/test-e2e/WorkGroupScratchMemory/dynamic_allocation.cpp

@@ -2,7 +2,7 @@
 // RUN: %{run} %t.out
 //
 
-// UNSUPPORTED: gpu-intel-gen12, cpu
+// UNSUPPORTED: gpu-intel-gen12
 // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/16072
 
 // Test work_group_dynamic extension with allocation size specified at runtime.
@@ -32,8 +32,8 @@ int main() {
     sycl_ext::work_group_scratch_size static_size(WgSize * RepeatWG *
                                                   sizeof(int));
     sycl_ext::properties properties{static_size};
-    Cgh.parallel_for(nd_range<1>(range<1>(Size), range<1>(WgSize)), properties,
-                     [=](nd_item<1> Item) {
+    Cgh.parallel_for(nd_range<1>(range<1>(WgSize * WgCount), range<1>(WgSize)),
+                     properties, [=](nd_item<1> Item) {
                        int *Ptr = reinterpret_cast<int *>(
                            sycl_ext::get_work_group_scratch_memory());
                        size_t GroupOffset =