Skip to content

Enable Metal in pointer test suite #958

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Draft
jhelferty-nv wants to merge 9 commits into
base: main
Choose a base branch
from
Draft

Enable Metal in pointer test suite #958

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
9 commits
Select commit Hold shift + click to select a range
8ed4ccf
Add SGL_LOCAL_RHI cmake option to use a local slang-rhi checkout
jhelferty-nv Apr 16, 2026
365dabe
Enable Metal in pointer test suite
jhelferty-nv Apr 25, 2026
d0aabdb
Disable `&` pointer tests
jhelferty-nv Apr 23, 2026
f07f5c3
Update slang-rhi submodule for Metal pointer residency support
jhelferty-nv Apr 27, 2026
0f11f8c
Add ResidencySet to Feature enum to match slang-rhi
jhelferty-nv Apr 30, 2026
e5239d7
Unskip Metal parameter block test
jhelferty-nv Apr 30, 2026
fd774f5
Fix to avoid argument order diagnostic
jhelferty-nv May 25, 2026
416a564
Realigned externals/lslang-rhi with origin/main
jhelferty-nv May 25, 2026
299a53b
Merge remote-tracking branch 'origin/main' into enable-metal-ptr
jhelferty-nv May 25, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 2 additions & 0 deletions AGENTS.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -174,6 +174,8 @@ cmake --build --preset windows-msvc-debug
| `SGL_LOCAL_SLANG` | OFF | Enable to use a local Slang build |
| `SGL_LOCAL_SLANG_DIR` | `../slang` | Path to the local Slang repository |
| `SGL_LOCAL_SLANG_BUILD_DIR` | `build/Debug` | Build directory within the Slang repo |
| `SGL_LOCAL_RHI` | OFF | Use a local slang-rhi instead of the submodule |
| `SGL_LOCAL_RHI_DIR` | `../slang-rhi` | Path to the local slang-rhi repository |

### Workflow

Expand Down
13 changes: 11 additions & 2 deletions external/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -215,6 +215,9 @@ install(FILES ${SLANG_DIR}/include/slang-image-format-defs.h DESTINATION ${CMAKE

# slang-rhi

set(SGL_LOCAL_RHI OFF CACHE BOOL "Use a local slang-rhi instead of the submodule")
set(SGL_LOCAL_RHI_DIR "${CMAKE_SOURCE_DIR}/../slang-rhi" CACHE PATH "Path to local slang-rhi")

set(SLANG_RHI_BUILD_TESTS OFF)
set(SLANG_RHI_BUILD_EXAMPLES OFF)
set(SLANG_RHI_FETCH_SLANG OFF)
Expand All @@ -226,10 +229,16 @@ set(SLANG_RHI_BUILD_FROM_SLANG_REPO ON)
set(SLANG_RHI_FETCH_SLANG OFF)
set(SLANG_RHI_SLANG_INCLUDE_DIR ${SLANG_INCLUDE_DIR} CACHE STRING "" FORCE)
set(SLANG_RHI_SLANG_BINARY_DIR ${SLANG_DIR} CACHE STRING "" FORCE)
add_subdirectory(slang-rhi)
if(SGL_LOCAL_RHI)
add_subdirectory(${SGL_LOCAL_RHI_DIR} ${CMAKE_BINARY_DIR}/slang-rhi)
set(SGL_RHI_SOURCE_DIR ${SGL_LOCAL_RHI_DIR})
else()
add_subdirectory(slang-rhi)
set(SGL_RHI_SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/slang-rhi)
endif()
Comment on lines 218 to +238
Copy link
Copy Markdown
Contributor

@skallweitNV skallweitNV Apr 30, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I would probably rename these to SGL_LOCAL_SLANG_RHI and SGL_LOCAL_SLANG_RHI_DIR to be more consistent with SGL_LOCAL_SLANG. Also rename SGL_RHI_SOURCE_DIR to SGL_SLANG_RHI_SOURCE_DIR


# Copy shared slang-rhi-device.h header to shaders directory (wrapped in sgl/device/rhi.slang)
sgl_copy_shader_file(${CMAKE_CURRENT_SOURCE_DIR}/slang-rhi/include/slang-rhi-device.h slang-rhi)
sgl_copy_shader_file(${SGL_RHI_SOURCE_DIR}/include/slang-rhi-device.h slang-rhi)

if(TARGET slang-rhi-nvapi)
get_target_property(NVAPI_DIR slang-rhi-nvapi INTERFACE_INCLUDE_DIRECTORIES)
Expand Down
2 changes: 0 additions & 2 deletions slangpy/tests/device/test_parameter_block.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,8 +13,6 @@
# to test the ParameterBlock binding.
@pytest.mark.parametrize("device_type", helpers.DEFAULT_DEVICE_TYPES)
def test_parameter_block(device_type: spy.DeviceType):
if device_type == spy.DeviceType.metal:
pytest.skip("Crash in slang-rhi due to invalid reflection data")

# Create device
device = helpers.get_device(type=device_type)
Expand Down
8 changes: 5 additions & 3 deletions slangpy/tests/slangpy_tests/test_pointers.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,10 +11,10 @@
from typing import Any, cast

# Filter default device types to only include those that support pointers
# TODO: Metal does support pointers but the is a slang bug leading to incorrect Metal shader code
# https://github.com/shader-slang/slang/issues/7605
POINTER_DEVICE_TYPES = [
x for x in helpers.DEFAULT_DEVICE_TYPES if x in [DeviceType.vulkan, DeviceType.cuda]
x
for x in helpers.DEFAULT_DEVICE_TYPES
if x in [DeviceType.vulkan, DeviceType.cuda, DeviceType.metal]
]


Expand All @@ -34,7 +34,7 @@
""",
)

res = function(100)

Check failure on line 37 in slangpy/tests/slangpy_tests/test_pointers.py

View workflow job for this annotation

GitHub Actions / build (macos, aarch64, clang, Release, 3.10) 

test_copy_pointer[DeviceType.metal]

RuntimeError: m_rhi_command_encoder->finish(rhi_command_buffer.writeRef()) failed with error: -2147467259 (SLANG_FAIL)

RHI messages:
slang: metal 32023.620: 1c7d72fbb205f30c2e60028a63056c511bb2216d(33): error :  type 'const constant CallData_default_0 *' is not valid for attribute 'buffer'
[[kernel]] void compute_main(uint3 sv_groupthreadid_0 [[thread_position_in_threadgroup]], uint3 flat_call_thread_id_0 [[thread_position_in_grid]], uint3 flat_call_group_id_0 [[threadgroup_position_in_grid]], CallData_default_0 constant* call_data_1 [[buffer(0)]])
                                                                                                                                                                                                                ^                                          ~~~~~~~~~
metal 32023.620: /Users/runner/work/slangpy/slangpy/slangpy/slang/core.slang(408): note :  type 'device int *device *' cannot be used in buffer pointee type
    int device* device* value_0;
    ^
metal 32023.620: /Users/runner/work/slangpy/slangpy/slangpy/slang/core.slang(410): note :  type 'RWValueRef_default_0' cannot be used in buffer pointee type
    RWValueRef_default_0 _result_0;
    ^
metal 32023.620: 1c7d72fbb205f30c2e60028a63056c511bb2216d(33): note :  type 'CallData_default_0' cannot be used in buffer pointee type
[[kernel]] void compute_main(uint3 sv_groupthreadid_0 [[thread_position_in_threadgroup]], uint3 flat_call_thread_id_0 [[thread_position_in_grid]], uint3 flat_call_group_id_0 [[threadgroup_position_in_grid]], CallData_default_0 constant* call_data_1 [[buffer(0)]])
                                                                                                                                                                                                                ^
metal 32023.620: 1c7d72fbb205f30c2e60028a63056c511bb2216d(33): warning :  unused variable 'sv_groupindex_0' [-Wunused-variable]
    uint sv_groupindex_0 = (sv_groupthreadid_0[int(2)] + sv_groupthreadid_0[int(1)]) * 32U + sv_groupthreadid_0[int(0)];
         ^
metal 32023.620: 1c7d72fbb205f30c2e60028a63056c511bb2216d(33): warning :  unused variable 'flat_call_group_id_1' [-Wunused-variable]
    int3 flat_call_group_id_1 = int3(flat_call_group_id_0);
         ^
2 warnings and 1 error generated.


/Users/runner/work/slangpy/slangpy/src/sgl/device/command.cpp:962 in function finish

Check failure on line 37 in slangpy/tests/slangpy_tests/test_pointers.py

View workflow job for this annotation

GitHub Actions / build (macos, aarch64, clang, Debug, 3.10) 

test_copy_pointer[DeviceType.metal]

RuntimeError: m_rhi_command_encoder->finish(rhi_command_buffer.writeRef()) failed with error: -2147467259 (SLANG_FAIL)

RHI messages:
slang: metal 32023.620: 1c7d72fbb205f30c2e60028a63056c511bb2216d(33): error :  type 'const constant CallData_default_0 *' is not valid for attribute 'buffer'
[[kernel]] void compute_main(uint3 sv_groupthreadid_0 [[thread_position_in_threadgroup]], uint3 flat_call_thread_id_0 [[thread_position_in_grid]], uint3 flat_call_group_id_0 [[threadgroup_position_in_grid]], CallData_default_0 constant* call_data_1 [[buffer(0)]])
                                                                                                                                                                                                                ^                                          ~~~~~~~~~
metal 32023.620: /Users/runner/work/slangpy/slangpy/slangpy/slang/core.slang(408): note :  type 'device int *device *' cannot be used in buffer pointee type
    int device* device* value_0;
    ^
metal 32023.620: /Users/runner/work/slangpy/slangpy/slangpy/slang/core.slang(410): note :  type 'RWValueRef_default_0' cannot be used in buffer pointee type
    RWValueRef_default_0 _result_0;
    ^
metal 32023.620: 1c7d72fbb205f30c2e60028a63056c511bb2216d(33): note :  type 'CallData_default_0' cannot be used in buffer pointee type
[[kernel]] void compute_main(uint3 sv_groupthreadid_0 [[thread_position_in_threadgroup]], uint3 flat_call_thread_id_0 [[thread_position_in_grid]], uint3 flat_call_group_id_0 [[threadgroup_position_in_grid]], CallData_default_0 constant* call_data_1 [[buffer(0)]])
                                                                                                                                                                                                                ^
metal 32023.620: 1c7d72fbb205f30c2e60028a63056c511bb2216d(33): warning :  unused variable 'sv_groupindex_0' [-Wunused-variable]
    uint sv_groupindex_0 = (sv_groupthreadid_0[int(2)] + sv_groupthreadid_0[int(1)]) * 32U + sv_groupthreadid_0[int(0)];
         ^
metal 32023.620: 1c7d72fbb205f30c2e60028a63056c511bb2216d(33): warning :  unused variable 'flat_call_group_id_1' [-Wunused-variable]
    int3 flat_call_group_id_1 = int3(flat_call_group_id_0);
         ^
2 warnings and 1 error generated.


/Users/runner/work/slangpy/slangpy/src/sgl/device/command.cpp:962 in function finish

assert res == 100

Expand Down Expand Up @@ -235,7 +235,7 @@
expected_values = rand_ints[indices]

# Call the function with the grid shape and the pointers buffer0
res = function(grid(shape=(num_ptrs,)), pointers_buffer, _result="numpy")

Check failure on line 238 in slangpy/tests/slangpy_tests/test_pointers.py

View workflow job for this annotation

GitHub Actions / build (macos, aarch64, clang, Debug, 3.10) 

test_raw_buffer_of_pointers[DeviceType.metal]

RuntimeError: m_rhi_command_encoder->finish(rhi_command_buffer.writeRef()) failed with error: -2147467259 (SLANG_FAIL)

RHI messages:
slang: metal 32023.620: ba1235e498d76d663514afeb5be19b8ad6b97360(38): error :  type 'const constant CallData_default_0 *' is not valid for attribute 'buffer'
[[kernel]] void compute_main(uint3 sv_groupthreadid_0 [[thread_position_in_threadgroup]], uint3 flat_call_thread_id_1 [[thread_position_in_grid]], uint3 flat_call_group_id_1 [[threadgroup_position_in_grid]], CallData_default_0 constant* call_data_1 [[buffer(0)]])
                                                                                                                                                                                                                ^                                          ~~~~~~~~~
metal 32023.620: /Users/runner/work/slangpy/slangpy/slangpy/slang/core.slang(200): note :  type 'device int *device *' cannot be used in buffer pointee type
    int device* device* value_0;
    ^
metal 32023.620: /Users/runner/work/slangpy/slangpy/slangpy/slang/tensorcommon.slang(12): note :  type 'StructuredBufferType_default_0' cannot be used in buffer pointee type
    StructuredBufferType_default_0 buffer_0;
    ^
metal 32023.620: ba1235e498d76d663514afeb5be19b8ad6b97360(38): note :  type 'CallData_default_0' cannot be used in buffer pointee type
[[kernel]] void compute_main(uint3 sv_groupthreadid_0 [[thread_position_in_threadgroup]], uint3 flat_call_thread_id_1 [[thread_position_in_grid]], uint3 flat_call_group_id_1 [[threadgroup_position_in_grid]], CallData_default_0 constant* call_data_1 [[buffer(0)]])
                                                                                                                                                                                                                ^
1 error generated.


/Users/runner/work/slangpy/slangpy/src/sgl/device/command.cpp:962 in function finish

assert np.array_equal(res, expected_values), f"Expected {expected_values}, got {res}"

Expand Down Expand Up @@ -592,6 +592,8 @@

@pytest.mark.parametrize("device_type", POINTER_DEVICE_TYPES)
def test_atomic_float_in_struct_access(device_type: DeviceType):
if device_type == DeviceType.metal:
pytest.skip("operator& is not supported on Metal and is being removed from Slang")

device = helpers.get_device(device_type)
function = helpers.create_function_from_module(
Expand Down
Loading