Skip to content

Commit 0d15753

Browse files
committed
Handle CUDA inline hint in ast_canopy
1 parent c7820d8 commit 0d15753

3 files changed

Lines changed: 38 additions & 1 deletion

File tree

ast_canopy/ast_canopy/api.py

Lines changed: 9 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -33,6 +33,11 @@
3333
"-nocudalib",
3434
]
3535

36+
# CUDA accepts this inline hint spelling, but Clang does not predefine it.
37+
_CUDA_PARSER_COMPAT_DEFINES: list[str] = [
38+
"__inline_hint__=",
39+
]
40+
3641

3742
def _get_shim_include_dir() -> str:
3843
"""Return the absolute path to the local shim include directory"""
@@ -477,7 +482,10 @@ def parse_declarations_from_source(
477482

478483
cuda_wrappers_dir = get_cuda_wrappers_include_dir(clang_resource_dir)
479484

480-
define_flags = [f"-D{define}" for define in defines]
485+
define_flags = [
486+
f"-D{define}"
487+
for define in [*_CUDA_PARSER_COMPAT_DEFINES, *defines]
488+
]
481489

482490
# The include paths are ordered a below:
483491
# 1. clang resource file include directory (via -isystem flag)
Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,8 @@
1+
// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
2+
// SPDX-License-Identifier: Apache-2.0
3+
4+
#define NVSHMEMI_DEVICE_INLINE __inline_hint__
5+
6+
__device__ NVSHMEMI_DEVICE_INLINE void inline_hint_void() {}
7+
8+
__device__ NVSHMEMI_DEVICE_INLINE int inline_hint_int() { return 1; }
Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,21 @@
1+
# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
2+
# SPDX-License-Identifier: Apache-2.0
3+
4+
import pytest
5+
6+
from ast_canopy import parse_declarations_from_source
7+
8+
9+
@pytest.fixture(scope="module")
10+
def sample_inline_hint_macro(data_folder):
11+
return data_folder / "inline_hint_macro_fix.cu"
12+
13+
14+
def test_inline_hint_macro(sample_inline_hint_macro):
15+
srcstr = str(sample_inline_hint_macro)
16+
17+
decls = parse_declarations_from_source(srcstr, [srcstr], "sm_80")
18+
functions = {func.name: func for func in decls.functions}
19+
20+
assert functions["inline_hint_void"].return_type.name == "void"
21+
assert functions["inline_hint_int"].return_type.name == "int"

0 commit comments

Comments
 (0)