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 ↗

(On Diff #65984)

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

74 ↗

(On Diff #65984)

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

75–76 ↗

(On Diff #65984)

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

test/CodeGenOpenCL/builtins-amdgcn.cl

207 ↗

(On Diff #65984)

Non-constant 3rd parameter is invalid

[wdng]

include/clang/Basic/BuiltinsAMDGPU.def

73–74 ↗

(On Diff #65984)

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 ↗

(On Diff #65984)

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

test/CodeGenOpenCL/builtins-amdgcn.cl

7 ↗

(On Diff #65984)

This should be removed

[wdng]

test/CodeGenOpenCL/builtins-amdgcn.cl

207 ↗

(On Diff #65984)

should the third constant be "const int c" ?

[arsenm]

test/CodeGenOpenCL/builtins-amdgcn.cl

207 ↗

(On Diff #65984)

Yes

207 ↗

(On Diff #65984)

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 ↗

(On Diff #66044)

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

test/CodeGenOpenCL/builtins-amdgcn.cl

210 ↗

(On Diff #66044)

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 ↗

(On Diff #66044)

The output type should be ulong

214 ↗

(On Diff #66044)

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

228 ↗

(On Diff #66044)

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 ↗

(On Diff #66134)

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 ↗

(On Diff #66143)

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.

cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def

 BUILTIN(__builtin_amdgcn_class, "bdi", "nc")
 BUILTIN(__builtin_amdgcn_classf, "bfi", "nc")
 BUILTIN(__builtin_amdgcn_cubeid, "ffff", "nc")
 BUILTIN(__builtin_amdgcn_cubesc, "ffff", "nc")
 BUILTIN(__builtin_amdgcn_cubetc, "ffff", "nc")
 BUILTIN(__builtin_amdgcn_cubema, "ffff", "nc")
 BUILTIN(__builtin_amdgcn_s_memtime, "LUi", "n")
 BUILTIN(__builtin_amdgcn_s_sleep, "vIi", "n")
+BUILTIN(__builtin_amdgcn_uicmp, "LUiUiUiIi", "nc")
+BUILTIN(__builtin_amdgcn_uicmpl, "LUiLUiLUiIi", "nc")
+BUILTIN(__builtin_amdgcn_sicmp, "LUiiiIi", "nc")
+BUILTIN(__builtin_amdgcn_sicmpl, "LUiLiLiIi", "nc")
+BUILTIN(__builtin_amdgcn_fcmp, "LUiddIi", "nc")
+BUILTIN(__builtin_amdgcn_fcmpf, "LUiffIi", "nc")
 
 //===----------------------------------------------------------------------===//
 // VI+ only builtins.
 //===----------------------------------------------------------------------===//
 
 TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "LUi", "n", "s-memrealtime")
 
 //===----------------------------------------------------------------------===//

cfe/trunk/lib/CodeGen/CGBuiltin.cpp

Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
   case AMDGPU::BI__builtin_amdgcn_frexp_expf: {
     return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_frexp_exp);
   }
   case AMDGPU::BI__builtin_amdgcn_fract:
   case AMDGPU::BI__builtin_amdgcn_fractf:
     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);
 
   case AMDGPU::BI__builtin_amdgcn_read_exec: {
     CallInst *CI = cast<CallInst>(
       EmitSpecialRegisterBuiltin(*this, E, Int64Ty, Int64Ty, true, "exec"));
     CI->setConvergent();

cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl

 // REQUIRES: amdgpu-registered-target
 // RUN: %clang_cc1 -triple amdgcn-unknown-amdhsa -target-cpu tahiti -verify -S -o - %s
 
 // FIXME: We only get one error if the functions are the other order in the
 // file.
 
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
 typedef unsigned long ulong;
+typedef unsigned int uint;
 
 ulong test_s_memrealtime()
 {
   return __builtin_amdgcn_s_memrealtime(); // expected-error {{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}}
 }
 
 void test_s_sleep(int x)
 {
   __builtin_amdgcn_s_sleep(x); // expected-error {{argument to '__builtin_amdgcn_s_sleep' must be a constant integer}}
 }
 
+void test_sicmp_i32(global ulong* out, int a, int b, uint c)
+{
+  *out = __builtin_amdgcn_sicmp(a, b, c); // expected-error {{argument to '__builtin_amdgcn_sicmp' must be a constant integer}}
+}
+
+void test_uicmp_i32(global ulong* out, uint a, uint b, uint c)
+{
+  *out = __builtin_amdgcn_uicmp(a, b, c); // expected-error {{argument to '__builtin_amdgcn_uicmp' must be a constant integer}}
+}
+
+void test_sicmp_i64(global ulong* out, long a, long b, uint c)
+{
+  *out = __builtin_amdgcn_sicmpl(a, b, c); // expected-error {{argument to '__builtin_amdgcn_sicmpl' must be a constant integer}}
+}
+
+void test_uicmp_i64(global ulong* out, ulong a, ulong b, uint c)
+{
+  *out = __builtin_amdgcn_uicmpl(a, b, c); // expected-error {{argument to '__builtin_amdgcn_uicmpl' must be a constant integer}}
+}
+
+void test_fcmp_f32(global ulong* out, float a, float b, uint c)
+{
+  *out = __builtin_amdgcn_fcmpf(a, b, c); // expected-error {{argument to '__builtin_amdgcn_fcmpf' must be a constant integer}}
+}
+
+void test_fcmp_f64(global ulong* out, double a, double b, uint c)
+{
+  *out = __builtin_amdgcn_fcmp(a, b, c); // expected-error {{argument to '__builtin_amdgcn_fcmp' must be a constant integer}}
+}
+

cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl

 // REQUIRES: amdgpu-registered-target
 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s
 
 #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)
 // CHECK-DAG: [[FLAG:%.+]] = extractvalue { double, i1 } %{{.+}}, 1
 // CHECK-DAG: [[VAL:%.+]] = extractvalue { double, i1 } %{{.+}}, 0
 // CHECK: [[FLAGEXT:%.+]] = zext i1 [[FLAG]] to i32
 // CHECK: store i32 [[FLAGEXT]]
 void test_div_scale_f64(global double* out, global int* flagout, double a, double b)

 
 // CHECK-LABEL: @test_lerp
 // CHECK: call i32 @llvm.amdgcn.lerp
 void test_lerp(global int* out, int a, int b, int c)
 {
   *out = __builtin_amdgcn_lerp(a, b, c);
 }
 
+// CHECK-LABEL: @test_sicmp_i32
+// CHECK: call i64 @llvm.amdgcn.icmp.i32(i32 %a, i32 %b, i32 32)
+void test_sicmp_i32(global ulong* out, int a, int b)
+{
+  *out = __builtin_amdgcn_sicmp(a, b, 32);
+}
+
+// CHECK-LABEL: @test_uicmp_i32
+// CHECK: call i64 @llvm.amdgcn.icmp.i32(i32 %a, i32 %b, i32 32)
+void test_uicmp_i32(global ulong* out, uint a, uint b)
+{
+  *out = __builtin_amdgcn_uicmp(a, b, 32);
+}
+
+// CHECK-LABEL: @test_sicmp_i64
+// CHECK: call i64 @llvm.amdgcn.icmp.i64(i64 %a, i64 %b, i32 38)
+void test_sicmp_i64(global ulong* out, long a, long b)
+{
+  *out = __builtin_amdgcn_sicmpl(a, b, 39-1);
+}
+
+// CHECK-LABEL: @test_uicmp_i64
+// CHECK: call i64 @llvm.amdgcn.icmp.i64(i64 %a, i64 %b, i32 35)
+void test_uicmp_i64(global ulong* out, ulong a, ulong b)
+{
+  *out = __builtin_amdgcn_uicmpl(a, b, 30+5);
+}
+
+// CHECK-LABEL: @test_fcmp_f32
+// CHECK: call i64 @llvm.amdgcn.fcmp.f32(float %a, float %b, i32 5)
+void test_fcmp_f32(global ulong* out, float a, float b)
+{
+  *out = __builtin_amdgcn_fcmpf(a, b, 5);
+}
+
+// CHECK-LABEL: @test_fcmp_f64
+// CHECK: call i64 @llvm.amdgcn.fcmp.f64(double %a, double %b, i32 6)
+void test_fcmp_f64(global ulong* 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)
 {
   *out = __builtin_amdgcn_classf(a, b);
 }
 
 // CHECK-LABEL: @test_class_f64