[PATCH] D153883: [Clang][OpenMP] Enable use of __kmpc_alloc_shared for VLAs defined in AMD GPU offloaded regions

Gheorghe-Teodor Bercea via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Tue Jun 27 08:11:04 PDT 2023


doru1004 created this revision.
doru1004 added reviewers: ronlieb, gregrodgers, carlo.bertolli, arsenm, jdoerfert, dhruvachak, ABataev.
doru1004 added a project: OpenMP.
Herald added subscribers: sunshaoce, guansong, yaxunl, jvesely.
Herald added a project: All.
doru1004 requested review of this revision.
Herald added subscribers: cfe-commits, jplehr, sstefan1, wdng.
Herald added a project: clang.

This patch enables the use of `___kmpc_alloc_shared` to allocate dynamically sized allocation on AMD GPUs. For example:

  #pragma omp target
  {
    int N = 10;
    double A[N];
    ...
  }

This will generate a pair of `__kmpc_alloc_shared / __kmpc_free_shared` to handle the allocation and deallocation of `A` inside the target region.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D153883

Files:
  clang/lib/CodeGen/CGDecl.cpp
  clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
  clang/lib/CodeGen/CGOpenMPRuntimeGPU.h
  clang/lib/CodeGen/CodeGenFunction.h
  clang/test/OpenMP/amdgcn_target_device_vla.cpp

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D153883.534992.patch
Type: text/x-patch
Size: 71449 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20230627/4f2f912e/attachment-0001.bin>


More information about the cfe-commits mailing list