[llvm] [NVPTX] Fix 64 bits rotations with large shift values (PR #89399)

[Nicolas Miller via llvm-commits]

================
@@ -1747,8 +1747,9 @@ def ROTL64reg_sw :
             ".reg .b64 %lhs;\n\t"
             ".reg .b64 %rhs;\n\t"
             ".reg .u32 %amt2;\n\t"
-            "shl.b64 \t%lhs, $src, $amt;\n\t"
-            "sub.u32 \t%amt2, 64, $amt;\n\t"
+            "and.b32 \t%amt2, $amt, 63;\n\t"
----------------
npmiller wrote:

That might make sense yeah.

I was also thinking of trying to lift the `reg_sw` implementations up to DAG level, they're just generic rotation implementations, we only need specific lowering for when we can use `shf`, but I thought I'd fix the correctness issue first.

https://github.com/llvm/llvm-project/pull/89399