Skip to content

[CUDA] Make target intrinsics work with ptx 8.7 #124818

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Jan 28, 2025
Merged

[CUDA] Make target intrinsics work with ptx 8.7 #124818

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 2 additions & 1 deletion clang/include/clang/Basic/BuiltinsNVPTX.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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>;
Expand Down
5 changes: 5 additions & 0 deletions clang/test/CodeGen/builtins-nvptx.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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))
Expand Down