Skip to content

Commit 7053fdf

Browse files
authored
[SYCLomatic] Fix the crash issue in asm migration parser when building SYCLomatic with assert on (#2805)
Signed-off-by: chenwei.sun <chenwei.sun@intel.com>
1 parent 77472b1 commit 7053fdf

File tree

2 files changed

+12
-1
lines changed

2 files changed

+12
-1
lines changed

clang/lib/DPCT/RulesAsm/Parser/AsmParser.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -102,7 +102,7 @@ void InlineAsmParser::addBuiltinIdentifier() {
102102

103103
InlineAsmBuiltinType *
104104
InlineAsmContext::getBuiltinType(InlineAsmBuiltinType::TypeKind Kind) {
105-
assert(Kind > 0 && Kind < InlineAsmBuiltinType::NUM_TYPES && "Unknown Kind");
105+
assert(Kind >= 0 && Kind < InlineAsmBuiltinType::NUM_TYPES && "Unknown Kind");
106106
if (AsmBuiltinTypes[Kind])
107107
return AsmBuiltinTypes[Kind];
108108

clang/test/dpct/asm/optimize.cu

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -176,3 +176,14 @@ __device__ int test5() {
176176
MACRO(7, s32 == 0, asm("add.s32.sat %0, %1, %2;" : "=r"(s32) : "r"(s32x), "r"(0)));
177177
return 0;
178178
}
179+
180+
// CHECK: inline void foo() {
181+
// CHECK-NEXT: #if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
182+
// CHECK-NEXT: asm volatile(".reg .b8 byte0;\n");
183+
// CHECK-NEXT: #else
184+
// CHECK-NEXT: uint8_t byte0;
185+
// CHECK-NEXT: #endif
186+
inline __device__ void foo() {
187+
asm volatile(".reg .b8 byte0;\n");
188+
}
189+

0 commit comments

Comments
 (0)