This is an archive of the discontinued LLVM Phabricator instance.

Differential D97009

[CUDA] fix builtin constraints for PTX 7.2
ClosedPublic

Authored by tra on Feb 18 2021, 4:39 PM.

Details

Reviewers
yaxunl
Commits
rG1a368ae3b78d: [CUDA] fix builtin constraints for PTX 7.2
Summary

This fixes build issues w/ CUDA-11 introduced by https://reviews.llvm.org/D95974

Diff Detail

Event Timeline

tra created this revision.Feb 18 2021, 4:39 PM
Herald added subscribers: bixia, jholewinski. · View Herald TranscriptFeb 18 2021, 4:39 PM
tra requested review of this revision.Feb 18 2021, 4:39 PM
Herald added a project: Restricted Project. · View Herald TranscriptFeb 18 2021, 4:39 PM
tra updated this revision to Diff 324812.Feb 18 2021, 4:44 PM
tra edited the summary of this revision. (Show Details)

Updated the test.

yaxunl added inline comments.Feb 18 2021, 5:39 PM
clang/include/clang/Basic/BuiltinsNVPTX.def
744–745

need to pop PTX72 ?

tra updated this revision to Diff 324829.Feb 18 2021, 5:51 PM

pop the macro

tra marked an inline comment as done.Feb 18 2021, 5:51 PM
tra added inline comments.
clang/include/clang/Basic/BuiltinsNVPTX.def
744–745

Good catch. Done.

yaxunl accepted this revision.Feb 18 2021, 5:52 PM

LGTM. Thanks.

This revision is now accepted and ready to land.Feb 18 2021, 5:52 PM
Harbormaster completed remote builds in B89824: Diff 324810.Feb 18 2021, 7:15 PM
Harbormaster completed remote builds in B89825: Diff 324812.Feb 18 2021, 7:32 PM
Harbormaster completed remote builds in B89833: Diff 324829.Feb 18 2021, 8:21 PM
rupprecht added a subscriber: rupprecht.Feb 19 2021, 9:40 AM

This fixes a build error we're seeing, so I'd like to land this in a bit if that's OK.

Closed by commit rG1a368ae3b78d: [CUDA] fix builtin constraints for PTX 7.2 (authored by tra, committed by rupprecht). · Explain WhyFeb 19 2021, 10:00 AM
This revision was automatically updated to reflect the committed changes.
rupprecht added a commit: rG1a368ae3b78d: [CUDA] fix builtin constraints for PTX 7.2.

Revision Contents

PathSize
clang/
include/
clang/
Basic/
5 lines
test/
CodeGen/
5 lines

Diff 325024

clang/include/clang/Basic/BuiltinsNVPTX.def

Show All 32 Lines
#pragma push_macro("PTX60") #pragma push_macro("PTX60")
#pragma push_macro("PTX61") #pragma push_macro("PTX61")
#pragma push_macro("PTX63") #pragma push_macro("PTX63")
#pragma push_macro("PTX64") #pragma push_macro("PTX64")
#pragma push_macro("PTX65") #pragma push_macro("PTX65")
#pragma push_macro("PTX70") #pragma push_macro("PTX70")
#pragma push_macro("PTX71") #pragma push_macro("PTX71")
#define PTX71 "ptx71" #pragma push_macro("PTX72")
#define PTX72 "ptx72"
#define PTX71 "ptx71|" PTX72
#define PTX70 "ptx70|" PTX71 #define PTX70 "ptx70|" PTX71
#define PTX65 "ptx65|" PTX70 #define PTX65 "ptx65|" PTX70
#define PTX64 "ptx64|" PTX65 #define PTX64 "ptx64|" PTX65
#define PTX63 "ptx63|" PTX64 #define PTX63 "ptx63|" PTX64
#define PTX61 "ptx61|" PTX63 #define PTX61 "ptx61|" PTX63
#define PTX60 "ptx60|" PTX61 #define PTX60 "ptx60|" PTX61
#pragma push_macro("AND") #pragma push_macro("AND")
▲ Show 20 Lines • Show All 684 Lines • ▼ Show 20 Lines
#pragma pop_macro("SM_80") #pragma pop_macro("SM_80")
#pragma pop_macro("SM_86") #pragma pop_macro("SM_86")
#pragma pop_macro("PTX60") #pragma pop_macro("PTX60")
#pragma pop_macro("PTX61") #pragma pop_macro("PTX61")
#pragma pop_macro("PTX63") #pragma pop_macro("PTX63")
#pragma pop_macro("PTX64") #pragma pop_macro("PTX64")
#pragma pop_macro("PTX65") #pragma pop_macro("PTX65")
#pragma pop_macro("PTX70") #pragma pop_macro("PTX70")
#pragma pop_macro("PTX71") #pragma pop_macro("PTX71")
#pragma pop_macro("PTX72")
yaxunlUnsubmitted
Not Done
ReplyInline Actions

need to pop PTX72 ?

yaxunl: need to pop PTX72 ?
traAuthorUnsubmitted
Done
ReplyInline Actions

Good catch. Done.

tra: Good catch. Done.

clang/test/CodeGen/builtins-nvptx-sm_70.cu

// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_70 \ // RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_70 \
// RUN: -fcuda-is-device -target-feature +ptx60 \ // RUN: -fcuda-is-device -target-feature +ptx60 \
// RUN: -S -emit-llvm -o - -x cuda %s \ // RUN: -S -emit-llvm -o - -x cuda %s \
// RUN: | FileCheck -check-prefix=CHECK_M16 %s // RUN: | FileCheck -check-prefix=CHECK_M16 %s
// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_70 \ // RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_70 \
// RUN: -fcuda-is-device -target-feature +ptx61 -DPTX61 \ // RUN: -fcuda-is-device -target-feature +ptx61 -DPTX61 \
// RUN: -S -emit-llvm -o - -x cuda %s \ // RUN: -S -emit-llvm -o - -x cuda %s \
// RUN: | FileCheck -check-prefixes=CHECK_M16,CHECK_M32_M8 %s // RUN: | FileCheck -check-prefixes=CHECK_M16,CHECK_M32_M8 %s
// Make sure builtins still work with the latest combination of GPU & PTX.
// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_86 \
// RUN: -fcuda-is-device -target-feature +ptx72 -DPTX61 \
// RUN: -S -emit-llvm -o - -x cuda %s \
// RUN: | FileCheck -check-prefixes=CHECK_M16,CHECK_M32_M8 %s
// RUN: %clang_cc1 -triple nvptx-unknown-unknown -target-cpu sm_60 \ // RUN: %clang_cc1 -triple nvptx-unknown-unknown -target-cpu sm_60 \
// RUN: -DPTX61 -fcuda-is-device -S -o /dev/null -x cuda -verify=pre-sm_70 %s // RUN: -DPTX61 -fcuda-is-device -S -o /dev/null -x cuda -verify=pre-sm_70 %s
// RUN: %clang_cc1 -triple nvptx-unknown-unknown \ // RUN: %clang_cc1 -triple nvptx-unknown-unknown \
// RUN: -target-cpu sm_70 -target-feature +ptx60 \ // RUN: -target-cpu sm_70 -target-feature +ptx60 \
// RUN: -DPTX61 -fcuda-is-device -S -o /dev/null -x cuda -verify=pre-ptx61 %s // RUN: -DPTX61 -fcuda-is-device -S -o /dev/null -x cuda -verify=pre-ptx61 %s
#if !defined(CUDA_VERSION) #if !defined(CUDA_VERSION)
#define __device__ __attribute__((device)) #define __device__ __attribute__((device))
▲ Show 20 Lines • Show All 451 Lines • Show Last 20 Lines