speaking for myself, even if we posit migration, I would not suppose that migrating data on first touch (best imaginable case for malloc()) is as fast as non-migrated access via cudaMalloc() in every respect. It’s not plausible to me. I do agree that once the page is migrated, access to data there is equivalent performance-wise to any other HBM access.
In a nutshell, ATS is a hardware mechanism for page mapping and translation. HMM is a software mechanism to do the same thing. Grace has ATS built into the hardware (related to the NVLink-C2C bus), so it can take advantage of that. For an x86_64 system, HMM would be the only option. I don’t know of performance comparison between the two cases, but my general expectation is that ATS is faster (and may have more utility) than HMM.
If my comments seem doubtful to you, another approach would be to use a profiler to understand the difference between the two cases. Since the only change you indicated you have made is to convert cudaMalloc to malloc, presumably that would be a clue to investigate; presumably the slowdown relates to using those allocations, somehow.