Add std::shared_ptr support in OpenMP offloading

[ikitayama]

Currently not supported, build doesn't succeed.

[ikitayama]

Cc @jdoerfert , could you take over from here?

[llvmbot]

@llvm/issue-subscribers-openmp

[ikitayama]

wonder in which Phabricator # this is discussed.

[jhuber6]

I don't think there is a phabricator review for this. std::unique_ptr works for CUDA offloading because we have a malloc / free function. The problem with std::shared_ptr is the atomic update of the reference count, this will probably use some GNU builtin we can't create an instruction for. I would guess that this would work if you copied the libstdc++ or libc++ implementation into your project manually and replaced the atomic update intrinsic with an OpenMP atomic operation. That being said, I'm not confident you would get good performance using a shared pointer on the GPU, the calls to malloc / free and atomic updates are going to be very expensive.

[VedantParanjape]

@jhuber6 @ikitayama what is the status of this issue? I would like to take stab at it.

[jhuber6]

This works with my current support for libc++ on the GPU. However it will be very slow because none of this is optimized yet, particularly the calls to malloc.

#include <memory>

int foo() {
  auto ptr = std::make_shared<int>();
  *ptr = 42;
  return *ptr;
}

int main() {
  int ret = 42;
#pragma omp target teams map(from : ret)
#pragma omp parallel
  ret = foo();
  return ret;
}

clang++ shared_ptr.cpp -fopenmp --offload-arch=gfx1030 -stdlib=libc++ -Xoffload-linker -lc -Xoffload-linker -lc++ -O2

[VedantParanjape]

@jhuber6 Thanks! So, I could work on optimizing this or are there any pending tasks that I could look at ? Which involve improving features in openMP offload for CUDA?

[jhuber6]

Thanks! So, I could work on optimizing this or are there any pending tasks that I could look at ? which involve improving features in openMP offload for CUDA?

I've been planning on writing a real malloc in the GPU libc. I don't have anything specific in mind but I guess you could play around with the C++ build. I still haven't done any GPU specific work, it's just straight port of the C++ code to the GPU, so it's not exactly fast. https://github.com/llvm/llvm-project/blob/main/offload/cmake/caches/Offload.cmake should show the CMake required to get the above demo working.

[VedantParanjape]

Thanks! So, I could work on optimizing this or are there any pending tasks that I could look at ? which involve improving features in openMP offload for CUDA?

I've been planning on writing a real malloc in the GPU libc. I don't have anything specific in mind but I guess you could play around with the C++ build. I still haven't done any GPU specific work, it's just straight port of the C++ code to the GPU, so it's not exactly fast. https://github.com/llvm/llvm-project/blob/main/offload/cmake/caches/Offload.cmake should show the CMake required to get the above demo working.

Ah, cool. implementing malloc would be a real fun project to do!