Device side new/delete or malloc/free incorrect?

[jiajunc516]

I'm porting my CUDA code to HIP platform. I'm doing device side new/delete in my kernel. It works fine on CUDA platform but an error occurs on HIP platform. Here is my code. "main.cpp"

#include <hip/hip_runtime.h>

struct myST {
    size_t a;
    size_t b;
    size_t c;
};

__global__ void my_kernel() {
    printf("sizeof(myST): %zu\n", sizeof(myST));

    myST* A = new myST();
    printf("[0]    p: %p\n", A);
    
    myST* B = new myST[2];
    printf("------------------------------------------------\n");
    for (int i = 0; i < 2; ++i) {
        printf("[%d]    p: %p\n", i, B + i);
    }

    delete A;

    char** C = new char*[10];
    printf("------------------------------------------------\n");
    for (int i = 0; i < 10; ++i) {
        printf("[%d]    p: %p\n", i, C + i);
    }
}

int main()
{
    hipLaunchKernelGGL(my_kernel, dim3(1), dim3(1), 0, 0);
    hipDeviceSynchronize();

    return 0;
}

As compiled with

hipcc -std=c++11 -D__HIP_ENABLE_DEVICE_MALLOC__ -D__HIP_SIZE_OF_PAGE=8  main.cpp -o main

The output is

sizeof(myST): 24
[0]    p: 0x7fd50b00f000
------------------------------------------------
[0]    p: 0x7fd50b00f018
[1]    p: 0x7fd50b00f030
------------------------------------------------
[0]    p: 0x7fd50b00f000
[1]    p: 0x7fd50b00f008
[2]    p: 0x7fd50b00f010
[3]    p: 0x7fd50b00f018 // the same location as B
[4]    p: 0x7fd50b00f020
[5]    p: 0x7fd50b00f028
[6]    p: 0x7fd50b00f030
[7]    p: 0x7fd50b00f038
[8]    p: 0x7fd50b00f040
[9]    p: 0x7fd50b00f048

As you can see, C, the array of char*, starting at the location of deleted object A, is overlapped with the memory location of B. I tried malloc/free instead and nothing changed. Is this a bug or am I doing something wrong? ps: This is on rocm-4.3.0.

[b-sumner]

HIP does not currently support malloc/free or non-placement new/delete.