add mappings #1403
add mappings #1403
Conversation
add std and cub::BLOCK_LOAD_WARP_TRANSPOSE Signed-off-by: Song <[email protected]>
| ("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?
|
Jenkins build for c63bab0cdb86c8560c5cd01889922dbff26622a5 commit finished as FAILURE |
|
Jenkins build for c63bab0cdb86c8560c5cd01889922dbff26622a5 commit finished as FAILURE |
|
Jenkins build for c63bab0cdb86c8560c5cd01889922dbff26622a5 commit finished as FAILURE |
|
Jenkins build for c63bab0cdb86c8560c5cd01889922dbff26622a5 commit finished as FAILURE Detected error during base docker image building: |
add std and cub::BLOCK_LOAD_WARP_TRANSPOSE
Jira ID: SWDEV-458189
Fixes #ISSUE_NUMBER