Runtime error when copying to constant device memory with HIP

[gzagaris]

Describe the bug

Copying data from the host to constant device memory using HIP throws a runtime error. The code works fine with CUDA , but, Umpire throws the following runtime error when compiled with HIP:

C++ exception with description "! Umpire runtime_error [.../Umpire/src/umpire/op/HipCopyOperation.cpp:31]: hipMemcpy( dest_ptr = 0xe3e480, src_ptr = 0x52fb90, length = 4096) failed with error: invalid argument

To Reproduce

Here is a code snippet that reproduces this behavior:

auto& rm = umpire::ResourceManager::getInstance();
auto const_allocator = rm.getAllocator("DEVICE_CONST");

static constexpr int N = 4;
static constexpr int BYTESIZE = N * sizeof(int);
static constexpr int TEST_VAL = 42;

auto host_allocator = rm.getAllocator("HOST");
int* HOST_DATA = static_cast< int* >(host_allocator.allocate(BYTESIZE));
for ( int i = 0; i < N; ++i ) {
    HOST_DATA[ i ] = TEST_VAL;
}

int* A_d = static_cast< int* >(const_allocator.allocate(BYTESIZE));
EXPECT_TRUE(A_d != nullptr);
rm.copy(A_d, HOST_DATA, BYTESIZE); // <------------- RUNTIME ERROR THROWN HERE!

I am compiling Umpire with -DENABLE_HIP=On -DUMPIRE_ENABLE_DEVICE_CONST=On.

Am I missing anything?

Expected behavior

I would have expected this to work and not throw a runtime error.

Compilers & Libraries (please complete the following information):

• Compiler & version: amdclang-16.0.0
• ROCM version: v5.6.0