From: Wei Ding Date: Fri, 5 Aug 2016 15:38:46 +0000 (+0000) Subject: AMDGPU : Add Clang builtin intrinsics for compare with the full X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=85ba22864280157c64f3acc490b9602282bba101;p=clang AMDGPU : Add Clang builtin intrinsics for compare with the full wavefront result. Differential Revision: http://reviews.llvm.org/D22934 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@277824 91177308-0d34-0410-b5e6-96231b3b80d8 --- diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def index ea63ea10f8..b4314e6b0f 100644 --- a/include/clang/Basic/BuiltinsAMDGPU.def +++ b/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/lib/CodeGen/CGBuiltin.cpp b/lib/CodeGen/CGBuiltin.cpp index ebccaecd78..5f47cb4e3d 100644 --- a/lib/CodeGen/CGBuiltin.cpp +++ b/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/test/CodeGenOpenCL/builtins-amdgcn-error.cl b/test/CodeGenOpenCL/builtins-amdgcn-error.cl index 89c3e490ec..5c676664c1 100644 --- a/test/CodeGenOpenCL/builtins-amdgcn-error.cl +++ b/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/test/CodeGenOpenCL/builtins-amdgcn.cl b/test/CodeGenOpenCL/builtins-amdgcn.cl index 6cac2a44bc..2347bc8e41 100644 --- a/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/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)