Skip to content

Commit 310f558

Browse files
authored
[CUDA] Make target intrinsics work with ptx 8.7 (#124818)
Fixes build break with CUDA-12.8 introduced in #123398
1 parent 964565c commit 310f558

File tree

2 files changed

+7
-1
lines changed

2 files changed

+7
-1
lines changed

clang/include/clang/Basic/BuiltinsNVPTX.td

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -48,8 +48,9 @@ class PTX<string version, PTXFeatures newer> : PTXFeatures {
4848
let Features = !strconcat("ptx", version, "|", newer.Features);
4949
}
5050

51-
let Features = "ptx86" in def PTX86 : PTXFeatures;
51+
let Features = "ptx87" in def PTX87 : PTXFeatures;
5252

53+
def PTX86 : PTX<"86", PTX87>;
5354
def PTX85 : PTX<"85", PTX86>;
5455
def PTX84 : PTX<"84", PTX85>;
5556
def PTX83 : PTX<"83", PTX84>;

clang/test/CodeGen/builtins-nvptx.c

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -28,6 +28,11 @@
2828
// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_89 -target-feature +ptx81 \
2929
// RUN: -fcuda-is-device -emit-llvm -o - -x cuda %s \
3030
// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX81_SM89 %s
31+
// ### The last run to check with the highest SM and PTX version available
32+
// ### to make sure target builtins are still accepted.
33+
// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_100a -target-feature +ptx87 \
34+
// RUN: -fcuda-is-device -emit-llvm -o - -x cuda %s \
35+
// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX81_SM89 %s
3136

3237
#define __device__ __attribute__((device))
3338
#define __global__ __attribute__((global))

0 commit comments

Comments
 (0)