Author: arsenm
Date: Fri Jul 15 16:33:02 2016
New Revision: 275622
URL: http://llvm.org/viewvc/llvm-project?rev=275622&view=rev
Log:
AMDGPU: Update for rsq intrinsic changes
Modified:
cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
cfe/trunk/lib/CodeGen/CGBuiltin.cpp
cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
cfe/trunk/test/CodeGenOpenCL/builtins-r600.cl
Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def?rev=275622&r1=275621&r2=275622&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original)
+++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Fri Jul 15 16:33:02 2016
@@ -96,12 +96,13 @@ BUILTIN(__builtin_r600_read_tidig_x, "Ui
BUILTIN(__builtin_r600_read_tidig_y, "Ui", "nc")
BUILTIN(__builtin_r600_read_tidig_z, "Ui", "nc")
+BUILTIN(__builtin_r600_recipsqrt_ieee, "dd", "nc")
+BUILTIN(__builtin_r600_recipsqrt_ieeef, "ff", "nc")
+
//===----------------------------------------------------------------------===//
// Legacy names with amdgpu prefix
//===----------------------------------------------------------------------===//
-BUILTIN(__builtin_amdgpu_rsq, "dd", "nc")
-BUILTIN(__builtin_amdgpu_rsqf, "ff", "nc")
BUILTIN(__builtin_amdgpu_ldexp, "ddi", "nc")
BUILTIN(__builtin_amdgpu_ldexpf, "ffi", "nc")
Modified: cfe/trunk/lib/CodeGen/CGBuiltin.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGBuiltin.cpp?rev=275622&r1=275621&r2=275622&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGBuiltin.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGBuiltin.cpp Fri Jul 15 16:33:02 2016
@@ -7671,19 +7671,6 @@ Value *CodeGenFunction::EmitAMDGPUBuilti
CI->setConvergent();
return CI;
}
- // Legacy amdgpu prefix
- case AMDGPU::BI__builtin_amdgpu_rsq:
- case AMDGPU::BI__builtin_amdgpu_rsqf: {
- if (getTarget().getTriple().getArch() == Triple::amdgcn)
- return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_rsq);
- return emitUnaryBuiltin(*this, E, Intrinsic::r600_rsq);
- }
- case AMDGPU::BI__builtin_amdgpu_ldexp:
- case AMDGPU::BI__builtin_amdgpu_ldexpf: {
- if (getTarget().getTriple().getArch() == Triple::amdgcn)
- return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_ldexp);
- return emitFPIntBuiltin(*this, E, Intrinsic::AMDGPU_ldexp);
- }
// amdgcn workitem
case AMDGPU::BI__builtin_amdgcn_workitem_id_x:
@@ -7693,13 +7680,24 @@ Value *CodeGenFunction::EmitAMDGPUBuilti
case AMDGPU::BI__builtin_amdgcn_workitem_id_z:
return emitRangedBuiltin(*this, Intrinsic::amdgcn_workitem_id_z, 0, 1024);
- // r600 workitem
+ // r600 intrinsics
+ case AMDGPU::BI__builtin_r600_recipsqrt_ieee:
+ case AMDGPU::BI__builtin_r600_recipsqrt_ieeef:
+ return emitUnaryBuiltin(*this, E, Intrinsic::r600_recipsqrt_ieee);
case AMDGPU::BI__builtin_r600_read_tidig_x:
return emitRangedBuiltin(*this, Intrinsic::r600_read_tidig_x, 0, 1024);
case AMDGPU::BI__builtin_r600_read_tidig_y:
return emitRangedBuiltin(*this, Intrinsic::r600_read_tidig_y, 0, 1024);
case AMDGPU::BI__builtin_r600_read_tidig_z:
return emitRangedBuiltin(*this, Intrinsic::r600_read_tidig_z, 0, 1024);
+
+ // Legacy amdgpu prefix
+ case AMDGPU::BI__builtin_amdgpu_ldexp:
+ case AMDGPU::BI__builtin_amdgpu_ldexpf: {
+ if (getTarget().getTriple().getArch() == Triple::amdgcn)
+ return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_ldexp);
+ return emitFPIntBuiltin(*this, E, Intrinsic::AMDGPU_ldexp);
+ }
default:
return nullptr;
}
Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl?rev=275622&r1=275621&r2=275622&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Fri Jul 15 16:33:02 2016
@@ -270,20 +270,6 @@ void test_read_exec(global ulong* out) {
// Legacy intrinsics with AMDGPU prefix
-// CHECK-LABEL: @test_legacy_rsq_f32
-// CHECK: call float @llvm.amdgcn.rsq.f32
-void test_legacy_rsq_f32(global float* out, float a)
-{
- *out = __builtin_amdgpu_rsqf(a);
-}
-
-// CHECK-LABEL: @test_legacy_rsq_f64
-// CHECK: call double @llvm.amdgcn.rsq.f64
-void test_legacy_rsq_f64(global double* out, double a)
-{
- *out = __builtin_amdgpu_rsq(a);
-}
-
// CHECK-LABEL: @test_legacy_ldexp_f32
// CHECK: call float @llvm.amdgcn.ldexp.f32
void test_legacy_ldexp_f32(global float* out, float a, int b)
Modified: cfe/trunk/test/CodeGenOpenCL/builtins-r600.cl
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-r600.cl?rev=275622&r1=275621&r2=275622&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-r600.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-r600.cl Fri Jul 15 16:33:02 2016
@@ -1,19 +1,19 @@
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple r600-unknown-unknown -target-cpu cypress -S
-emit-llvm -o - %s | FileCheck %s
-// CHECK-LABEL: @test_rsq_f32
-// CHECK: call float @llvm.r600.rsq.f32
-void test_rsq_f32(global float* out, float a)
+// CHECK-LABEL: @test_recipsqrt_ieee_f32
+// CHECK: call float @llvm.r600.recipsqrt.ieee.f32
+void test_recipsqrt_ieee_f32(global float* out, float a)
{
- *out = __builtin_amdgpu_rsqf(a);
+ *out = __builtin_r600_recipsqrt_ieeef(a);
}
#if cl_khr_fp64
-// XCHECK-LABEL: @test_rsq_f64
-// XCHECK: call double @llvm.r600.rsq.f64
-void test_rsq_f64(global double* out, double a)
+// XCHECK-LABEL: @test_recipsqrt_ieee_f64
+// XCHECK: call double @llvm.r600.recipsqrt.ieee.f64
+void test_recipsqrt_ieee_f64(global double* out, double a)
{
- *out = __builtin_amdgpu_rsq(a);
+ *out = __builtin_r600_recipsqrt_ieee(a);
}
#endif
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
