BUILTIN(__builtin_amdgcn_rcpf, "ff", "nc")
BUILTIN(__builtin_amdgcn_rsq, "dd", "nc")
BUILTIN(__builtin_amdgcn_rsqf, "ff", "nc")
-BUILTIN(__builtin_amdgcn_rsq_clamped, "dd", "nc")
-BUILTIN(__builtin_amdgcn_rsq_clampedf, "ff", "nc")
+BUILTIN(__builtin_amdgcn_rsq_clamp, "dd", "nc")
+BUILTIN(__builtin_amdgcn_rsq_clampf, "ff", "nc")
BUILTIN(__builtin_amdgcn_ldexp, "ddi", "nc")
BUILTIN(__builtin_amdgcn_ldexpf, "ffi", "nc")
BUILTIN(__builtin_amdgcn_class, "bdi", "nc")
case AMDGPU::BI__builtin_amdgcn_rsq:
case AMDGPU::BI__builtin_amdgcn_rsqf:
return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_rsq);
- case AMDGPU::BI__builtin_amdgcn_rsq_clamped:
- case AMDGPU::BI__builtin_amdgcn_rsq_clampedf:
- return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_rsq_clamped);
+ case AMDGPU::BI__builtin_amdgcn_rsq_clamp:
+ case AMDGPU::BI__builtin_amdgcn_rsq_clampf:
+ return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_rsq_clamp);
case AMDGPU::BI__builtin_amdgcn_ldexp:
case AMDGPU::BI__builtin_amdgcn_ldexpf:
return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_ldexp);
*out = __builtin_amdgcn_rsq(a);
}
-// CHECK-LABEL: @test_rsq_clamped_f32
-// CHECK: call float @llvm.amdgcn.rsq.clamped.f32
-void test_rsq_clamped_f32(global float* out, float a)
+// CHECK-LABEL: @test_rsq_clamp_f32
+// CHECK: call float @llvm.amdgcn.rsq.clamp.f32
+void test_rsq_clamp_f32(global float* out, float a)
{
- *out = __builtin_amdgcn_rsq_clampedf(a);
+ *out = __builtin_amdgcn_rsq_clampf(a);
}
-// CHECK-LABEL: @test_rsq_clamped_f64
-// CHECK: call double @llvm.amdgcn.rsq.clamped.f64
-void test_rsq_clamped_f64(global double* out, double a)
+// CHECK-LABEL: @test_rsq_clamp_f64
+// CHECK: call double @llvm.amdgcn.rsq.clamp.f64
+void test_rsq_clamp_f64(global double* out, double a)
{
- *out = __builtin_amdgcn_rsq_clamped(a);
+ *out = __builtin_amdgcn_rsq_clamp(a);
}
// CHECK-LABEL: @test_ldexp_f32