[llvm] 0cbcf17 - [llvm][nvptx] Add sm_90a

[Guray Ozen via llvm-commits]

Author: Guray Ozen
Date: 2023-07-24T15:11:54+02:00
New Revision: 0cbcf17cb8be4a2c6f1b55e36cba31070c674045

URL: https://github.com/llvm/llvm-project/commit/0cbcf17cb8be4a2c6f1b55e36cba31070c674045
DIFF: https://github.com/llvm/llvm-project/commit/0cbcf17cb8be4a2c6f1b55e36cba31070c674045.diff

LOG: [llvm][nvptx] Add sm_90a

This works adds `sm_90a` as nvptx target. `sm_90a` is required to generate  wgmma and setmaxnreg instructions.

Here is information about "a" prefix in PTX document:
Target architectures with suffix “a”, such as sm_90a, include architecture-accelerated features that are supported on the specified architecture only, hence such targets do not follow the onion layer model. Therefore, PTX code generated for such targets cannot be run on later generation devices. Architecture-accelerated features can only be used with targets that support these features.

Reviewed By: tra

Differential Revision: https://reviews.llvm.org/D155851

Added: 
    

Modified: 
    llvm/lib/Target/NVPTX/NVPTX.td

Removed: 
    


################################################################################
diff  --git a/llvm/lib/Target/NVPTX/NVPTX.td b/llvm/lib/Target/NVPTX/NVPTX.td
index f867c531181420..02fa2a4ee81ec5 100644
--- a/llvm/lib/Target/NVPTX/NVPTX.td
+++ b/llvm/lib/Target/NVPTX/NVPTX.td
@@ -28,6 +28,8 @@ class FeatureSM<int version>:
    SubtargetFeature<"sm_"# version, "SmVersion",
                     "" # version,
                     "Target SM " # version>;
+def SM90a: FeatureSM<90>;
+
 class FeaturePTX<int version>:
    SubtargetFeature<"ptx"# version, "PTXVersion",
                     "" # version,
@@ -68,6 +70,7 @@ def : Proc<"sm_86", [SM86, PTX71]>;
 def : Proc<"sm_87", [SM87, PTX74]>;
 def : Proc<"sm_89", [SM89, PTX78]>;
 def : Proc<"sm_90", [SM90, PTX78]>;
+def : Proc<"sm_90a", [SM90a, PTX80]>;
 
 def NVPTXInstrInfo : InstrInfo {
 }