This is an archive of the discontinued LLVM Phabricator instance.

Differential D151361

[CUDA] bump supported CUDA version to 12.1/11.8
ClosedPublic

Authored by tra on May 24 2023, 12:17 PM.

Details

Reviewers
jlebar
Commits
rGffb635cb2d4e: [CUDA] bump supported CUDA version to 12.1/11.8

Diff Detail

Event Timeline

tra created this revision.May 24 2023, 12:17 PM
Herald added a project: Restricted Project. · View Herald TranscriptMay 24 2023, 12:17 PM
Herald added subscribers: mattd, gchakrabarti, asavonic and 3 others. · View Herald Transcript
tra added a child revision: D151009: [NVPTX] add new sm90-specific intrinsics..May 24 2023, 12:18 PM
tra published this revision for review.May 24 2023, 12:24 PM
tra added a reviewer: jlebar.
Herald added projects: Restricted Project, Restricted Project. · View Herald TranscriptMay 24 2023, 12:24 PM
Herald added subscribers: llvm-commits, cfe-commits, MaskRay, jholewinski. · View Herald Transcript
tra added a parent revision: D151359: [CUDA] Relax restrictions on variadics in host-side compilation..May 24 2023, 2:35 PM
jlebar accepted this revision.May 24 2023, 3:16 PM
This revision is now accepted and ready to land.May 24 2023, 3:16 PM
Harbormaster completed remote builds in B234288: Diff 525291.May 24 2023, 5:55 PM
Closed by commit rGffb635cb2d4e: [CUDA] bump supported CUDA version to 12.1/11.8 (authored by tra). · Explain WhyMay 25 2023, 11:58 AM
This revision was automatically updated to reflect the committed changes.
tra added a commit: rGffb635cb2d4e: [CUDA] bump supported CUDA version to 12.1/11.8.
jhuber6 mentioned this in D151735: [libc] Implement basic `malloc` and `free` support on the GPU.May 30 2023, 11:02 AM
bader added a subscriber: bader.Jun 15 2023, 2:19 PM
bader added inline comments.
clang/docs/ReleaseNotes.rst
590

@tra, could you update llvm/docs/CompileCudaWithLLVM.rst as well, please?

tra added inline comments.Jun 15 2023, 2:28 PM
clang/docs/ReleaseNotes.rst
590

Done in d028188412fa54774e2c60e21f0929a0fede93bb

bader added inline comments.Jun 15 2023, 2:44 PM
clang/docs/ReleaseNotes.rst
590

Great. Thank you!

Revision Contents

PathSize
clang/
docs/
1 line
include/
clang/
Basic/
8 lines
6 lines
lib/
Basic/
2 lines
Driver/
ToolChains/
6 lines
llvm/
lib/
Target/
NVPTX/
97 lines

Diff 525735

clang/docs/ReleaseNotes.rst

Show First 20 Lines • Show All 581 Lines • ▼ Show 20 Lines- Added ``attribute(riscv_rvv_vector_bits(__riscv_v_fixed_vlen))`` to allow
the size of a RVV (RISC-V Vector) scalable type to be specified. This allows the size of a RVV (RISC-V Vector) scalable type to be specified. This allows
RVV scalable vector types to be used in structs or in global variables. RVV scalable vector types to be used in structs or in global variables.
CUDA/HIP Language Changes CUDA/HIP Language Changes
^^^^^^^^^^^^^^^^^^^^^^^^^ ^^^^^^^^^^^^^^^^^^^^^^^^^
CUDA Support CUDA Support
^^^^^^^^^^^^ ^^^^^^^^^^^^
- Clang now supports CUDA SDK up to 12.1
baderUnsubmitted
Not Done
ReplyInline Actions

@tra, could you update llvm/docs/CompileCudaWithLLVM.rst as well, please?

bader: @tra, could you update llvm/docs/CompileCudaWithLLVM.rst as well, please?
traAuthorUnsubmitted
Done
ReplyInline Actions

Done in d028188412fa54774e2c60e21f0929a0fede93bb

tra: Done in d028188412fa54774e2c60e21f0929a0fede93bb
baderUnsubmitted
Not Done
ReplyInline Actions

Great. Thank you!

bader: Great. Thank you!
AIX Support AIX Support
^^^^^^^^^^^ ^^^^^^^^^^^
- Add an AIX-only link-time option, `-mxcoff-build-id=0xHEXSTRING`, to allow users - Add an AIX-only link-time option, `-mxcoff-build-id=0xHEXSTRING`, to allow users
to embed a hex id in their binary such that it's readable by the program itself. to embed a hex id in their binary such that it's readable by the program itself.
This option is an alternative to the `--build-id=0xHEXSTRING` GNU linker option This option is an alternative to the `--build-id=0xHEXSTRING` GNU linker option
which is currently not supported by the AIX linker. which is currently not supported by the AIX linker.
▲ Show 20 Lines • Show All 100 Lines • Show Last 20 Lines

clang/include/clang/Basic/BuiltinsNVPTX.def

Show First 20 Lines • Show All 48 Lines • ▼ Show 20 Lines
#pragma push_macro("PTX71") #pragma push_macro("PTX71")
#pragma push_macro("PTX72") #pragma push_macro("PTX72")
#pragma push_macro("PTX73") #pragma push_macro("PTX73")
#pragma push_macro("PTX74") #pragma push_macro("PTX74")
#pragma push_macro("PTX75") #pragma push_macro("PTX75")
#pragma push_macro("PTX76") #pragma push_macro("PTX76")
#pragma push_macro("PTX77") #pragma push_macro("PTX77")
#pragma push_macro("PTX78") #pragma push_macro("PTX78")
#define PTX78 "ptx78" #pragma push_macro("PTX80")
#pragma push_macro("PTX81")
#define PTX81 "ptx81"
#define PTX80 "ptx80|" PTX81
#define PTX78 "ptx78|" PTX80
#define PTX77 "ptx77|" PTX78 #define PTX77 "ptx77|" PTX78
#define PTX76 "ptx76|" PTX77 #define PTX76 "ptx76|" PTX77
#define PTX75 "ptx75|" PTX76 #define PTX75 "ptx75|" PTX76
#define PTX74 "ptx74|" PTX75 #define PTX74 "ptx74|" PTX75
#define PTX73 "ptx73|" PTX74 #define PTX73 "ptx73|" PTX74
#define PTX72 "ptx72|" PTX73 #define PTX72 "ptx72|" PTX73
#define PTX71 "ptx71|" PTX72 #define PTX71 "ptx71|" PTX72
#define PTX70 "ptx70|" PTX71 #define PTX70 "ptx70|" PTX71
▲ Show 20 Lines • Show All 941 Lines • ▼ Show 20 Lines
#pragma pop_macro("PTX71") #pragma pop_macro("PTX71")
#pragma pop_macro("PTX72") #pragma pop_macro("PTX72")
#pragma pop_macro("PTX73") #pragma pop_macro("PTX73")
#pragma pop_macro("PTX74") #pragma pop_macro("PTX74")
#pragma pop_macro("PTX75") #pragma pop_macro("PTX75")
#pragma pop_macro("PTX76") #pragma pop_macro("PTX76")
#pragma pop_macro("PTX77") #pragma pop_macro("PTX77")
#pragma pop_macro("PTX78") #pragma pop_macro("PTX78")
#pragma pop_macro("PTX80")
#pragma pop_macro("PTX81")

clang/include/clang/Basic/Cuda.h

Show All 31 Linesenum class CudaVersion {
CUDA_111, CUDA_111,
CUDA_112, CUDA_112,
CUDA_113, CUDA_113,
CUDA_114, CUDA_114,
CUDA_115, CUDA_115,
CUDA_116, CUDA_116,
CUDA_117, CUDA_117,
CUDA_118, CUDA_118,
FULLY_SUPPORTED = CUDA_115, CUDA_120,
CUDA_121,
FULLY_SUPPORTED = CUDA_118,
PARTIALLY_SUPPORTED = PARTIALLY_SUPPORTED =
CUDA_118, // Partially supported. Proceed with a warning. CUDA_121, // Partially supported. Proceed with a warning.
NEW = 10000, // Too new. Issue a warning, but allow using it. NEW = 10000, // Too new. Issue a warning, but allow using it.
}; };
const char *CudaVersionToString(CudaVersion V); const char *CudaVersionToString(CudaVersion V);
// Input is "Major.Minor" // Input is "Major.Minor"
CudaVersion CudaStringToVersion(const llvm::Twine &S); CudaVersion CudaStringToVersion(const llvm::Twine &S);
enum class CudaArch { enum class CudaArch {
UNUSED, UNUSED,
▲ Show 20 Lines • Show All 105 Lines • Show Last 20 Lines

clang/lib/Basic/Cuda.cpp

Show All 31 Linesstatic const CudaVersionMapEntry CudaNameVersionMap[] = {
CUDA_ENTRY(11, 1), CUDA_ENTRY(11, 1),
CUDA_ENTRY(11, 2), CUDA_ENTRY(11, 2),
CUDA_ENTRY(11, 3), CUDA_ENTRY(11, 3),
CUDA_ENTRY(11, 4), CUDA_ENTRY(11, 4),
CUDA_ENTRY(11, 5), CUDA_ENTRY(11, 5),
CUDA_ENTRY(11, 6), CUDA_ENTRY(11, 6),
CUDA_ENTRY(11, 7), CUDA_ENTRY(11, 7),
CUDA_ENTRY(11, 8), CUDA_ENTRY(11, 8),
CUDA_ENTRY(12, 0),
CUDA_ENTRY(12, 1),
{"", CudaVersion::NEW, llvm::VersionTuple(std::numeric_limits<int>::max())}, {"", CudaVersion::NEW, llvm::VersionTuple(std::numeric_limits<int>::max())},
{"unknown", CudaVersion::UNKNOWN, {}} // End of list tombstone. {"unknown", CudaVersion::UNKNOWN, {}} // End of list tombstone.
}; };
#undef CUDA_ENTRY #undef CUDA_ENTRY
const char *CudaVersionToString(CudaVersion V) { const char *CudaVersionToString(CudaVersion V) {
for (auto *I = CudaNameVersionMap; I->Version != CudaVersion::UNKNOWN; ++I) for (auto *I = CudaNameVersionMap; I->Version != CudaVersion::UNKNOWN; ++I)
if (I->Version == V) if (I->Version == V)
▲ Show 20 Lines • Show All 195 Lines • Show Last 20 Lines

clang/lib/Driver/ToolChains/Cuda.cpp

Show First 20 Lines • Show All 68 Lines • ▼ Show 20 LinesCudaVersion getCudaVersion(uint32_t raw_version) {
if (raw_version < 11060) if (raw_version < 11060)
return CudaVersion::CUDA_115; return CudaVersion::CUDA_115;
if (raw_version < 11070) if (raw_version < 11070)
return CudaVersion::CUDA_116; return CudaVersion::CUDA_116;
if (raw_version < 11080) if (raw_version < 11080)
return CudaVersion::CUDA_117; return CudaVersion::CUDA_117;
if (raw_version < 11090) if (raw_version < 11090)
return CudaVersion::CUDA_118; return CudaVersion::CUDA_118;
if (raw_version < 12010)
return CudaVersion::CUDA_120;
if (raw_version < 12020)
return CudaVersion::CUDA_121;
return CudaVersion::NEW; return CudaVersion::NEW;
} }
CudaVersion parseCudaHFile(llvm::StringRef Input) { CudaVersion parseCudaHFile(llvm::StringRef Input) {
// Helper lambda which skips the words if the line starts with them or returns // Helper lambda which skips the words if the line starts with them or returns
// std::nullopt otherwise. // std::nullopt otherwise.
auto StartsWithWords = auto StartsWithWords =
[](llvm::StringRef Line, [](llvm::StringRef Line,
▲ Show 20 Lines • Show All 578 Lines • ▼ Show 20 Linesvoid NVPTX::getNVPTXTargetFeatures(const Driver &D, const llvm::Triple &Triple,
// by new PTX version, so we need to raise PTX level to enable them in NVPTX // by new PTX version, so we need to raise PTX level to enable them in NVPTX
// back-end. // back-end.
const char *PtxFeature = nullptr; const char *PtxFeature = nullptr;
switch (CudaInstallation.version()) { switch (CudaInstallation.version()) {
#define CASE_CUDA_VERSION(CUDA_VER, PTX_VER) \ #define CASE_CUDA_VERSION(CUDA_VER, PTX_VER) \
case CudaVersion::CUDA_##CUDA_VER: \ case CudaVersion::CUDA_##CUDA_VER: \
PtxFeature = "+ptx" #PTX_VER; \ PtxFeature = "+ptx" #PTX_VER; \
break; break;
CASE_CUDA_VERSION(121, 81);
CASE_CUDA_VERSION(120, 80);
CASE_CUDA_VERSION(118, 78); CASE_CUDA_VERSION(118, 78);
CASE_CUDA_VERSION(117, 77); CASE_CUDA_VERSION(117, 77);
CASE_CUDA_VERSION(116, 76); CASE_CUDA_VERSION(116, 76);
CASE_CUDA_VERSION(115, 75); CASE_CUDA_VERSION(115, 75);
CASE_CUDA_VERSION(114, 74); CASE_CUDA_VERSION(114, 74);
CASE_CUDA_VERSION(113, 73); CASE_CUDA_VERSION(113, 73);
CASE_CUDA_VERSION(112, 72); CASE_CUDA_VERSION(112, 72);
CASE_CUDA_VERSION(111, 71); CASE_CUDA_VERSION(111, 71);
▲ Show 20 Lines • Show All 354 Lines • Show Last 20 Lines

llvm/lib/Target/NVPTX/NVPTX.td

Show All 18 Lines
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// Subtarget Features. // Subtarget Features.
// - We use the SM version number instead of explicit feature table. // - We use the SM version number instead of explicit feature table.
// - Need at least one feature to avoid generating zero sized array by // - Need at least one feature to avoid generating zero sized array by
// TableGen in NVPTXGenSubtarget.inc. // TableGen in NVPTXGenSubtarget.inc.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// SM Versions class FeatureSM<int version>:
def SM20 : SubtargetFeature<"sm_20", "SmVersion", "20", SubtargetFeature<"sm_"# version, "SmVersion",
"Target SM 2.0">; "" # version,
def SM21 : SubtargetFeature<"sm_21", "SmVersion", "21", "Target SM " # version>;
"Target SM 2.1">; class FeaturePTX<int version>:
def SM30 : SubtargetFeature<"sm_30", "SmVersion", "30", SubtargetFeature<"ptx"# version, "PTXVersion",
"Target SM 3.0">; "" # version,
def SM32 : SubtargetFeature<"sm_32", "SmVersion", "32", "Use PTX version " # version>;
"Target SM 3.2">;
def SM35 : SubtargetFeature<"sm_35", "SmVersion", "35", foreach version = [20, 21, 30, 32, 35, 37, 50, 52, 53,
"Target SM 3.5">; 60, 61, 62, 70, 72, 75, 80, 86, 87, 89, 90] in
def SM37 : SubtargetFeature<"sm_37", "SmVersion", "37", def SM#version: FeatureSM<version>;
"Target SM 3.7">;
def SM50 : SubtargetFeature<"sm_50", "SmVersion", "50", foreach version = [32, 40, 41, 42, 43, 50, 60, 61, 63, 64, 65,
"Target SM 5.0">; 70, 71, 72, 73, 74, 75, 76, 77, 78, 80, 81] in
def SM52 : SubtargetFeature<"sm_52", "SmVersion", "52", def PTX#version: FeaturePTX<version>;
"Target SM 5.2">;
def SM53 : SubtargetFeature<"sm_53", "SmVersion", "53",
"Target SM 5.3">;
def SM60 : SubtargetFeature<"sm_60", "SmVersion", "60",
"Target SM 6.0">;
def SM61 : SubtargetFeature<"sm_61", "SmVersion", "61",
"Target SM 6.1">;
def SM62 : SubtargetFeature<"sm_62", "SmVersion", "62",
"Target SM 6.2">;
def SM70 : SubtargetFeature<"sm_70", "SmVersion", "70",
"Target SM 7.0">;
def SM72 : SubtargetFeature<"sm_72", "SmVersion", "72",
"Target SM 7.2">;
def SM75 : SubtargetFeature<"sm_75", "SmVersion", "75",
"Target SM 7.5">;
def SM80 : SubtargetFeature<"sm_80", "SmVersion", "80",
"Target SM 8.0">;
def SM86 : SubtargetFeature<"sm_86", "SmVersion", "86",
"Target SM 8.6">;
def SM87 : SubtargetFeature<"sm_87", "SmVersion", "87",
"Target SM 8.7">;
def SM89 : SubtargetFeature<"sm_89", "SmVersion", "89",
"Target SM 8.9">;
def SM90 : SubtargetFeature<"sm_90", "SmVersion", "90",
"Target SM 9.0">;
// PTX Versions
def PTX32 : SubtargetFeature<"ptx32", "PTXVersion", "32",
"Use PTX version 3.2">;
def PTX40 : SubtargetFeature<"ptx40", "PTXVersion", "40",
"Use PTX version 4.0">;
def PTX41 : SubtargetFeature<"ptx41", "PTXVersion", "41",
"Use PTX version 4.1">;
def PTX42 : SubtargetFeature<"ptx42", "PTXVersion", "42",
"Use PTX version 4.2">;
def PTX43 : SubtargetFeature<"ptx43", "PTXVersion", "43",
"Use PTX version 4.3">;
def PTX50 : SubtargetFeature<"ptx50", "PTXVersion", "50",
"Use PTX version 5.0">;
def PTX60 : SubtargetFeature<"ptx60", "PTXVersion", "60",
"Use PTX version 6.0">;
def PTX61 : SubtargetFeature<"ptx61", "PTXVersion", "61",
"Use PTX version 6.1">;
def PTX63 : SubtargetFeature<"ptx63", "PTXVersion", "63",
"Use PTX version 6.3">;
def PTX64 : SubtargetFeature<"ptx64", "PTXVersion", "64",
"Use PTX version 6.4">;
def PTX65 : SubtargetFeature<"ptx65", "PTXVersion", "65",
"Use PTX version 6.5">;
def PTX70 : SubtargetFeature<"ptx70", "PTXVersion", "70",
"Use PTX version 7.0">;
def PTX71 : SubtargetFeature<"ptx71", "PTXVersion", "71",
"Use PTX version 7.1">;
def PTX72 : SubtargetFeature<"ptx72", "PTXVersion", "72",
"Use PTX version 7.2">;
def PTX73 : SubtargetFeature<"ptx73", "PTXVersion", "73",
"Use PTX version 7.3">;
def PTX74 : SubtargetFeature<"ptx74", "PTXVersion", "74",
"Use PTX version 7.4">;
def PTX75 : SubtargetFeature<"ptx75", "PTXVersion", "75",
"Use PTX version 7.5">;
def PTX76 : SubtargetFeature<"ptx76", "PTXVersion", "76",
"Use PTX version 7.6">;
def PTX77 : SubtargetFeature<"ptx77", "PTXVersion", "77",
"Use PTX version 7.7">;
def PTX78 : SubtargetFeature<"ptx78", "PTXVersion", "78",
"Use PTX version 7.8">;
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// NVPTX supported processors. // NVPTX supported processors.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
class Proc<string Name, list<SubtargetFeature> Features> class Proc<string Name, list<SubtargetFeature> Features>
: Processor<Name, NoItineraries, Features>; : Processor<Name, NoItineraries, Features>;
Show All 27 Lines