This is an archive of the discontinued LLVM Phabricator instance.

Differential D79449

[CUDA] Make NVVM builtins available with CUDA-11 & PTX6.5
ClosedPublic

Authored by tra on May 5 2020, 2:57 PM.

Details

Reviewers
timshen
Commits
rG844096b996a0: [CUDA] Make NVVM builtins available with CUDA-11/PTX6.5

Diff Detail

Event Timeline

tra created this revision.May 5 2020, 2:57 PM
Herald added a project: Restricted Project. · View Herald TranscriptMay 5 2020, 2:57 PM
Herald added subscribers: sanjoy.google, bixia, yaxunl, jholewinski. · View Herald Transcript
tra updated this revision to Diff 262236.May 5 2020, 3:07 PM

Test that the builtins work all the way to sm_80/ptx65

tra updated this revision to Diff 262237.May 5 2020, 3:09 PM

Actually added sm_80 predicate. It's not used on any builtins yet, but it makes
sense to add it right now along with ptx65.

timshen accepted this revision.May 5 2020, 3:40 PM
This revision is now accepted and ready to land.May 5 2020, 3:40 PM
Harbormaster completed remote builds in B55852: Diff 262236.May 5 2020, 3:40 PM
Harbormaster completed remote builds in B55851: Diff 262234.
Harbormaster completed remote builds in B55853: Diff 262237.May 5 2020, 4:12 PM
Closed by commit rG844096b996a0: [CUDA] Make NVVM builtins available with CUDA-11/PTX6.5 (authored by tra). · Explain WhyMay 5 2020, 4:13 PM
This revision was automatically updated to reflect the committed changes.
tra mentioned this in D79515: [CUDA] Enable existing builtins for PTX7.0 as well..May 6 2020, 12:31 PM

Revision Contents

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

Diff 262254

clang/include/clang/Basic/BuiltinsNVPTX.def

Show All 14 Lines
#if defined(BUILTIN) && !defined(TARGET_BUILTIN) #if defined(BUILTIN) && !defined(TARGET_BUILTIN)
# define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE) BUILTIN(ID, TYPE, ATTRS) # define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE) BUILTIN(ID, TYPE, ATTRS)
#endif #endif
#pragma push_macro("SM_70") #pragma push_macro("SM_70")
#pragma push_macro("SM_72") #pragma push_macro("SM_72")
#pragma push_macro("SM_75") #pragma push_macro("SM_75")
#define SM_75 "sm_75" #pragma push_macro("SM_80")
#define SM_80 "sm_80"
#define SM_75 "sm_75|" SM_80
#define SM_72 "sm_72|" SM_75 #define SM_72 "sm_72|" SM_75
#define SM_70 "sm_70|" SM_72 #define SM_70 "sm_70|" SM_72
#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")
#define PTX64 "ptx64" #pragma push_macro("PTX65")
#define PTX65 "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
// Special Registers // Special Registers
▲ Show 20 Lines • Show All 673 Lines • ▼ Show 20 Lines
#undef BUILTIN #undef BUILTIN
#undef TARGET_BUILTIN #undef TARGET_BUILTIN
#pragma pop_macro("AND") #pragma pop_macro("AND")
#pragma pop_macro("SM_60") #pragma pop_macro("SM_60")
#pragma pop_macro("SM_70") #pragma pop_macro("SM_70")
#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("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")

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: -fcuda-is-device -target-feature +ptx65 \
// 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