add mappings #1403
add mappings #1403
Conversation
add std and cub::BLOCK_LOAD_WARP_TRANSPOSE Signed-off-by: Song <xiaosong@amd.com>
| ("std::fabs", ("std::fabs")), | ||
| ("std::fmod", ("std::fmod")), | ||
| ("std::remainder", ("std::remainder")), | ||
| ("std::frexp", ("std::frexp")), |
There was a problem hiding this comment.
@jeffdaily I think we specifically chose to map the std:: functions to anonymous namespace due to some device libs issue? Do you recall?
There was a problem hiding this comment.
For these math operations we do not want to use the std implementations of them but instead the implementations provided by HIP. That's what I recall.
There was a problem hiding this comment.
Are we facing problems by not mapping them to std:: namespace?
There was a problem hiding this comment.
@jeffdaily Have tested on the rocm6.0 with nightly pytorch build and rocm6.2 provided by @jithunnair-amd, the hip can't handle it properly. got some error like this "/root/AISS-Data-Science/mamba/causal-conv1d/csrc/hip/causal_conv1d_fwd.hip:36:11: error: no matching function for call to 'max'
: ::max({sizeof(typename BlockLoadT::TempStorage), sizeof(typename BlockStoreT::TempStorage)});"
There was a problem hiding this comment.
It can be reproduced by this ticket https://ontrack-internal.amd.com/browse/SWDEV-458189
There was a problem hiding this comment.
I was working on a similar issue recently. It appears that these std mappings are handled differently depending on whether we are talking about device vs host functions. I am not sure how it will behave in cases when a function is defined as both, as it is in https://github.com/state-spaces/mamba/blob/main/csrc/selective_scan/uninitialized_copy.cuh, for example.
It definitely fails to map properly in its present state, but I am not sure that simply adding the mapping won't break something else.
Maybe the parser gets confused for such functions?
There was a problem hiding this comment.
TLDR: I suspect it might be due to a function being defined as both host and device, while these mappings are supposed to be replaced in device function only. Maybe the hipifyer does not handle the "both" case properly?
add std and cub::BLOCK_LOAD_WARP_TRANSPOSE
Jira ID: SWDEV-458189
Fixes #ISSUE_NUMBER