Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

hipcc does not hipify __nvvm_get_smem_pointer function #1457

Open
terU3760 opened this issue Apr 24, 2022 · 3 comments
Open

hipcc does not hipify __nvvm_get_smem_pointer function #1457

terU3760 opened this issue Apr 24, 2022 · 3 comments

Comments

@terU3760
Copy link

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
Copy link

@b-sumner
Copy link

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
Copy link
Author

terU3760 commented Apr 25, 2022

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?

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

No branches or pull requests

3 participants