[clang] [CIR][AMDGPU] Add support for AMDGCN sin builtins (PR #197784)
Ayokunle Amodu via cfe-commits
cfe-commits at lists.llvm.org
Wed Jul 15 12:52:43 PDT 2026
https://github.com/ayokunle321 updated https://github.com/llvm/llvm-project/pull/197784
>From 6a1b69f0261fdbc10820a3105d92258a28a4d6e0 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Thu, 14 May 2026 15:04:15 -0400
Subject: [PATCH 1/5] add amdgcn sin builtin
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 39 ++++++++---
.../CodeGenHIP/builtins-amdgcn-gfx1250.hip | 26 ++++++++
.../CIR/CodeGenHIP/builtins-amdgcn-vi.hip | 65 +++++++++++++++++++
clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip | 8 +++
4 files changed, 129 insertions(+), 9 deletions(-)
create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 0a7ba0c194400..e8cbc2b4f90a2 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -211,12 +211,33 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
return mlir::Value{};
}
case AMDGPU::BI__builtin_amdgcn_permlane16:
- case AMDGPU::BI__builtin_amdgcn_permlanex16:
+ case AMDGPU::BI__builtin_amdgcn_permlanex16: {
+ llvm::StringRef intrinsicName =
+ builtinId == AMDGPU::BI__builtin_amdgcn_permlane16
+ ? "amdgcn.permlane16"
+ : "amdgcn.permlanex16";
+ mlir::Value src0 = emitScalarExpr(expr->getArg(0));
+ mlir::Value src1 = emitScalarExpr(expr->getArg(1));
+ mlir::Value src2 = emitScalarExpr(expr->getArg(2));
+ mlir::Value src3 = emitScalarExpr(expr->getArg(3));
+ mlir::Value src4 = emitScalarExpr(expr->getArg(4));
+ mlir::Value src5 = emitScalarExpr(expr->getArg(5));
+ mlir::Value result =
+ cir::LLVMIntrinsicCallOp::create(builder, getLoc(expr->getExprLoc()),
+ builder.getStringAttr(intrinsicName),
+ src0.getType(),
+ {src0, src1, src2, src3, src4, src5})
+ .getResult();
+ return result;
+ }
case AMDGPU::BI__builtin_amdgcn_permlane64: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ mlir::Value src = emitScalarExpr(expr->getArg(0));
+ mlir::Value result =
+ cir::LLVMIntrinsicCallOp::create(
+ builder, getLoc(expr->getExprLoc()),
+ builder.getStringAttr("amdgcn.permlane64"), src.getType(), {src})
+ .getResult();
+ return result;
}
case AMDGPU::BI__builtin_amdgcn_readlane:
return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.readlane")
@@ -284,10 +305,10 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
case AMDGPU::BI__builtin_amdgcn_sinf:
case AMDGPU::BI__builtin_amdgcn_sinh:
case AMDGPU::BI__builtin_amdgcn_sin_bf16: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ mlir::Value src = emitScalarExpr(expr->getArg(0));
+ return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()),
+ "amdgcn.sin", src.getType(),
+ mlir::ValueRange{src});
}
case AMDGPU::BI__builtin_amdgcn_cosf:
case AMDGPU::BI__builtin_amdgcn_cosh:
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
new file mode 100644
index 0000000000000..a36b22b66da74
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
@@ -0,0 +1,26 @@
+#include "../CodeGenCUDA/Inputs/cuda.h"
+
+// 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
+
+//===----------------------------------------------------------------------===//
+// Test AMDGPU builtins
+//===----------------------------------------------------------------------===//
+
+// CIR-LABEL: @_Z13test_sin_bf16PDF16bDF16b
+// CIR: cir.call_llvm_intrinsic "amdgcn.sin" {{.*}} : (!cir.bf16) -> !cir.bf16
+// LLVM: define{{.*}} void @_Z13test_sin_bf16PDF16bDF16b
+// LLVM: call{{.*}} bfloat @llvm.amdgcn.sin.bf16(bfloat %{{.*}})
+__device__ void test_sin_bf16(__bf16* out, __bf16 a) {
+ *out = __builtin_amdgcn_sin_bf16(a);
+}
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip
new file mode 100644
index 0000000000000..6d166c80855a1
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip
@@ -0,0 +1,65 @@
+#include "../CodeGenCUDA/Inputs/cuda.h"
+
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
+// RUN: -target-cpu tonga -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
+
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
+// RUN: -target-cpu gfx900 -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
+
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
+// RUN: -target-cpu gfx1010 -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
+
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
+// RUN: -target-cpu gfx1012 -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 tonga -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 -fclangir \
+// RUN: -target-cpu gfx900 -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 -fclangir \
+// RUN: -target-cpu gfx1010 -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 -fclangir \
+// RUN: -target-cpu gfx1012 -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 tonga -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 gfx900 -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 gfx1010 -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 gfx1012 -fcuda-is-device -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
+
+//===----------------------------------------------------------------------===//
+// Test AMDGPU builtins
+//===----------------------------------------------------------------------===//
+
+// CIR-LABEL: @_Z12test_sin_f16PDF16_DF16_
+// CIR: cir.call_llvm_intrinsic "amdgcn.sin" {{.*}} : (!cir.f16) -> !cir.f16
+// LLVM: define{{.*}} void @_Z12test_sin_f16PDF16_DF16_
+// LLVM: call{{.*}} half @llvm.{{((amdgcn.){0,1})}}sin.f16(half %{{.*}})
+__device__ void test_sin_f16(_Float16* out, _Float16 a) {
+ *out = __builtin_amdgcn_sinh(a);
+}
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
index ca708bca8587b..fe4fc45eb041d 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
@@ -111,3 +111,11 @@ __device__ void test_readfirstlane(int* out, int a) {
__device__ void test_dispatch_ptr(__attribute__((address_space(4))) void ** out) {
*out = (__attribute__((address_space(4))) void *)__builtin_amdgcn_dispatch_ptr();
}
+
+// CIR-LABEL: @_Z13test_sinf_f32Pff
+// CIR: cir.call_llvm_intrinsic "amdgcn.sin" {{.*}} : (!cir.float) -> !cir.float
+// LLVM: define{{.*}} void @_Z13test_sinf_f32Pff
+// LLVM: call{{.*}} float @llvm.amdgcn.sin.f32(float %{{.*}})
+__device__ void test_sinf_f32(float* out, float a) {
+ *out = __builtin_amdgcn_sinf(a);
+}
>From 796f135aa2563a233ab99a00694c0e477dd11807 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Thu, 14 May 2026 15:11:34 -0400
Subject: [PATCH 2/5] fix style
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 5 ++---
1 file changed, 2 insertions(+), 3 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index e8cbc2b4f90a2..9b938bf1c8fbd 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -306,9 +306,8 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
case AMDGPU::BI__builtin_amdgcn_sinh:
case AMDGPU::BI__builtin_amdgcn_sin_bf16: {
mlir::Value src = emitScalarExpr(expr->getArg(0));
- return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()),
- "amdgcn.sin", src.getType(),
- mlir::ValueRange{src});
+ return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), "amdgcn.sin",
+ src.getType(), mlir::ValueRange{src});
}
case AMDGPU::BI__builtin_amdgcn_cosf:
case AMDGPU::BI__builtin_amdgcn_cosh:
>From 7155a7c403ba5128483d98602c4b7185a34b541c Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Thu, 14 May 2026 15:14:24 -0400
Subject: [PATCH 3/5] fix style
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 31 +++----------------
1 file changed, 5 insertions(+), 26 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 9b938bf1c8fbd..064ec14e62c9b 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -211,33 +211,12 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
return mlir::Value{};
}
case AMDGPU::BI__builtin_amdgcn_permlane16:
- case AMDGPU::BI__builtin_amdgcn_permlanex16: {
- llvm::StringRef intrinsicName =
- builtinId == AMDGPU::BI__builtin_amdgcn_permlane16
- ? "amdgcn.permlane16"
- : "amdgcn.permlanex16";
- mlir::Value src0 = emitScalarExpr(expr->getArg(0));
- mlir::Value src1 = emitScalarExpr(expr->getArg(1));
- mlir::Value src2 = emitScalarExpr(expr->getArg(2));
- mlir::Value src3 = emitScalarExpr(expr->getArg(3));
- mlir::Value src4 = emitScalarExpr(expr->getArg(4));
- mlir::Value src5 = emitScalarExpr(expr->getArg(5));
- mlir::Value result =
- cir::LLVMIntrinsicCallOp::create(builder, getLoc(expr->getExprLoc()),
- builder.getStringAttr(intrinsicName),
- src0.getType(),
- {src0, src1, src2, src3, src4, src5})
- .getResult();
- return result;
- }
+ case AMDGPU::BI__builtin_amdgcn_permlanex16:
case AMDGPU::BI__builtin_amdgcn_permlane64: {
- mlir::Value src = emitScalarExpr(expr->getArg(0));
- mlir::Value result =
- cir::LLVMIntrinsicCallOp::create(
- builder, getLoc(expr->getExprLoc()),
- builder.getStringAttr("amdgcn.permlane64"), src.getType(), {src})
- .getResult();
- return result;
+ cgm.errorNYI(expr->getSourceRange(),
+ std::string("unimplemented AMDGPU builtin call: ") +
+ getContext().BuiltinInfo.getName(builtinId));
+ return mlir::Value{};
}
case AMDGPU::BI__builtin_amdgcn_readlane:
return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.readlane")
>From 3494b898a66d14663b5e33e8bec292dd3bdd5f92 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Sat, 16 May 2026 17:45:14 -0400
Subject: [PATCH 4/5] remove regex matching in CHECK line
---
clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip
index 6d166c80855a1..5f2b6cc2489bd 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip
@@ -59,7 +59,7 @@
// CIR-LABEL: @_Z12test_sin_f16PDF16_DF16_
// CIR: cir.call_llvm_intrinsic "amdgcn.sin" {{.*}} : (!cir.f16) -> !cir.f16
// LLVM: define{{.*}} void @_Z12test_sin_f16PDF16_DF16_
-// LLVM: call{{.*}} half @llvm.{{((amdgcn.){0,1})}}sin.f16(half %{{.*}})
+// LLVM: call{{.*}} half @llvm.amdgcn.sin.f16(half %{{.*}})
__device__ void test_sin_f16(_Float16* out, _Float16 a) {
*out = __builtin_amdgcn_sinh(a);
}
>From 78fb56cf312cf88d4e7a6815e5d898782eb871dc Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Fri, 10 Jul 2026 15:40:14 -0400
Subject: [PATCH 5/5] switch manual codegen for function
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 4 +-
.../CodeGenHIP/builtins-amdgcn-gfx1250.hip | 4 +-
.../CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip | 8 +++
.../CIR/CodeGenHIP/builtins-amdgcn-vi.hip | 65 -------------------
4 files changed, 11 insertions(+), 70 deletions(-)
delete mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 064ec14e62c9b..ff508db32ce41 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -284,9 +284,7 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
case AMDGPU::BI__builtin_amdgcn_sinf:
case AMDGPU::BI__builtin_amdgcn_sinh:
case AMDGPU::BI__builtin_amdgcn_sin_bf16: {
- mlir::Value src = emitScalarExpr(expr->getArg(0));
- return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), "amdgcn.sin",
- src.getType(), mlir::ValueRange{src});
+ return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.sin").getValue();
}
case AMDGPU::BI__builtin_amdgcn_cosf:
case AMDGPU::BI__builtin_amdgcn_cosh:
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
index a36b22b66da74..08aafe539599f 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
@@ -1,5 +1,3 @@
-#include "../CodeGenCUDA/Inputs/cuda.h"
-
// 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
@@ -13,6 +11,8 @@
// 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
//===----------------------------------------------------------------------===//
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
index c21359ff7a88b..494a33af91918 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
@@ -63,3 +63,11 @@
__device__ void test_div_fixup_f16(_Float16* out, _Float16 a, _Float16 b, _Float16 c) {
*out = __builtin_amdgcn_div_fixuph(a, b, c);
}
+
+// CIR-LABEL: @_Z12test_sin_f16PDF16_DF16_
+// CIR: cir.call_llvm_intrinsic "amdgcn.sin" {{.*}} : (!cir.f16) -> !cir.f16
+// LLVM: define{{.*}} void @_Z12test_sin_f16PDF16_DF16_
+// LLVM: call{{.*}} half @llvm.amdgcn.sin.f16(half %{{.*}})
+__device__ void test_sin_f16(_Float16* out, _Float16 a) {
+ *out = __builtin_amdgcn_sinh(a);
+}
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip
deleted file mode 100644
index 5f2b6cc2489bd..0000000000000
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip
+++ /dev/null
@@ -1,65 +0,0 @@
-#include "../CodeGenCUDA/Inputs/cuda.h"
-
-// REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
-// RUN: -target-cpu tonga -fcuda-is-device -emit-cir %s -o %t.cir
-// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
-
-// REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
-// RUN: -target-cpu gfx900 -fcuda-is-device -emit-cir %s -o %t.cir
-// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
-
-// REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
-// RUN: -target-cpu gfx1010 -fcuda-is-device -emit-cir %s -o %t.cir
-// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
-
-// REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
-// RUN: -target-cpu gfx1012 -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 tonga -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 -fclangir \
-// RUN: -target-cpu gfx900 -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 -fclangir \
-// RUN: -target-cpu gfx1010 -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 -fclangir \
-// RUN: -target-cpu gfx1012 -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 tonga -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 gfx900 -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 gfx1010 -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 gfx1012 -fcuda-is-device -emit-llvm %s -o %t.ll
-// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
-
-//===----------------------------------------------------------------------===//
-// Test AMDGPU builtins
-//===----------------------------------------------------------------------===//
-
-// CIR-LABEL: @_Z12test_sin_f16PDF16_DF16_
-// CIR: cir.call_llvm_intrinsic "amdgcn.sin" {{.*}} : (!cir.f16) -> !cir.f16
-// LLVM: define{{.*}} void @_Z12test_sin_f16PDF16_DF16_
-// LLVM: call{{.*}} half @llvm.amdgcn.sin.f16(half %{{.*}})
-__device__ void test_sin_f16(_Float16* out, _Float16 a) {
- *out = __builtin_amdgcn_sinh(a);
-}
More information about the cfe-commits
mailing list