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 <[email protected]> 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 <[email protected]> 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 <[email protected]> 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 <[email protected]> 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 <[email protected]> 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); -} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
