[SYCLomatic] Relax the syntax restriction of asm builtin identifier (#1951)

Signed-off-by: Wang, Yihan <[email protected]>

Author: Yihan Wang

clang/lib/DPCT/Asm/AsmLexer.cpp

@@ -195,6 +195,10 @@ bool InlineAsmLexer::lexTokenInternal(InlineAsmToken &Result) {
     } else if (Char == '|') {
       ++CurPtr;
       Kind = asmtok::pipe;
+    } else if (isIdentifierContinue(Char)) {
+      // PTX allows % builtin-identifier.
+      Result.setFlag(InlineAsmToken::NeedsCleaning);
+      return lexIdentifierContinue(Result, CurPtr);
     } else {
       // error: unknown escaped characher.
       return false;

clang/test/dpct/asm/syntax.cu

@@ -17,6 +17,12 @@ __global__ void gpu_ptx(int *d_ptr, int length) {
       unsigned int WARP_SZ;
       // CHECK: laneid = item_ct1.get_sub_group().get_local_linear_id();
       asm("mov.u32 %0, %%laneid;" : "=r"(laneid));
+      
+      // CHECK: laneid = item_ct1.get_sub_group().get_local_linear_id();
+      asm("mov.u32 %0, %laneid;" : "=r"(laneid));
+
+      // CHECK: warpid = item_ct1.get_sub_group().get_group_linear_id();
+      asm("mov.u32 %0, %%warpid;" : "=r"(warpid));
 
       // CHECK: warpid = item_ct1.get_sub_group().get_group_linear_id();
       asm("mov.u32 %0, %%warpid;" : "=r"(warpid));