[lang] added __global__ storage specifier

include/triton/lang/token.h

@@ -142,7 +142,7 @@ public:
     STATIC,
     THREAD,	// _Thread_local
     AUTO,
-    REGISTER,
+    GLOBAL,
 
     // STORAGE CLASS SPECIFIER END
     BREAK,
@@ -236,7 +236,7 @@ public:
   bool IsIdentifier() const { return IDENTIFIER == tag_; }
   bool IsEOF() const { return tag_ == Token::END; }
   bool IsTypeSpecQual() const { return CONST <= tag_ && tag_ <= ENUM; }
-  bool IsDecl() const { return CONST <= tag_ && tag_ <= REGISTER; }
+  bool IsDecl() const { return CONST <= tag_ && tag_ <= GLOBAL; }
   static const char* Lexeme(int tag) {
     auto iter = tagLexemeMap_.find(tag);
     if (iter == tagLexemeMap_.end())

include/triton/lang/type.h

@@ -40,7 +40,7 @@ enum {
   S_STATIC = 0x04,
   S_THREAD = 0x08,
   S_AUTO = 0x10,
-  S_REGISTER = 0x20,
+  S_GLOBAL = 0x20,
 
   // Type specifier
   T_SIGNED = 0x40,

lib/lang/parser.cc

@@ -1000,6 +1000,10 @@ QualType Parser::ParseDeclSpec(int* storageSpec, int* funcSpec, int* alignSpec)
       EnsureAndSetStorageSpec(tok, storageSpec, S_EXTERN);
       break;
 
+    case Token::GLOBAL:
+      EnsureAndSetStorageSpec(tok, storageSpec, S_GLOBAL);
+      break;
+
     case Token::STATIC:
       if (!storageSpec)
         Error(tok, ERR_FUNC_SPEC);
@@ -1020,10 +1024,6 @@ QualType Parser::ParseDeclSpec(int* storageSpec, int* funcSpec, int* alignSpec)
       EnsureAndSetStorageSpec(tok, storageSpec, S_AUTO);
       break;
 
-    case Token::REGISTER:
-      EnsureAndSetStorageSpec(tok, storageSpec, S_REGISTER);
-      break;
-
     // Type qualifier
     case Token::CONST:    qualSpec |= Qualifier::CONST;    break;
     case Token::RESTRICT: qualSpec |= Qualifier::RESTRICT; break;

lib/lang/token.cc

@@ -7,6 +7,7 @@
 static MemPoolImp<Token> tokenPool;
 
 const std::unordered_map<std::string, int> Token::kwTypeMap_ {
+  { "__global__", Token::GLOBAL },
   { "auto", Token::AUTO },
   { "break", Token::BREAK },
   { "case", Token::CASE },
@@ -30,7 +31,6 @@ const std::unordered_map<std::string, int> Token::kwTypeMap_ {
   { "newaxis", Token::NEWAXIS },
   { "signed", Token::SIGNED },
   { "unsigned", Token::UNSIGNED },
-  { "register", Token::REGISTER },
   { "restrict", Token::RESTRICT },
   { "return", Token::RETURN },
   { "short", Token::SHORT },
@@ -121,6 +121,7 @@ const std::unordered_map<int, const char*> Token::tagLexemeMap_ {
   { Token::EXTERN, "extern" },
   { Token::FLOAT, "float" },
   { Token::FOR, "for" },
+  { Token::GLOBAL, "global" },
   { Token::GOTO, "goto" },
   { Token::IF, "if" },
   { Token::INLINE, "inline" },
@@ -129,7 +130,6 @@ const std::unordered_map<int, const char*> Token::tagLexemeMap_ {
   { Token::NEWAXIS, "newaxis" },
   { Token::SIGNED, "signed" },
   { Token::UNSIGNED, "unsigned" },
-  { Token::REGISTER, "register" },
   { Token::RESTRICT, "restrict" },
   { Token::RETURN, "return" },
   { Token::SHORT, "short" },

tests/common/src/dot.h

@@ -34,11 +34,11 @@ R"(
 #define SHAPE_B TK, TN
 #endif
 
-void dot(TYPE * A, TYPE * B, TYPE * C,
-         int M, int N, int K,
-         int lda __multipleof(8),
-         int ldb __multipleof(8),
-         int ldc) {
+__global__ void dot(TYPE * A, TYPE * B, TYPE * C,
+                    int M, int N, int K,
+                    int lda __multipleof(8),
+                    int ldb __multipleof(8),
+                    int ldc) {
   // prologue
   int ridx = get_program_id(0);
   int ridy = get_program_id(1);