https://github.com/Artem-B created
https://github.com/llvm/llvm-project/pull/124818
Fixes build break with CUDA-12.8 introduced in #123398
>From ebb322550e72a4c7c58d990340701248d6f61c95 Mon Sep 17 00:00:00 2001
From: Artem Belevich <t...@google.com>
Date: Tue, 28 Jan 2025 10:45:16 -0800
Subject: [PATCH] [CUDA] Make target intrinsics work with ptx 8.7
---
clang/include/clang/Basic/BuiltinsNVPTX.td | 3 ++-
clang/test/CodeGen/builtins-nvptx.c | 5 +++++
2 files changed, 7 insertions(+), 1 deletion(-)
diff --git a/clang/include/clang/Basic/BuiltinsNVPTX.td
b/clang/include/clang/Basic/BuiltinsNVPTX.td
index b43e8ba57f7a0b..9d24a992563a45 100644
--- a/clang/include/clang/Basic/BuiltinsNVPTX.td
+++ b/clang/include/clang/Basic/BuiltinsNVPTX.td
@@ -48,8 +48,9 @@ class PTX<string version, PTXFeatures newer> : PTXFeatures {
let Features = !strconcat("ptx", version, "|", newer.Features);
}
-let Features = "ptx86" in def PTX86 : PTXFeatures;
+let Features = "ptx87" in def PTX87 : PTXFeatures;
+def PTX86 : PTX<"86", PTX87>;
def PTX85 : PTX<"85", PTX86>;
def PTX84 : PTX<"84", PTX85>;
def PTX83 : PTX<"83", PTX84>;
diff --git a/clang/test/CodeGen/builtins-nvptx.c
b/clang/test/CodeGen/builtins-nvptx.c
index 26c465eef306a0..ffa41c85c2734f 100644
--- a/clang/test/CodeGen/builtins-nvptx.c
+++ b/clang/test/CodeGen/builtins-nvptx.c
@@ -28,6 +28,11 @@
// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown
-target-cpu sm_89 -target-feature +ptx81 \
// RUN: -fcuda-is-device -emit-llvm -o - -x cuda %s \
// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX81_SM89 %s
+// ### The last run to check with the highest SM and PTX version available
+// ### to make sure target builtins are still accepted.
+// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown
-target-cpu sm_100a -target-feature +ptx87 \
+// RUN: -fcuda-is-device -emit-llvm -o - -x cuda %s \
+// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX81_SM89 %s
#define __device__ __attribute__((device))
#define __global__ __attribute__((global))
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
