Author: arsenm Date: Mon Jan 30 21:42:07 2017 New Revision: 293600 URL: http://llvm.org/viewvc/llvm-project?rev=293600&view=rev Log: AMDGPU: Add builtin for fmed3 intrinsic
Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def cfe/trunk/lib/CodeGen/CGBuiltin.cpp 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=293600&r1=293599&r2=293600&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original) +++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Mon Jan 30 21:42:07 2017 @@ -81,6 +81,7 @@ BUILTIN(__builtin_amdgcn_sicmpl, "LUiLiL BUILTIN(__builtin_amdgcn_fcmp, "LUiddIi", "nc") BUILTIN(__builtin_amdgcn_fcmpf, "LUiffIi", "nc") BUILTIN(__builtin_amdgcn_ds_swizzle, "iiIi", "nc") +BUILTIN(__builtin_amdgcn_fmed3f, "ffff", "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=293600&r1=293599&r2=293600&view=diff ============================================================================== --- cfe/trunk/lib/CodeGen/CGBuiltin.cpp (original) +++ cfe/trunk/lib/CodeGen/CGBuiltin.cpp Mon Jan 30 21:42:07 2017 @@ -8397,7 +8397,8 @@ Value *CodeGenFunction::EmitAMDGPUBuilti case AMDGPU::BI__builtin_amdgcn_classf: case AMDGPU::BI__builtin_amdgcn_classh: return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_class); - + case AMDGPU::BI__builtin_amdgcn_fmed3f: + return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_fmed3); case AMDGPU::BI__builtin_amdgcn_read_exec: { CallInst *CI = cast<CallInst>( EmitSpecialRegisterBuiltin(*this, E, Int64Ty, Int64Ty, true, "exec")); Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl?rev=293600&r1=293599&r2=293600&view=diff ============================================================================== --- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl (original) +++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Mon Jan 30 21:42:07 2017 @@ -396,6 +396,13 @@ void test_get_local_id(int d, global int } } +// CHECK-LABEL: @test_fmed3_f32 +// CHECK: call float @llvm.amdgcn.fmed3.f32( +void test_fmed3_f32(global float* out, float a, float b, float c) +{ + *out = __builtin_amdgcn_fmed3f(a, b, c); +} + // CHECK-DAG: [[WI_RANGE]] = !{i32 0, i32 1024} // CHECK-DAG: attributes #[[NOUNWIND_READONLY:[0-9]+]] = { nounwind readonly } // CHECK-DAG: attributes #[[READ_EXEC_ATTRS]] = { convergent } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits