================
@@ -0,0 +1,42 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
+// RUN: -target-cpu gfx1250 -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
+// RUN: -target-cpu gfx1250 -fcuda-is-device -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \
+// RUN: -target-cpu gfx1250 -fcuda-is-device -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
+
+#define __device__ __attribute__((device))
+
+//===----------------------------------------------------------------------===//
+// Test AMDGPU builtins
+//===----------------------------------------------------------------------===//
+
+// CIR-LABEL: @_Z14test_tanhf_f32Pff
+// CIR: cir.call_llvm_intrinsic "amdgcn.tanh" {{.*}} : (!cir.float) ->
!cir.float
+// LLVM: define{{.*}} void @_Z14test_tanhf_f32Pff
+// LLVM: call{{.*}} float @llvm.amdgcn.tanh.f32(float %{{.*}})
+__device__ void test_tanhf_f32(float* out, float a) {
+ *out = __builtin_amdgcn_tanhf(a);
+}
+
+// CIR-LABEL: @_Z13test_tanh_f16PDF16_DF16_
+// CIR: cir.call_llvm_intrinsic "amdgcn.tanh" {{.*}} : (!cir.f16) -> !cir.f16
+// LLVM: define{{.*}} void @_Z13test_tanh_f16PDF16_DF16_
+// LLVM: call{{.*}} half @llvm.{{((amdgcn.){0,1})}}tanh.f16(half %{{.*}})
----------------
ranapratap55 wrote:
This pattern also matches the generic `llvm.tanh.f16`. Tighten to match OGCG
exactly. Check `_f32` and `_bf16` cases for ref.
https://github.com/llvm/llvm-project/pull/197852
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
