This is an archive of the discontinued LLVM Phabricator instance.

Differential D79515

[CUDA] Enable existing builtins for PTX7.0 as well.
ClosedPublic

Authored by tra on May 6 2020, 12:31 PM.

Details

Reviewers
timshen
Commits
rG314f99e7d42d: [CUDA] Enable existing builtins for PTX7.0 as well.
Summary

I've missed PXT7.0 in D79449

Diff Detail

Event Timeline

tra created this revision.May 6 2020, 12:31 PM
Herald added a project: Restricted Project. · View Herald TranscriptMay 6 2020, 12:31 PM
Herald added subscribers: sanjoy.google, bixia, yaxunl, jholewinski. · View Herald Transcript
timshen accepted this revision.May 6 2020, 1:34 PM
timshen added inline comments.
clang/test/CodeGen/builtins-nvptx-ptx60.cu
6

Keep ptx65 as well?

This revision is now accepted and ready to land.May 6 2020, 1:34 PM
Harbormaster completed remote builds in B55969: Diff 262453.May 6 2020, 1:35 PM
tra updated this revision to Diff 262473.May 6 2020, 2:10 PM

Updates test.

tra marked an inline comment as done.May 6 2020, 2:10 PM
Harbormaster completed remote builds in B55979: Diff 262473.May 6 2020, 2:44 PM
Closed by commit rG314f99e7d42d: [CUDA] Enable existing builtins for PTX7.0 as well. (authored by tra). · Explain WhyMay 6 2020, 2:44 PM
This revision was automatically updated to reflect the committed changes.

Revision Contents

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

Diff 262482

clang/include/clang/Basic/BuiltinsNVPTX.def

Show All 28 Lines
#pragma push_macro("SM_60") #pragma push_macro("SM_60")
#define SM_60 "sm_60|sm_61|sm_62|" SM_70 #define SM_60 "sm_60|sm_61|sm_62|" SM_70
#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")
#define PTX65 "ptx65" #pragma push_macro("PTX70")
#define PTX70 "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")
#define AND(a, b) a "," b #define AND(a, b) a "," b
▲ Show 20 Lines • Show All 680 Lines • ▼ Show 20 Lines
#pragma pop_macro("SM_72") #pragma pop_macro("SM_72")
#pragma pop_macro("SM_75") #pragma pop_macro("SM_75")
#pragma pop_macro("SM_80") #pragma pop_macro("SM_80")
#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")

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

// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_60 \ // RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_60 \
// 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 %s // RUN: | FileCheck -check-prefix=CHECK %s
// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_80 \ // RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_80 \
// RUN: -fcuda-is-device -target-feature +ptx65 \ // RUN: -fcuda-is-device -target-feature +ptx65 \
timshenUnsubmitted
Done
ReplyInline Actions

Keep ptx65 as well?

timshen: Keep ptx65 as well?
// RUN: -S -emit-llvm -o - -x cuda %s \ // RUN: -S -emit-llvm -o - -x cuda %s \
// RUN: | FileCheck -check-prefix=CHECK %s // RUN: | FileCheck -check-prefix=CHECK %s
// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_80 \
// RUN: -fcuda-is-device -target-feature +ptx70 \
// RUN: -S -emit-llvm -o - -x cuda %s \
// RUN: | FileCheck -check-prefix=CHECK %s
// RUN: %clang_cc1 -triple nvptx-unknown-unknown -target-cpu sm_60 \ // RUN: %clang_cc1 -triple nvptx-unknown-unknown -target-cpu sm_60 \
// RUN: -fcuda-is-device -S -o /dev/null -x cuda -verify %s // RUN: -fcuda-is-device -S -o /dev/null -x cuda -verify %s
#define __device__ __attribute__((device)) #define __device__ __attribute__((device))
#define __global__ __attribute__((global)) #define __global__ __attribute__((global))
#define __shared__ __attribute__((shared)) #define __shared__ __attribute__((shared))
#define __constant__ __attribute__((constant)) #define __constant__ __attribute__((constant))
▲ Show 20 Lines • Show All 85 Lines • Show Last 20 Lines