Add fmoreimprovments

Author: Victor Lomuller

llvm/lib/SYCLLowerIR/LowerWGLocalMemory.cpp

@@ -131,13 +131,6 @@ lowerDynamicLocalMemCallDirect(CallInst *CI, Triple TT,
     IRBuilder<> Builder(CI);
     if (TT.isSPIROrSPIRV())
       return Builder.CreateLoad(CI->getType(), LocalMemPlaceholder);
-    Value *ArgAlign = CI->getArgOperand(0);
-    Align RequestedAlignment{
-        cast<llvm::ConstantInt>(ArgAlign)->getZExtValue()};
-    MaybeAlign CurrentAlignment = LocalMemPlaceholder->getAlign();
-    if (!CurrentAlignment.has_value() ||
-        (CurrentAlignment.value() < RequestedAlignment))
-      LocalMemPlaceholder->setAlignment(RequestedAlignment);
 
     return Builder.CreatePointerCast(LocalMemPlaceholder, CI->getType());
   }();
@@ -211,6 +204,8 @@ static bool dynamicWGLocalMemory(Module &M) {
         GlobalVariable::NotThreadLocal, // ThreadLocalMode
         LocalAS                         // AddressSpace
     );
+    if (!TT.isSPIROrSPIRV())
+      LocalMemArrayGV->setAlignment(Align{128});
   }
   lowerLocalMemCall(DLMFunc, [&](CallInst *CI) {
     lowerDynamicLocalMemCallDirect(CI, TT, LocalMemArrayGV);

llvm/test/SYCLLowerIR/work_group_static.ll

@@ -6,18 +6,18 @@
 target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
 target triple = "spir64-unknown-unknown"
 
-; CHECK-DAG: @__sycl_dynamicLocalMemoryPlaceholder_GV = linkonce_odr addrspace(3) global ptr addrspace(3) undef
+; CHECK: @__sycl_dynamicLocalMemoryPlaceholder_GV = linkonce_odr addrspace(3) global ptr addrspace(3) undef
 
 ; Function Attrs: convergent norecurse
-; CHECK-DAG: @_ZTS7KernelA(ptr addrspace(1) %0, ptr addrspace(3) noalias "sycl-implicit-local-arg" %[[IMPLICT_ARG:[a-zA-Z0-9]+]]{{.*}} !kernel_arg_addr_space ![[ADDR_SPACE_MD:[0-9]+]]
+; CHECK: @_ZTS7KernelA(ptr addrspace(1) %0, ptr addrspace(3) noalias "sycl-implicit-local-arg" %[[IMPLICT_ARG:[a-zA-Z0-9]+]]{{.*}} !kernel_arg_addr_space ![[ADDR_SPACE_MD:[0-9]+]]
 define weak_odr dso_local spir_kernel void @_ZTS7KernelA(ptr addrspace(1) %0) local_unnamed_addr #0 !kernel_arg_addr_space !5 {
 entry:
-; CHECK-DAG: store ptr addrspace(3) %[[IMPLICT_ARG]], ptr addrspace(3) @__sycl_dynamicLocalMemoryPlaceholder_GV
-; CHECK-DAG: %[[LD1:[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr addrspace(3) @__sycl_dynamicLocalMemoryPlaceholder_GV
+; CHECK: store ptr addrspace(3) %[[IMPLICT_ARG]], ptr addrspace(3) @__sycl_dynamicLocalMemoryPlaceholder_GV
+; CHECK: %[[LD1:[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr addrspace(3) @__sycl_dynamicLocalMemoryPlaceholder_GV
   %1 = tail call spir_func ptr addrspace(3) @__sycl_dynamicLocalMemoryPlaceholder(i64 128) #1
-; CHECK-DAG: getelementptr inbounds i8, ptr addrspace(3) %[[LD1]]
+; CHECK: getelementptr inbounds i8, ptr addrspace(3) %[[LD1]]
   %2 = getelementptr inbounds i8, ptr addrspace(3) %1, i64 4
-; CHECK-DAG: %[[LD2:[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr addrspace(3) @__sycl_dynamicLocalMemoryPlaceholder_GV
+; CHECK: %[[LD2:[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr addrspace(3) @__sycl_dynamicLocalMemoryPlaceholder_GV
   %3 = tail call spir_func ptr addrspace(3) @__sycl_dynamicLocalMemoryPlaceholder(i64 4) #1
   ret void
 }

llvm/test/SYCLLowerIR/work_group_static_nv.ll

@@ -0,0 +1,38 @@
+; RUN: opt -S -sycllowerwglocalmemory -bugpoint-enable-legacy-pm < %s | FileCheck %s
+; RUN: opt -S -passes=sycllowerwglocalmemory < %s | FileCheck %s
+
+; CHECK-NOT: __sycl_dynamicLocalMemoryPlaceholder
+
+target triple = "nvptx64-nvidia-cuda"
+
+; CHECK: @__sycl_dynamicLocalMemoryPlaceholder_GV = external addrspace(3) global [0 x i8], align 128
+
+; Function Attrs: convergent norecurse
+; CHECK: @_ZTS7KernelA(ptr addrspace(1) %0)
+define void @_ZTS7KernelA(ptr addrspace(1) %0) local_unnamed_addr #0 !kernel_arg_addr_space !5 {
+entry:
+; CHECK: getelementptr inbounds i8, ptr addrspace(3) @__sycl_dynamicLocalMemoryPlaceholder_GV
+  %1 = tail call spir_func ptr addrspace(3) @__sycl_dynamicLocalMemoryPlaceholder(i64 128) #1
+  %2 = getelementptr inbounds i8, ptr addrspace(3) %1, i64 4
+  %3 = tail call spir_func ptr addrspace(3) @__sycl_dynamicLocalMemoryPlaceholder(i64 4) #1
+  ret void
+}
+
+; Function Attrs: convergent
+declare dso_local spir_func ptr addrspace(3) @__sycl_dynamicLocalMemoryPlaceholder(i64) local_unnamed_addr #1
+
+attributes #0 = { convergent norecurse "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" "sycl-work-group-static"="1" }
+attributes #1 = { convergent norecurse }
+
+!llvm.module.flags = !{!0}
+!opencl.spir.version = !{!1}
+!spirv.Source = !{!2}
+!llvm.ident = !{!3}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 1, i32 2}
+!2 = !{i32 4, i32 100000}
+!3 = !{!"clang version 13.0.0"}
+!4 = !{}
+; ![[ADDR_SPACE_MD]] = !{i32 1, i32 3}
+!5 = !{i32 1}

sycl/doc/extensions/proposed/sycl_ext_oneapi_work_group_static.asciidoc

@@ -195,10 +195,7 @@ to a dynamically allocated buffer in the device local memory.
 
 [source,c++]
 ----
-template <typename T>
-sycl::multi_ptr<T, access::address_space::local_space,
-                           access::decorated::no>
-get_dynamic_work_group_memory()
+void* get_dynamic_work_group_memory()
 ----
 _Constraints_: `T` must be trivially constructible and trivially destructible.
 
@@ -209,9 +206,7 @@ The size of the allocation is unknown at compile-time,
 and must be communicated to the SYCL implementation via the
 `work_group_static_memory_size` property. Every call to
 `get_dynamic_work_group_memory` returns the same allocation
-in device local memory, regardless of `T`. For example, two call declared
-as `get_dynamic_work_group_memory<int>` and
-`get_dynamic_work_group_memory<float>` will be associated with the same shared allocation.
+in device local memory.
 
 === Kernel properties
 
@@ -290,8 +285,8 @@ using namespace syclex = sycl::ext::oneapi::experimental;
 q.parallel_for(sycl::nd_range<1>{N, M},
   syclex::properties{syclex::work_group_static_size(M * sizeof(int))},
   [=](sycl::nd_item<1> it) {
- auto ptr= get_dynamic_work_group_memory<int>();
- auto ptr2= get_dynamic_work_group_memory<float>();
+ auto ptr= get_dynamic_work_group_memory();
+ auto ptr2= get_dynamic_work_group_memory();
 });
 ----
 
@@ -302,13 +297,9 @@ This non-normative section provides information about one possible
 implementation of this extension.  It is not part of the specification of the
 extension's API.
 
-For class types and bounded arrays, the class can be implemented using
-the WG scope attribute used by hierarchical implementation.
-Note, however, that this requires to expose the attribute to the users.
-
 For `get_dynamic_work_group_memory`,
-and the implementation may need to generate some additional code to
-appropriately initialize the pointer(s) returns by the call.
+the implementation may need to generate some additional code to
+appropriately initialize the pointer(s) returned by the call.
 Alternatively, it may be possible to initialize the pointer to the beginning
 of the device's local memory region (if that value is known). Either way, the
 implementation must account for the existence of one or more `local_accessor`

sycl/include/sycl/access/access.hpp

@@ -23,7 +23,7 @@ __sycl_allocateLocalMemory(std::size_t Size, std::size_t Alignment);
 // Request a placeholder for a dynamically-sized buffer in local address space
 // at kernel scope. Required for work_group_static.
 extern "C" __DPCPP_SYCL_EXTERNAL __attribute__((opencl_local)) std::uint8_t *
-__sycl_dynamicLocalMemoryPlaceholder(std::size_t Alignment);
+__sycl_dynamicLocalMemoryPlaceholder();
 #endif
 namespace sycl {
 inline namespace _V1 {

sycl/include/sycl/ext/oneapi/work_group_static.hpp

@@ -46,7 +46,7 @@ template <typename T> class __SYCL_WG_SCOPE work_group_static final {
   template <class TArg = T>
   typename std::enable_if<!std::is_array_v<TArg>, work_group_static &>::type
   operator=(const T &value) noexcept {
-    *getDecorated() = value;
+    data = value;
     return *this;
   }
 
@@ -56,16 +56,10 @@ template <typename T> class __SYCL_WG_SCOPE work_group_static final {
   T data;
 };
 
-template <typename T>
-std::enable_if_t<
-    std::is_trivially_destructible_v<T> && std::is_trivially_constructible_v<T>,
-    multi_ptr<T, access::address_space::local_space, access::decorated::no>>
-    __SYCL_ALWAYS_INLINE get_dynamic_work_group_memory() {
+__SYCL_ALWAYS_INLINE
+inline void* get_dynamic_work_group_memory() {
 #ifdef __SYCL_DEVICE_ONLY__
-  return multi_ptr<T, access::address_space::local_space,
-                   access::decorated::no>{
-      reinterpret_cast<__attribute__((opencl_local)) T *>(
-          __sycl_dynamicLocalMemoryPlaceholder(alignof(T)))};
+  return __sycl_dynamicLocalMemoryPlaceholder();
 #else
   throw sycl::exception(
       sycl::errc::feature_not_supported,

sycl/test-e2e/WorkGroupStaticMemory/copy_dynamic_size.cpp

@@ -17,7 +17,7 @@ void copy_via_smem(DataType *a, DataType *b, sycl::nd_item<1> it) {
   // And then puts in back into B
 
   DataType *smem_ptr =
-      sycl_ext::get_dynamic_work_group_memory<DataType>().get();
+      reinterpret_cast<DataType *>(sycl_ext::get_dynamic_work_group_memory());
   auto threadIdx_x = it.get_local_linear_id();
 
   smem_ptr[threadIdx_x] = a[threadIdx_x];

sycl/test-e2e/WorkGroupStaticMemory/dynamic_alloc_local_accessor.cpp

@@ -34,10 +34,8 @@ int main() {
         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) {
-                       multi_ptr<int, access::address_space::local_space,
-                                 sycl::access::decorated::no>
-                           Ptr = sycl::ext::oneapi::experimental::
-                               get_dynamic_work_group_memory<int>();
+                       int *Ptr = reinterpret_cast<int *>(
+                           sycl_ext::get_dynamic_work_group_memory());
                        size_t GroupOffset =
                            Item.get_group_linear_id() * ElemPerWG;
                        for (size_t I = 0; I < RepeatWG; ++I) {

sycl/test-e2e/WorkGroupStaticMemory/dynamic_alloc_ptr_alias.cpp

@@ -30,35 +30,32 @@ int main() {
     sycl_ext::work_group_static_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) {
-          multi_ptr<int, access::address_space::local_space,
-                    sycl::access::decorated::no>
-              Ptr = sycl::ext::oneapi::experimental::
-                  get_dynamic_work_group_memory<int>();
-          size_t GroupOffset = Item.get_group_linear_id() * ElemPerWG;
-          for (size_t I = 0; I < RepeatWG; ++I) {
-            Ptr[WgSize * I + Item.get_local_linear_id()] =
-                Item.get_local_linear_id();
-          }
+    Cgh.parallel_for(nd_range<1>(range<1>(Size), range<1>(WgSize)), properties,
+                     [=](nd_item<1> Item) {
+                       int *Ptr = reinterpret_cast<int *>(
+                           sycl_ext::get_dynamic_work_group_memory());
+                       size_t GroupOffset =
+                           Item.get_group_linear_id() * ElemPerWG;
+                       for (size_t I = 0; I < RepeatWG; ++I) {
+                         Ptr[WgSize * I + Item.get_local_linear_id()] =
+                             Item.get_local_linear_id();
+                       }
 
-          Item.barrier();
-          // Check that multiple calls return the same pointer.
-          multi_ptr<unsigned int, access::address_space::local_space,
-                    sycl::access::decorated::no>
-              PtrAlias = sycl::ext::oneapi::experimental::
-                  get_dynamic_work_group_memory<unsigned int>();
+                       Item.barrier();
+                       // Check that multiple calls return the same pointer.
+                       unsigned int *PtrAlias =
+                           reinterpret_cast<unsigned int *>(
+                               sycl_ext::get_dynamic_work_group_memory());
 
-          for (size_t I = 0; I < RepeatWG; ++I) {
-            // Check that the memory is accessible from other
-            // work-items
-            size_t BaseIdx = GroupOffset + (I * WgSize);
-            size_t LocalIdx = Item.get_local_linear_id() ^ 1;
-            size_t GlobalIdx = BaseIdx + LocalIdx;
-            Acc[GlobalIdx] = PtrAlias[WgSize * I + LocalIdx];
-          }
-        });
+                       for (size_t I = 0; I < RepeatWG; ++I) {
+                         // Check that the memory is accessible from other
+                         // work-items
+                         size_t BaseIdx = GroupOffset + (I * WgSize);
+                         size_t LocalIdx = Item.get_local_linear_id() ^ 1;
+                         size_t GlobalIdx = BaseIdx + LocalIdx;
+                         Acc[GlobalIdx] = PtrAlias[WgSize * I + LocalIdx];
+                       }
+                     });
   });
 
   host_accessor Acc(Buf, read_only);

sycl/test-e2e/WorkGroupStaticMemory/dynamic_allocation.cpp

@@ -31,10 +31,8 @@ int main() {
     sycl_ext::properties properties{static_size};
     Cgh.parallel_for(nd_range<1>(range<1>(Size), range<1>(WgSize)), properties,
                      [=](nd_item<1> Item) {
-                       multi_ptr<int, access::address_space::local_space,
-                                 sycl::access::decorated::no>
-                           Ptr = sycl::ext::oneapi::experimental::
-                               get_dynamic_work_group_memory<int>();
+                       int *Ptr = reinterpret_cast<int *>(
+                           sycl_ext::get_dynamic_work_group_memory());
                        size_t GroupOffset =
                            Item.get_group_linear_id() * ElemPerWG;
                        for (size_t I = 0; I < RepeatWG; ++I) {