[PATCH] D117118: [NVPTX] Fix shr/and pair replace with bfe

[Georgi Mirazchiyski via Phabricator via llvm-commits]

gmirazchiyski added a comment.

I can simplify this but I don't feel it's ideal to strip the checks to a plain `CHECK-NOT: bfe` only, because the test is intended to test that `bfe` is not replacing `shr/and` pair due to the overflow issue.
For that reason I've split the test in two identical labels but one only checking for no `bfe` and the other for the presence of the `shr/and` pair.

So the following version I think should be good.

  ; CHECK-LABEL: no_bfe_on_64bit_overflow
  define i64 @no_bfe_on_64bit_overflow(i64 %a) {
  ; CHECK-NOT: bfe.u64 %rd{{[0-9]+}}, %rd{{[0-9]+}}, 63, 3
    %val0 = ashr i64 %a, 63
    %val1 = and i64 %val0, 7
    ret i64 %val1
  }
  
  ; CHECK-LABEL: no_bfe_on_64bit_overflow_shr_and_pair
  define i64 @no_bfe_on_64bit_overflow_shr_and_pair(i64 %a) {
  ; CHECK: shr.s64 %rd{{[0-9]+}}, %rd{{[0-9]+}}, 63
  ; CHECK: and.b64 %rd{{[0-9]+}}, %rd{{[0-9]+}}, 7
    %val0 = ashr i64 %a, 63
    %val1 = and i64 %val0, 7
    ret i64 %val1
  }
  
  While not the prettiest, it will at least test the complete thing as we intended.
  
  With the applied changes, the assembly for the 64-bit test look like:

And the assembly looks like follows:

  	// .globl	no_bfe_on_64bit_overflow // -- Begin function no_bfe_on_64bit_overflow
  .visible .func  (.param .b64 func_retval0) no_bfe_on_64bit_overflow(
  	.param .b64 no_bfe_on_64bit_overflow_param_0
  )                                       // @no_bfe_on_64bit_overflow
  {
  	.reg .b64 	%rd<4>;
  
  // %bb.0:
  	ld.param.u64 	%rd1, [no_bfe_on_64bit_overflow_param_0];
  	shr.s64 	%rd2, %rd1, 63;
  	and.b64  	%rd3, %rd2, 7;
  	st.param.b64 	[func_retval0+0], %rd3;
  	ret;
                                          // -- End function
  }

So everything will be tested correctly and if it ever regresses again will report nicely but at the cost of duplicating the labels in order for us to be able to run the `CHECK-NOT` exclusively.

@tra Thanks for your feedback! What do you think about this approach instead of just simplifying the test?



================
Comment at: llvm/test/CodeGen/NVPTX/bfe.ll:34
+
+; CHECK: bfe3
+define i32 @bfe3(i32 %a) {
----------------
tra wrote:
> `CHECK-LABEL:`
> 
> I'd also give it a more meaningful name. `no_bfe_on_32bit_oveflow`.
Sure! Looks more clear that way.


================
Comment at: llvm/test/CodeGen/NVPTX/bfe.ll:36
+define i32 @bfe3(i32 %a) {
+; CHECK-NOT: bfe %r{{[0-9]+}}, %r{{[0-9]+}}, 31, 4
+; CHECK: shr
----------------
tra wrote:
> This would never match as the actual instruction would be `bfe.u32`
Yeah, thanks! This needed changed to `bfe.u32` and the check for the 64-bit test-case changed to `bfe.u64` to match.


================
Comment at: llvm/test/CodeGen/NVPTX/bfe.ll:46
+define i64 @bfe4(i64 %a) {
+; CHECK-NOT: bfe %r{{[0-9]+}}, %r{{[0-9]+}}, 63, 3
+; CHECK: shr
----------------
tra wrote:
> This will also never match.  The instruction will be `bfe.u64` and the registers are 64-bit `%r*d*`.
> 
> Please run the tests with unmodified llvm and make sure it does fail when it detects bfe. 
> 
> As it happens, it will still fail when it would fail to find `shr`, but it would still be great to correctly test for what we intended to test.
> As an alternative, you could simplify it all to `CHECK-NOT: bfe`
> This will also never match. The instruction will be bfe.u64 and the registers are 64-bit %r*d*.
Good catch!
Yes, the expression also wasn't correct as you pointed out. I checked exactly what to match the produced `bfe` to and it has to be as follows:
```
; CHECK-NOT: bfe.u64 %rd{{[0-9]+}}, %rd{{[0-9]+}}, 63, 3
```

> Please run the tests with unmodified llvm and make sure it does fail when it detects bfe.
I have run those and it does fail when `bfe` is detected. Without the llvm changes, `bfe.u64` is produced and unfortunately failed by reporting it's not being able find `shr` rather than because the `bfe` instruction was present.

That is not great because we're not communicating the issue too well in case this ever regresses in the future. This is because `CHECK`s seem to take precedence over `CHECK-NOT`s.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D117118/new/

https://reviews.llvm.org/D117118