search

ROCm / hcc

HCC is an Open Source, Optimizing C++ Compiler for Heterogeneous Compute currently for the ROCm GPU Computing Platform
https://github.com/RadeonOpenCompute/hcc/wiki
Other
433 stars 108 forks source link

hipcc does not hipify __nvvm_get_smem_pointer function #1457

Open terU3760 opened 2 years ago

terU3760 commented 2 years ago

In many CUDA related project, we can see the line of code as the following:

extern "C" __device__ uint32_t __nvvm_get_smem_pointer(void *ptr);

It is used to convert the shared memory address into an 32 bit int address. But when I tried to use hipcc to build the project containing such line, it just doesn't hipify this line i.e. compile OK but link error:

lld: error: undefined hidden symbol: __nvvm_get_smem_pointer

So what is the equivalent on AMD gpu platform of the function __nvvm_get_smem_pointer in Nvidia CUDA library?

searlmc1 commented 2 years ago
b-sumner commented 2 years ago

This looks like an internal function to me, well outside the scope of the cuda programming language. Why are those "many Cuda related project" choosing such a risky direction? What is wrong with just using ordinary pointers?

terU3760 commented 2 years ago

This looks like an internal function to me, well outside the scope of the cuda programming language. Why are those "many Cuda related project" choosing such a risky direction? What is wrong with just using ordinary pointers?

@b-sumner "many Cuda related project", Yes, there a few of them, such as cutlass, apex, etc. A lot of applications using shared memory mechanism of CUDA may use it. So by your answer, I guess you mean, in AMD gpu platform, just replace this __nvvm_get_smem_pointer function with an simple type conversion from void* to uint32_t is enough? Or there are more things needed to be done other than just a simple type conversion?