hcc icon indicating copy to clipboard operation
hcc copied to clipboard

hipcc does not hipify __nvvm_get_smem_pointer function

Open terU3760 opened this issue 2 years ago • 3 comments

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?

terU3760 avatar Apr 24 '22 16:04 terU3760

  • @b-sumner , @yxsamliu , @emankov

searlmc1 avatar Apr 24 '22 21:04 searlmc1

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 avatar Apr 24 '22 23:04 b-sumner

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?

terU3760 avatar Apr 25 '22 15:04 terU3760