AMDGPU : Add Clang builtin intrinsics for compare with the full wavefront result

[wdng]

Summary

AMDGPU : Add Clang builtin intrinsics for compare with the full wavefront result.

[arsenm]

Tests should be added for non-constant compare type operands being rejected in test/CodeGenOpenCL/builtins-amdgcn-error.cl

include/clang/Basic/BuiltinsAMDGPU.def

73–74

There should be 4 intrinsics defined for icmp. We need uicmp, uicmpl, sicmp, sicmpl. I'm not sure what W is.

74

The bultin suffix for long should be l. The other builtins also use L instead of W for int64

75–76

All of these intrinsics should have the final parameter marked with I to require it be an immediate

test/CodeGenOpenCL/builtins-amdgcn.cl

207

Non-constant 3rd parameter is invalid

[wdng]

include/clang/Basic/BuiltinsAMDGPU.def

73–74

According to the definition on https://code.woboq.org/llvm/clang/include/clang/Basic/Builtins.def.html W is for int64_t

[arsenm]

include/clang/Basic/BuiltinsAMDGPU.def

73–74

L is for long which is also 64-bit, we don't need to use W

test/CodeGenOpenCL/builtins-amdgcn.cl

7

This should be removed

[wdng]

test/CodeGenOpenCL/builtins-amdgcn.cl

207

should the third constant be "const int c" ?

[arsenm]

test/CodeGenOpenCL/builtins-amdgcn.cl

207

Yes

207

Maybe unsigned

[wdng]

1. Modified intrinsic parameters' data type
2. Modified the LIT test functions' last parameter to const uint.

[wdng]

Fixed ref assembly errors.

[arsenm]

include/clang/Basic/BuiltinsAMDGPU.def

73–78

The return types look wrong. They should all be ulong, not signed

test/CodeGenOpenCL/builtins-amdgcn.cl

210

The tests seem to be in a random order. Can you sort them so that the same operation on the same type is the next one? This is missing uicmp_i32

212

The output type should be ulong

214

This isn't a valid icmp predicate type. Ideally there would be an error for this but that can be another patch

228

This test is using an unsigned predicate value for the signed version

[wdng]

1. Change return data type to unsigned long.
2. Fix incorrect use of cond code.

[arsenm]

include/clang/Basic/BuiltinsAMDGPU.def

75–76

This looks inconsistent. You don't need the 'S'es

[wdng]

Removed 'S'es.

[arsenm]

Still missing test that non-constants for 3rd operand are rejected

test/CodeGenOpenCL/builtins-amdgcn.cl

233

Unused c parameter

[wdng]

In D22934#501396, @arsenm wrote:

Still missing test that non-constants for 3rd operand are rejected

If I write a test with the third parameter as an non-constant, it fails at the compilation phase:

llvm/tools/clang/test/CodeGenOpenCL/builtins-amdgcn.cl:249:10: error: argument to '__builtin_amdgcn_fcmp' must be a constant integer

*out = __builtin_amdgcn_fcmp(a, b, c);
       ^                           ~

1 error generated.
FileCheck error: '-' is empty.

So how to write a FileCheck to detect this? Thanks!

[arsenm]

In D22934#502183, @wdng wrote:

In D22934#501396, @arsenm wrote:

Still missing test that non-constants for 3rd operand are rejected

If I write a test with the third parameter as an non-constant, it fails at the compilation phase:

llvm/tools/clang/test/CodeGenOpenCL/builtins-amdgcn.cl:249:10: error: argument to '__builtin_amdgcn_fcmp' must be a constant integer

*out = __builtin_amdgcn_fcmp(a, b, c);
       ^                           ~

1 error generated.
FileCheck error: '-' is empty.

So how to write a FileCheck to detect this? Thanks!

Yes, that's the point. The tests in builtins-amdgcn-error.cl do this, you should add these tests there

[wdng]

A failure test is added.

[arsenm]

test/CodeGenOpenCL/builtins-amdgcn-error.cl

20–23 ↗

(On Diff #66484)

All of the new builtins should be tested here

[wdng]

Add all of the new builtins fail tests.

[arsenm]

LGTM with the last test cleanup

test/CodeGenOpenCL/builtins-amdgcn-error.cl

21 ↗

(On Diff #66703)

It doesn't really matter, but the parameter types here should be int.

31 ↗

(On Diff #66703)

Same, should be long

[jdoerfert]

Wrong review/tab. Apologies.

include/clang/Basic/BuiltinsAMDGPU.def

 BUILTIN(__builtin_amdgcn_cubema, "ffff", "nc")
 BUILTIN(__builtin_amdgcn_s_memtime, "LUi", "n")
 BUILTIN(__builtin_amdgcn_s_sleep, "vIi", "n")
+BUILTIN(__builtin_amdgcn_uicmp, "SiUiUiIi", "nc")
+BUILTIN(__builtin_amdgcn_uicmpl, "LiLUiLUiIi", "nc")
+BUILTIN(__builtin_amdgcn_sicmp, "SiSiSiIi", "nc")
+BUILTIN(__builtin_amdgcn_sicmpl, "LiLiLiIi", "nc")
+BUILTIN(__builtin_amdgcn_fcmp, "LiddIi", "nc")
+BUILTIN(__builtin_amdgcn_fcmpf, "LiffIi", "nc")
 
 //===----------------------------------------------------------------------===//
 // VI+ only builtins.

lib/CodeGen/CGBuiltin.cpp

     return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_fract);
   case AMDGPU::BI__builtin_amdgcn_lerp:
     return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_lerp);
+  case AMDGPU::BI__builtin_amdgcn_uicmp:
+  case AMDGPU::BI__builtin_amdgcn_uicmpl:
+  case AMDGPU::BI__builtin_amdgcn_sicmp:
+  case AMDGPU::BI__builtin_amdgcn_sicmpl:
+    return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_icmp);
+  case AMDGPU::BI__builtin_amdgcn_fcmp:
+  case AMDGPU::BI__builtin_amdgcn_fcmpf:
+    return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_fcmp);
   case AMDGPU::BI__builtin_amdgcn_class:
   case AMDGPU::BI__builtin_amdgcn_classf:
     return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_class);

test/CodeGenOpenCL/builtins-amdgcn.cl

 #pragma OPENCL EXTENSION cl_khr_fp64 : enable
 
 typedef unsigned long ulong;
+typedef unsigned int uint;
 
 // CHECK-LABEL: @test_div_scale_f64
 // CHECK: call { double, i1 } @llvm.amdgcn.div.scale.f64(double %a, double %b, i1 true)

   *out = __builtin_amdgcn_lerp(a, b, c);
 }
 
+// CHECK-LABEL: @test_uicmp_i32
+// CHECK: call i64 @llvm.amdgcn.icmp.i32(global long* out, uint a, uint b, i32 32)
+void test_uicmp_i32(global long* out, uint a, uint b)
+{
+  *out = __builtin_amdgcn_uicmp(a, b, 32);
+}
+
+// CHECK-LABEL: @test_uicmp_i64
+// CHECK: call i64 @llvm.amdgcn.icmp.i32(global long* out, ulong a, ulong b, i32 4+5)
+void test_uicmp_i64(global long* out, ulong a, ulong b)
+{
+  *out = __builtin_amdgcn_uicmpl(a, b, 4+5);
+}
+
+// CHECK-LABEL: @test_sicmp_i32
+// CHECK: call i64 @llvm.amdgcn.icmp.i32(global long* out, uint a, uint b, i32 32)
+void test_sicmp_i32(global long* out, uint a, uint b)
+{
+  *out = __builtin_amdgcn_sicmp(a, b, 32);
+}
+
+// CHECK-LABEL: @test_sicmp_i64
+// CHECK: call i64 @llvm.amdgcn.icmp.i32(global long* out, ulong a, ulong b, i32 33+1)
+void test_sicmp_i64(global long* out, ulong a, ulong b)
+{
+  *out = __builtin_amdgcn_sicmpl(a, b, 33+1);
+}
+
+// CHECK-LABEL: @test_fcmp_f32
+// CHECK: call i64 @llvm.amdgcn.fcmp.f32(global long* out, float a, float b, i32 5)
+void test_fcmp_f32(global long* out, float a, float b, const uint c)
+{
+  *out = __builtin_amdgcn_fcmpf(a, b, 5);
+}
+
+// CHECK-LABEL: @test_fcmp_f64
+// CHECK: call i64 @llvm.amdgcn.fcmp.f64(global long* out, double a, double b, i32 3+3)
+void test_fcmp_f64(global long* out, double a, double b)
+{
+  *out = __builtin_amdgcn_fcmp(a, b, 3+3);
+}
+
 // CHECK-LABEL: @test_class_f32
 // CHECK: call i1 @llvm.amdgcn.class.f32
 void test_class_f32(global float* out, float a, int b)