================
@@ -0,0 +1,42 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgpu13.10-amd-amdhsa -x hip -std=c++11 -fclangir \
+// RUN: -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
+
+// RUN: %clang_cc1 -triple amdgpu13.10-amd-amdhsa -x hip -std=c++11 -fclangir \
+// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
+
+// RUN: %clang_cc1 -triple amdgpu13.10-amd-amdhsa -x hip -std=c++11 \
+// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
+
+// FIXME: CIR doesn't propagate the 'contract' fast-math flag to LLVM IR
+// calls yet, so LLVM check lines use {{.*}} to tolerate the difference
+// between CIR (no flags) and classic codegen ('contract').
+
+#define __device__ __attribute__((device))
+#define global __attribute__((address_space(1)))
+
+typedef _Float16 v2f16 __attribute__((ext_vector_type(2)));
+
+// CIR-LABEL:
@_Z36test_raw_ptr_buffer_atomic_fmin_fmaxPU3AS1fPU3AS1du22__amdgpu_buffer_rsrc_tfdii
+// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.atomic.fmin" {{.*}} ->
!cir.float
+// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.atomic.fmax" {{.*}} ->
!cir.float
+// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.atomic.fmin" {{.*}} ->
!cir.double
+// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.atomic.fmax" {{.*}} ->
!cir.double
+// LLVM: define{{.*}} void
@_Z36test_raw_ptr_buffer_atomic_fmin_fmaxPU3AS1fPU3AS1du22__amdgpu_buffer_rsrc_tfdii
+// LLVM: call{{.*}} float @llvm.amdgcn.raw.ptr.buffer.atomic.fmin.f32(float
{{.*}})
+// LLVM: call{{.*}} float @llvm.amdgcn.raw.ptr.buffer.atomic.fmax.f32(float
{{.*}})
+// LLVM: call{{.*}} double @llvm.amdgcn.raw.ptr.buffer.atomic.fmin.f64(double
{{.*}})
+// LLVM: call{{.*}} double @llvm.amdgcn.raw.ptr.buffer.atomic.fmax.f64(double
{{.*}})
----------------
steffenlarsen wrote:
Could you please add at the very least the types of the arguments to the checks?
https://github.com/llvm/llvm-project/pull/223137
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
