[llvm] [AMDGPU] Exclude CDNA-3 parts from cube instruction feature (PR #215484)
Arseniy Obolenskiy via llvm-commits
llvm-commits at lists.llvm.org
Tue Aug 11 01:19:59 PDT 2026
https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/215484
>From b09b5ea91694c5f6e7ca3a9de9cd409b98c55f7a Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Tue, 11 Aug 2026 09:24:40 +0200
Subject: [PATCH 1/2] [AMDGPU] Exclude CDNA-3 parts from cube instruction
feature
---
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 2 +
llvm/lib/Target/AMDGPU/AMDGPU.td | 6 +-
.../AMDGPU/llvm.amdgcn.cube-insts-cdna.ll | 97 +++++++++++++++++++
3 files changed, 103 insertions(+), 2 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cube-insts-cdna.ll
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 9101187fc6838..19d931c9085ca 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -556,6 +556,7 @@ def int_amdgcn_fmed3 :
[LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>]
>;
+let TargetFeatures = "cube-insts" in {
def int_amdgcn_cubeid : ClangBuiltin<"__builtin_amdgcn_cubeid">,
PureIntrinsic<[llvm_float_ty],
[llvm_float_ty, llvm_float_ty, llvm_float_ty]
@@ -575,6 +576,7 @@ def int_amdgcn_cubetc : ClangBuiltin<"__builtin_amdgcn_cubetc">,
PureIntrinsic<[llvm_float_ty],
[llvm_float_ty, llvm_float_ty, llvm_float_ty]
>;
+} // TargetFeatures = "cube-insts"
// v_ffbh_i32, as opposed to v_ffbh_u32. For v_ffbh_u32, llvm.ctlz
// should be used.
diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td
index 213d9fc1393de..5c1c5d136146e 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPU.td
+++ b/llvm/lib/Target/AMDGPU/AMDGPU.td
@@ -1631,7 +1631,7 @@ def FeatureGFX9 : GCNSubtargetFeatureGeneration<"GFX9",
FeatureUnalignedBufferAccess, FeatureUnalignedScratchAccess,
FeatureUnalignedDSAccess, FeatureNegativeScratchOffsetBug, FeatureGWS,
FeatureDefaultComponentZero,FeatureVmemWriteVgprInOrder, FeatureVMemToLDSLoad,
- FeatureCubeInsts, FeatureLerpInst, FeatureSadInsts, FeatureQsadInsts,
+ FeatureLerpInst, FeatureSadInsts, FeatureQsadInsts,
FeatureMsadInsts, FeatureMqsadPkInsts, FeatureMqsadInsts,
FeatureCvtNormInsts, FeatureCvtPkNormVOP2Insts,
FeatureCvtPkNormVOP3Insts, FeatureDX10ClampAndIEEEMode,
@@ -1846,6 +1846,7 @@ def FeatureISAVersion9_0_Consumer_Common : FeatureSet<
FeatureDsSrc2Insts,
FeatureExtendedImageInsts,
FeatureGDS,
+ FeatureCubeInsts,
FeaturePopsExitingWaveID])>;
def FeatureISAVersion9_Generic : FeatureSet<
@@ -1854,7 +1855,8 @@ def FeatureISAVersion9_Generic : FeatureSet<
def FeatureISAVersion9_0_MI_Common : FeatureSet<
!listconcat(FeatureISAVersion9_0_Common.Features,
- [FeatureFmaMixInsts,
+ [FeatureCubeInsts,
+ FeatureFmaMixInsts,
FeatureDLInsts,
FeatureDot1Insts,
FeatureDot2Insts,
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cube-insts-cdna.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cube-insts-cdna.ll
new file mode 100644
index 0000000000000..561e26d5cb9a8
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cube-insts-cdna.ll
@@ -0,0 +1,97 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 4
+; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx900 < %s | FileCheck %s -check-prefix=GFX900-SDAG
+; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx900 < %s | FileCheck %s -check-prefix=GFX900-GISEL
+; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx908 < %s | FileCheck %s -check-prefix=GFX908-SDAG
+; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx908 < %s | FileCheck %s -check-prefix=GFX908-GISEL
+; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx90a < %s | FileCheck %s -check-prefix=GFX90A-SDAG
+; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx90a < %s | FileCheck %s -check-prefix=GFX90A-GISEL
+; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx940 < %s 2>&1 | FileCheck %s -check-prefix=NOCUBE
+; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx942 < %s 2>&1 | FileCheck %s -check-prefix=NOCUBE
+; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx950 < %s 2>&1 | FileCheck %s -check-prefix=NOCUBE
+
+; The v_cube* ALU is graphics-pipe hardware, absent on the CDNA-3 compute-only
+; parts (gfx940, gfx942, gfx950). It is present on gfx900 (consumer) and on
+; the earlier CDNA-1/2 parts (gfx908, gfx90a).
+; NOCUBE: error: {{.*}} llvm.amdgcn.cubeid requires target feature
+
+declare float @llvm.amdgcn.cubeid(float, float, float)
+declare float @llvm.amdgcn.cubema(float, float, float)
+declare float @llvm.amdgcn.cubesc(float, float, float)
+declare float @llvm.amdgcn.cubetc(float, float, float)
+
+define amdgpu_ps <4 x float> @cube_vvv(float %a, float %b, float %c) {
+; GFX900-SDAG-LABEL: cube_vvv:
+; GFX900-SDAG: ; %bb.0:
+; GFX900-SDAG-NEXT: v_cubeid_f32 v6, v0, v1, v2
+; GFX900-SDAG-NEXT: v_cubema_f32 v4, v0, v1, v2
+; GFX900-SDAG-NEXT: v_cubesc_f32 v5, v0, v1, v2
+; GFX900-SDAG-NEXT: v_cubetc_f32 v3, v0, v1, v2
+; GFX900-SDAG-NEXT: v_mov_b32_e32 v0, v6
+; GFX900-SDAG-NEXT: v_mov_b32_e32 v1, v4
+; GFX900-SDAG-NEXT: v_mov_b32_e32 v2, v5
+; GFX900-SDAG-NEXT: ; return to shader part epilog
+;
+; GFX900-GISEL-LABEL: cube_vvv:
+; GFX900-GISEL: ; %bb.0:
+; GFX900-GISEL-NEXT: v_cubeid_f32 v6, v0, v1, v2
+; GFX900-GISEL-NEXT: v_cubema_f32 v4, v0, v1, v2
+; GFX900-GISEL-NEXT: v_cubesc_f32 v5, v0, v1, v2
+; GFX900-GISEL-NEXT: v_cubetc_f32 v3, v0, v1, v2
+; GFX900-GISEL-NEXT: v_mov_b32_e32 v0, v6
+; GFX900-GISEL-NEXT: v_mov_b32_e32 v1, v4
+; GFX900-GISEL-NEXT: v_mov_b32_e32 v2, v5
+; GFX900-GISEL-NEXT: ; return to shader part epilog
+;
+; GFX908-SDAG-LABEL: cube_vvv:
+; GFX908-SDAG: ; %bb.0:
+; GFX908-SDAG-NEXT: v_cubeid_f32 v6, v0, v1, v2
+; GFX908-SDAG-NEXT: v_cubema_f32 v4, v0, v1, v2
+; GFX908-SDAG-NEXT: v_cubesc_f32 v5, v0, v1, v2
+; GFX908-SDAG-NEXT: v_cubetc_f32 v3, v0, v1, v2
+; GFX908-SDAG-NEXT: v_mov_b32_e32 v0, v6
+; GFX908-SDAG-NEXT: v_mov_b32_e32 v1, v4
+; GFX908-SDAG-NEXT: v_mov_b32_e32 v2, v5
+; GFX908-SDAG-NEXT: ; return to shader part epilog
+;
+; GFX908-GISEL-LABEL: cube_vvv:
+; GFX908-GISEL: ; %bb.0:
+; GFX908-GISEL-NEXT: v_cubeid_f32 v6, v0, v1, v2
+; GFX908-GISEL-NEXT: v_cubema_f32 v4, v0, v1, v2
+; GFX908-GISEL-NEXT: v_cubesc_f32 v5, v0, v1, v2
+; GFX908-GISEL-NEXT: v_cubetc_f32 v3, v0, v1, v2
+; GFX908-GISEL-NEXT: v_mov_b32_e32 v0, v6
+; GFX908-GISEL-NEXT: v_mov_b32_e32 v1, v4
+; GFX908-GISEL-NEXT: v_mov_b32_e32 v2, v5
+; GFX908-GISEL-NEXT: ; return to shader part epilog
+;
+; GFX90A-SDAG-LABEL: cube_vvv:
+; GFX90A-SDAG: ; %bb.0:
+; GFX90A-SDAG-NEXT: v_cubeid_f32 v6, v0, v1, v2
+; GFX90A-SDAG-NEXT: v_cubema_f32 v4, v0, v1, v2
+; GFX90A-SDAG-NEXT: v_cubesc_f32 v5, v0, v1, v2
+; GFX90A-SDAG-NEXT: v_cubetc_f32 v3, v0, v1, v2
+; GFX90A-SDAG-NEXT: v_mov_b32_e32 v0, v6
+; GFX90A-SDAG-NEXT: v_mov_b32_e32 v1, v4
+; GFX90A-SDAG-NEXT: v_mov_b32_e32 v2, v5
+; GFX90A-SDAG-NEXT: ; return to shader part epilog
+;
+; GFX90A-GISEL-LABEL: cube_vvv:
+; GFX90A-GISEL: ; %bb.0:
+; GFX90A-GISEL-NEXT: v_cubeid_f32 v6, v0, v1, v2
+; GFX90A-GISEL-NEXT: v_cubema_f32 v4, v0, v1, v2
+; GFX90A-GISEL-NEXT: v_cubesc_f32 v5, v0, v1, v2
+; GFX90A-GISEL-NEXT: v_cubetc_f32 v3, v0, v1, v2
+; GFX90A-GISEL-NEXT: v_mov_b32_e32 v0, v6
+; GFX90A-GISEL-NEXT: v_mov_b32_e32 v1, v4
+; GFX90A-GISEL-NEXT: v_mov_b32_e32 v2, v5
+; GFX90A-GISEL-NEXT: ; return to shader part epilog
+ %cubeid = call float @llvm.amdgcn.cubeid(float %a, float %b, float %c)
+ %cubema = call float @llvm.amdgcn.cubema(float %a, float %b, float %c)
+ %cubesc = call float @llvm.amdgcn.cubesc(float %a, float %b, float %c)
+ %cubetc = call float @llvm.amdgcn.cubetc(float %a, float %b, float %c)
+ %vec0 = insertelement <4 x float> poison, float %cubeid, i32 0
+ %vec1 = insertelement <4 x float> %vec0, float %cubema, i32 1
+ %vec2 = insertelement <4 x float> %vec1, float %cubesc, i32 2
+ %vec3 = insertelement <4 x float> %vec2, float %cubetc, i32 3
+ ret <4 x float> %vec3
+}
>From 4fde2669f3c3f00026c394b5f616e53812deeb60 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Tue, 11 Aug 2026 10:19:44 +0200
Subject: [PATCH 2/2] Address comments
---
.../AMDGPU/llvm.amdgcn.cube-insts-cdna.ll | 97 -------------------
.../AMDGPU/llvm.amdgcn.cube-insts.err.ll | 40 ++++++++
2 files changed, 40 insertions(+), 97 deletions(-)
delete mode 100644 llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cube-insts-cdna.ll
create mode 100644 llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cube-insts.err.ll
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cube-insts-cdna.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cube-insts-cdna.ll
deleted file mode 100644
index 561e26d5cb9a8..0000000000000
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cube-insts-cdna.ll
+++ /dev/null
@@ -1,97 +0,0 @@
-; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 4
-; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx900 < %s | FileCheck %s -check-prefix=GFX900-SDAG
-; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx900 < %s | FileCheck %s -check-prefix=GFX900-GISEL
-; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx908 < %s | FileCheck %s -check-prefix=GFX908-SDAG
-; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx908 < %s | FileCheck %s -check-prefix=GFX908-GISEL
-; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx90a < %s | FileCheck %s -check-prefix=GFX90A-SDAG
-; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx90a < %s | FileCheck %s -check-prefix=GFX90A-GISEL
-; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx940 < %s 2>&1 | FileCheck %s -check-prefix=NOCUBE
-; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx942 < %s 2>&1 | FileCheck %s -check-prefix=NOCUBE
-; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx950 < %s 2>&1 | FileCheck %s -check-prefix=NOCUBE
-
-; The v_cube* ALU is graphics-pipe hardware, absent on the CDNA-3 compute-only
-; parts (gfx940, gfx942, gfx950). It is present on gfx900 (consumer) and on
-; the earlier CDNA-1/2 parts (gfx908, gfx90a).
-; NOCUBE: error: {{.*}} llvm.amdgcn.cubeid requires target feature
-
-declare float @llvm.amdgcn.cubeid(float, float, float)
-declare float @llvm.amdgcn.cubema(float, float, float)
-declare float @llvm.amdgcn.cubesc(float, float, float)
-declare float @llvm.amdgcn.cubetc(float, float, float)
-
-define amdgpu_ps <4 x float> @cube_vvv(float %a, float %b, float %c) {
-; GFX900-SDAG-LABEL: cube_vvv:
-; GFX900-SDAG: ; %bb.0:
-; GFX900-SDAG-NEXT: v_cubeid_f32 v6, v0, v1, v2
-; GFX900-SDAG-NEXT: v_cubema_f32 v4, v0, v1, v2
-; GFX900-SDAG-NEXT: v_cubesc_f32 v5, v0, v1, v2
-; GFX900-SDAG-NEXT: v_cubetc_f32 v3, v0, v1, v2
-; GFX900-SDAG-NEXT: v_mov_b32_e32 v0, v6
-; GFX900-SDAG-NEXT: v_mov_b32_e32 v1, v4
-; GFX900-SDAG-NEXT: v_mov_b32_e32 v2, v5
-; GFX900-SDAG-NEXT: ; return to shader part epilog
-;
-; GFX900-GISEL-LABEL: cube_vvv:
-; GFX900-GISEL: ; %bb.0:
-; GFX900-GISEL-NEXT: v_cubeid_f32 v6, v0, v1, v2
-; GFX900-GISEL-NEXT: v_cubema_f32 v4, v0, v1, v2
-; GFX900-GISEL-NEXT: v_cubesc_f32 v5, v0, v1, v2
-; GFX900-GISEL-NEXT: v_cubetc_f32 v3, v0, v1, v2
-; GFX900-GISEL-NEXT: v_mov_b32_e32 v0, v6
-; GFX900-GISEL-NEXT: v_mov_b32_e32 v1, v4
-; GFX900-GISEL-NEXT: v_mov_b32_e32 v2, v5
-; GFX900-GISEL-NEXT: ; return to shader part epilog
-;
-; GFX908-SDAG-LABEL: cube_vvv:
-; GFX908-SDAG: ; %bb.0:
-; GFX908-SDAG-NEXT: v_cubeid_f32 v6, v0, v1, v2
-; GFX908-SDAG-NEXT: v_cubema_f32 v4, v0, v1, v2
-; GFX908-SDAG-NEXT: v_cubesc_f32 v5, v0, v1, v2
-; GFX908-SDAG-NEXT: v_cubetc_f32 v3, v0, v1, v2
-; GFX908-SDAG-NEXT: v_mov_b32_e32 v0, v6
-; GFX908-SDAG-NEXT: v_mov_b32_e32 v1, v4
-; GFX908-SDAG-NEXT: v_mov_b32_e32 v2, v5
-; GFX908-SDAG-NEXT: ; return to shader part epilog
-;
-; GFX908-GISEL-LABEL: cube_vvv:
-; GFX908-GISEL: ; %bb.0:
-; GFX908-GISEL-NEXT: v_cubeid_f32 v6, v0, v1, v2
-; GFX908-GISEL-NEXT: v_cubema_f32 v4, v0, v1, v2
-; GFX908-GISEL-NEXT: v_cubesc_f32 v5, v0, v1, v2
-; GFX908-GISEL-NEXT: v_cubetc_f32 v3, v0, v1, v2
-; GFX908-GISEL-NEXT: v_mov_b32_e32 v0, v6
-; GFX908-GISEL-NEXT: v_mov_b32_e32 v1, v4
-; GFX908-GISEL-NEXT: v_mov_b32_e32 v2, v5
-; GFX908-GISEL-NEXT: ; return to shader part epilog
-;
-; GFX90A-SDAG-LABEL: cube_vvv:
-; GFX90A-SDAG: ; %bb.0:
-; GFX90A-SDAG-NEXT: v_cubeid_f32 v6, v0, v1, v2
-; GFX90A-SDAG-NEXT: v_cubema_f32 v4, v0, v1, v2
-; GFX90A-SDAG-NEXT: v_cubesc_f32 v5, v0, v1, v2
-; GFX90A-SDAG-NEXT: v_cubetc_f32 v3, v0, v1, v2
-; GFX90A-SDAG-NEXT: v_mov_b32_e32 v0, v6
-; GFX90A-SDAG-NEXT: v_mov_b32_e32 v1, v4
-; GFX90A-SDAG-NEXT: v_mov_b32_e32 v2, v5
-; GFX90A-SDAG-NEXT: ; return to shader part epilog
-;
-; GFX90A-GISEL-LABEL: cube_vvv:
-; GFX90A-GISEL: ; %bb.0:
-; GFX90A-GISEL-NEXT: v_cubeid_f32 v6, v0, v1, v2
-; GFX90A-GISEL-NEXT: v_cubema_f32 v4, v0, v1, v2
-; GFX90A-GISEL-NEXT: v_cubesc_f32 v5, v0, v1, v2
-; GFX90A-GISEL-NEXT: v_cubetc_f32 v3, v0, v1, v2
-; GFX90A-GISEL-NEXT: v_mov_b32_e32 v0, v6
-; GFX90A-GISEL-NEXT: v_mov_b32_e32 v1, v4
-; GFX90A-GISEL-NEXT: v_mov_b32_e32 v2, v5
-; GFX90A-GISEL-NEXT: ; return to shader part epilog
- %cubeid = call float @llvm.amdgcn.cubeid(float %a, float %b, float %c)
- %cubema = call float @llvm.amdgcn.cubema(float %a, float %b, float %c)
- %cubesc = call float @llvm.amdgcn.cubesc(float %a, float %b, float %c)
- %cubetc = call float @llvm.amdgcn.cubetc(float %a, float %b, float %c)
- %vec0 = insertelement <4 x float> poison, float %cubeid, i32 0
- %vec1 = insertelement <4 x float> %vec0, float %cubema, i32 1
- %vec2 = insertelement <4 x float> %vec1, float %cubesc, i32 2
- %vec3 = insertelement <4 x float> %vec2, float %cubetc, i32 3
- ret <4 x float> %vec3
-}
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cube-insts.err.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cube-insts.err.ll
new file mode 100644
index 0000000000000..c38c2d7d14158
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cube-insts.err.ll
@@ -0,0 +1,40 @@
+; RUN: not llc -global-isel=0 -mtriple=amdgpu9.4 -mcpu=gfx942 -filetype=null < %s 2>&1 | FileCheck %s -check-prefix=GFX942
+; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgpu9.4 -mcpu=gfx942 -filetype=null < %s 2>&1 | FileCheck %s -check-prefix=GFX942
+; RUN: not llc -global-isel=0 -mtriple=amdgpu9.4 -mcpu=gfx950 -filetype=null < %s 2>&1 | FileCheck %s -check-prefix=GFX950
+; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgpu9.4 -mcpu=gfx950 -filetype=null < %s 2>&1 | FileCheck %s -check-prefix=GFX950
+
+; The v_cube* ALU is graphics-pipe hardware, absent on the CDNA-3 compute-only
+; parts gfx942 and gfx950.
+
+; GFX942: error: {{.*}} llvm.amdgcn.cubeid requires target feature 'cube-insts'
+; GFX950: error: {{.*}} llvm.amdgcn.cubeid requires target feature 'cube-insts'
+define float @test_cubeid(float %a, float %b, float %c) {
+ %result = call float @llvm.amdgcn.cubeid(float %a, float %b, float %c)
+ ret float %result
+}
+
+; GFX942: error: {{.*}} llvm.amdgcn.cubema requires target feature 'cube-insts'
+; GFX950: error: {{.*}} llvm.amdgcn.cubema requires target feature 'cube-insts'
+define float @test_cubema(float %a, float %b, float %c) {
+ %result = call float @llvm.amdgcn.cubema(float %a, float %b, float %c)
+ ret float %result
+}
+
+; GFX942: error: {{.*}} llvm.amdgcn.cubesc requires target feature 'cube-insts'
+; GFX950: error: {{.*}} llvm.amdgcn.cubesc requires target feature 'cube-insts'
+define float @test_cubesc(float %a, float %b, float %c) {
+ %result = call float @llvm.amdgcn.cubesc(float %a, float %b, float %c)
+ ret float %result
+}
+
+; GFX942: error: {{.*}} llvm.amdgcn.cubetc requires target feature 'cube-insts'
+; GFX950: error: {{.*}} llvm.amdgcn.cubetc requires target feature 'cube-insts'
+define float @test_cubetc(float %a, float %b, float %c) {
+ %result = call float @llvm.amdgcn.cubetc(float %a, float %b, float %c)
+ ret float %result
+}
+
+declare float @llvm.amdgcn.cubeid(float, float, float)
+declare float @llvm.amdgcn.cubema(float, float, float)
+declare float @llvm.amdgcn.cubesc(float, float, float)
+declare float @llvm.amdgcn.cubetc(float, float, float)
More information about the llvm-commits
mailing list