CHIP-SPV / chipStar

chipStar is a tool for compiling and running HIP/CUDA on SPIR-V via OpenCL or Level Zero APIs.
Other
227 stars 34 forks source link

Can we really support "__managed__" feature ? #955

Open joyongzhu opened 2 weeks ago

joyongzhu commented 2 weeks ago
#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 commented 2 weeks ago

I found a history issue:#624

pvelesko commented 2 weeks ago

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