Can we really support "__managed__" feature ？

[joyongzhu]

#include <iostream>
#include <hip/hip_runtime.h>

__managed__ int data[10];

__global__ void kernel() {
    int tid = threadIdx.x;
    data[tid] *= 2; 
}

int main() {
    for (int i = 0; i < 10; i++) {
        data[i] = i;
    }

    hipLaunchKernelGGL(kernel, dim3(1), dim3(10), 0, 0);
    hipDeviceSynchronize(); 

    for (int i = 0; i < 10; i++) {
        std::cout << data[i] << " ";
    }
    std::cout << std::endl;

    return 0;
}

I ran this code but got an incorrect result, so I checked the Chipstar code and found one line:

//spirv_hip_runtime.h
#define __managed__ __device__

So we can't really support it ?
Can we use some other way to fix it？such as opencl api clSVMAlloc.

[joyongzhu]

I found a history issue:#624

[pvelesko]

Implementing this feature would require writing an additional LLVM pass that would insert allocations. This is doable just a question of priorities.

Can you share what code you are working on that requires this? @joyongzhu

[joyongzhu]

hi bro，code is on the top。normally，I can set a value for a managed variable in host code，just like writing a normal cpp code。

And，on chipstar，if I want to get a managed variable，I must use hipMemcpyFromSymbol，but on AMD，I can use it directly。

[pvelesko]

Currently we're prioritizing features by the application needs as we are trying to enable as many applications as possible. For example, currently we're working on enabling OpenMM and fixing bugs/features associated with that application.