Optimise the gridDim.n * blockDim.m idiom

[AlexVlx]

No description provided.

[whchung]

+@asroy +@deven-amd for awareness.

Although this PR doesn't solve the fundamental issue of hipGridDim_x/y/z but it does address a common pattern seen among all Eigen kernels on ROCm.

[AlexVlx]

Not quite sure what the fundamental problem is if I'm honest. If it's the integral division required to get the number of groups, we can probably address that one fairly easily too, although I suspect that this idiom is far more commonly encountered.

[whchung]

Yes, the fundamental issue is the integer division used in the implementation of __ockl_get_num_groups, and also the apparent omission of enable_sgpr_grid _workgroup_count_X/Y/Z

https://llvm.org/docs/AMDGPUUsage.html#initial-kernel-execution-state

Indeed with this PR merged it should already address quite a lot of inefficient codes in many HIP kernels as the idiom is used literally everywhere.

[b-sumner]

Is it really the fundamental problem? In a separate context, I thought you said that the performance impact is minimal.

[AlexVlx]

Yeah I don't know how fundamental that one is, but "fixing" it should be straightforward. Let us discuss separately.

[whchung]

@b-sumner the performance impact was small, for the kernels I was studying at that moment. So I wasn't pushing for a change. And this PR should nicely optimize away issues found in those kernels.

But there are other kernels from other workloads which directly access hipGridDim_x/y/z without following this idiom in performance-critical loops. Again it's hard to quantify the potential impact yet without truly dive down into the ISA for those kernels. So I'm not requesting having a change either, yet. We might revisit this issue again in a couple of weeks when we have better understanding on these kernels.

[b-sumner]

@whchung I agree it is a problem if the division is not lifted out of a performance critical loop. But the fundamental issue then is not the specifics of the computation, but rather that an invariant computation is not being identified and lifted out of that loop.

[whchung]

@b-sumner perhaps we are interpreting the term "fundamental issue" differently.

To me the "fundamental issue" is that reading hipGridDim_x/y/z involves integer division and makes derivative values reside in VGPR, not in SGPR. On the other hand, all other indexing values in HIP/ROCm-Device-Libs are nicely stored in SGPRs.

Should v_readlane be used in ROCm-Device-Libs implementation, to put hipGridDim_x/y/z values back into SGPR we may be able to at least keep most derivative values in SGPR, which may reduce some instructions and VGPR register usage.

---

As a side note, we've been starting to upstream our works on MLIR and the very first step is to enable integrating ROCm-Device-Libs into MLIR, so functions written in GPU dialect could be properly indexed on ROCm with ROCm-Device-Libs.
tensorflow/mlir#146

[b-sumner]

@whchung I think you're observing side effects of the base or "fundamental" problem which has not been identified. Using readfirstlane in the device libs would simply be a workaround that could affect other optimizations and I would oppose it until found necessary after the fundamental problem has been identified.

[whchung]

@b-sumner I agree we postpone this discussion until further understanding the nature of these new kernels which don't follow idiom fixed in this PR.

[bensander]

LGTM