diff options
| -rw-r--r-- | clang/include/clang/Basic/BuiltinsAMDGPU.def | 6 | ||||
| -rw-r--r-- | clang/lib/CodeGen/CGBuiltin.cpp | 8 | ||||
| -rw-r--r-- | clang/test/CodeGenOpenCL/builtins-amdgcn-error.cl | 32 | ||||
| -rw-r--r-- | clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 43 |
4 files changed, 89 insertions, 0 deletions
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index ea63ea10f84..b4314e6b0ff 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -70,6 +70,12 @@ 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. diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index ebccaecd78c..5f47cb4e3d7 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -7689,6 +7689,14 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, 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); diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-error.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-error.cl index 89c3e490ecd..5c676664c1a 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-error.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-error.cl @@ -4,7 +4,9 @@ // 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() { @@ -16,3 +18,33 @@ 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}} +} + diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index 6cac2a44bc7..2347bc8e41c 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -4,6 +4,7 @@ #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) @@ -199,6 +200,48 @@ 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) |
