Author: wdng Date: Fri Aug 5 10:38:46 2016 New Revision: 277824 URL: http://llvm.org/viewvc/llvm-project?rev=277824&view=rev Log: AMDGPU : Add Clang builtin intrinsics for compare with the full wavefront result.
Differential Revision: http://reviews.llvm.org/D22934 Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def cfe/trunk/lib/CodeGen/CGBuiltin.cpp cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def?rev=277824&r1=277823&r2=277824&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original) +++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Fri Aug 5 10:38:46 2016 @@ -70,6 +70,12 @@ BUILTIN(__builtin_amdgcn_cubetc, "ffff", 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. Modified: cfe/trunk/lib/CodeGen/CGBuiltin.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGBuiltin.cpp?rev=277824&r1=277823&r2=277824&view=diff ============================================================================== --- cfe/trunk/lib/CodeGen/CGBuiltin.cpp (original) +++ cfe/trunk/lib/CodeGen/CGBuiltin.cpp Fri Aug 5 10:38:46 2016 @@ -7689,6 +7689,14 @@ Value *CodeGenFunction::EmitAMDGPUBuilti 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); Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl?rev=277824&r1=277823&r2=277824&view=diff ============================================================================== --- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl (original) +++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl Fri Aug 5 10:38:46 2016 @@ -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}} +} + Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl?rev=277824&r1=277823&r2=277824&view=diff ============================================================================== --- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl (original) +++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Fri Aug 5 10:38:46 2016 @@ -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, i *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) _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
