gemini-code-assist[bot] commented on code in PR #283: URL: https://github.com/apache/tvm-ffi/pull/283#discussion_r2552365260
########## examples/cubin_launcher/src/lib_embedded.cc: ########## @@ -0,0 +1,129 @@ +/* + * Licensed to the Apache Software Foundation (ASF) under one + * or more contributor license agreements. See the NOTICE file + * distributed with this work for additional information + * regarding copyright ownership. The ASF licenses this file + * to you under the Apache License, Version 2.0 (the + * "License"); you may not use this file except in compliance + * with the License. You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, + * software distributed under the License is distributed on an + * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY + * KIND, either express or implied. See the License for the + * specific language governing permissions and limitations + * under the License. + */ +/*! + * \file examples/cubin_launcher/src/lib_embedded.cc + * \brief TVM-FFI library with embedded CUBIN kernels. + * + * This library exports TVM-FFI functions to launch CUDA kernels from + * embedded CUBIN data. + */ + +#include <tvm/ffi/container/tensor.h> +#include <tvm/ffi/error.h> +#include <tvm/ffi/extra/c_env_api.h> +#include <tvm/ffi/extra/cubin_launcher.h> +#include <tvm/ffi/function.h> + +#include <cstdint> + +// External symbols for embedded CUBIN data (linked via objcopy) +extern "C" const char __cubin_data[]; +extern "C" const char __cubin_data_end[]; + +// Calculate size from the symbols +static const uint64_t cubin_data_size = + reinterpret_cast<const char*>(&__cubin_data_end) - reinterpret_cast<const char*>(&__cubin_data); + +// Global CUBIN module and kernels (initialized on first use) +static tvm::ffi::CubinModule* g_cubin_module = nullptr; +static tvm::ffi::CubinKernel* g_add_one_kernel = nullptr; +static tvm::ffi::CubinKernel* g_mul_two_kernel = nullptr; + +// Initialize the CUBIN module and kernels +void InitializeCubinModule() { + if (g_cubin_module == nullptr) { + g_cubin_module = new tvm::ffi::CubinModule(__cubin_data, cubin_data_size); + g_add_one_kernel = new tvm::ffi::CubinKernel((*g_cubin_module)["add_one_cuda"]); + g_mul_two_kernel = new tvm::ffi::CubinKernel((*g_cubin_module)["mul_two_cuda"]); + } +} Review Comment:  The current implementation of global module and kernel management has two issues: 1. **Resource Leak**: The global `CubinModule` and `CubinKernel` objects are allocated with `new` but are never deallocated with `delete`. This leads to a memory leak and, more importantly, a leak of CUDA resources (`CUlibrary`) as their destructors are never called. 2. **Race Condition**: The `InitializeCubinModule` function is not thread-safe. If multiple threads call it for the first time concurrently, it can lead to race conditions, potentially causing multiple initializations or crashes. To address these issues, you should use `std::unique_ptr` for automatic resource management and `std::call_once` for thread-safe initialization. Please also add `#include <memory>` and `#include <mutex>` at the top of the file. ```suggestion static std::unique_ptr<tvm::ffi::CubinModule> g_cubin_module; static std::unique_ptr<tvm::ffi::CubinKernel> g_add_one_kernel; static std::unique_ptr<tvm::ffi::CubinKernel> g_mul_two_kernel; static std::once_flag g_init_flag; // Initialize the CUBIN module and kernels void InitializeCubinModule() { std::call_once(g_init_flag, []() { g_cubin_module = std::make_unique<tvm::ffi::CubinModule>(__cubin_data, cubin_data_size); g_add_one_kernel = std::make_unique<tvm::ffi::CubinKernel>((*g_cubin_module)["add_one_cuda"]); g_mul_two_kernel = std::make_unique<tvm::ffi::CubinKernel>((*g_cubin_module)["mul_two_cuda"]); }); } ``` ########## examples/cubin_launcher/example_embeded_cubin.py: ########## @@ -0,0 +1,112 @@ +#!/usr/bin/env python3 +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. +"""Example script for embedded CUBIN library. Review Comment:  The filename `example_embeded_cubin.py` contains a typo ("embeded"). It should be "embedded". Please consider renaming the file to `example_embedded_cubin.py` and updating all references to it for consistency. ########## examples/cubin_launcher/README.md: ########## @@ -0,0 +1,79 @@ +<!--- Licensed to the Apache Software Foundation (ASF) under one --> +<!--- or more contributor license agreements. See the NOTICE file --> +<!--- distributed with this work for additional information --> +<!--- regarding copyright ownership. The ASF licenses this file --> +<!--- to you under the Apache License, Version 2.0 (the --> +<!--- "License"); you may not use this file except in compliance --> +<!--- with the License. You may obtain a copy of the License at --> + +<!--- http://www.apache.org/licenses/LICENSE-2.0 --> + +<!--- Unless required by applicable law or agreed to in writing, --> +<!--- software distributed under the License is distributed on an --> +<!--- "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY --> +<!--- KIND, either express or implied. See the License for the --> +<!--- specific language governing permissions and limitations --> +<!--- under the License. --> + +# CUBIN Launcher + +## Overview + +Demonstrates loading and executing CUDA kernels from CUBIN files using TVM-FFI. The `cubin_launcher.h` header wraps CUDA Driver API to provide lightweight CUBIN module and kernel management. + +## Techniques + +The implementation uses CUDA Driver API Library Management: + +- **`cuLibraryLoadData()`** - Load CUBIN from memory buffer +- **`cuLibraryGetKernel()`** - Get kernel handle by name +- **`cuKernelGetFunction()`** - Get function handle for current CUDA context +- **`cuLaunchKernel()`** - Launch kernel with grid/block dimensions + +Key features: + +- Multi-GPU support via CUDA primary contexts +- RAII-based resource management (CubinModule, CubinKernel) +- CUBIN embedding at compile time (via `ld` + `objcopy`) +- TVM-FFI integration for tensor argument passing + +## Examples + +### 1. Embedded CUBIN (TVM-FFI Library) + +`example_embeded_cubin.py` - CUBIN linked into shared library at build time. Review Comment:  There's a typo in the filename. It should be `embedded` instead of `embeded`. This typo appears in multiple places, including the filename itself. It would be best to correct it everywhere for consistency and correctness. ```suggestion `example_embedded_cubin.py` - CUBIN linked into shared library at build time. ``` ########## examples/cubin_launcher/README.md: ########## @@ -0,0 +1,79 @@ +<!--- Licensed to the Apache Software Foundation (ASF) under one --> +<!--- or more contributor license agreements. See the NOTICE file --> +<!--- distributed with this work for additional information --> +<!--- regarding copyright ownership. The ASF licenses this file --> +<!--- to you under the Apache License, Version 2.0 (the --> +<!--- "License"); you may not use this file except in compliance --> +<!--- with the License. You may obtain a copy of the License at --> + +<!--- http://www.apache.org/licenses/LICENSE-2.0 --> + +<!--- Unless required by applicable law or agreed to in writing, --> +<!--- software distributed under the License is distributed on an --> +<!--- "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY --> +<!--- KIND, either express or implied. See the License for the --> +<!--- specific language governing permissions and limitations --> +<!--- under the License. --> + +# CUBIN Launcher + +## Overview + +Demonstrates loading and executing CUDA kernels from CUBIN files using TVM-FFI. The `cubin_launcher.h` header wraps CUDA Driver API to provide lightweight CUBIN module and kernel management. + +## Techniques + +The implementation uses CUDA Driver API Library Management: + +- **`cuLibraryLoadData()`** - Load CUBIN from memory buffer +- **`cuLibraryGetKernel()`** - Get kernel handle by name +- **`cuKernelGetFunction()`** - Get function handle for current CUDA context +- **`cuLaunchKernel()`** - Launch kernel with grid/block dimensions + +Key features: + +- Multi-GPU support via CUDA primary contexts +- RAII-based resource management (CubinModule, CubinKernel) +- CUBIN embedding at compile time (via `ld` + `objcopy`) +- TVM-FFI integration for tensor argument passing + +## Examples + +### 1. Embedded CUBIN (TVM-FFI Library) + +`example_embeded_cubin.py` - CUBIN linked into shared library at build time. + +```bash +cd build +cmake .. +make +cd .. +python examples/cubin_launcher/example_embeded_cubin.py Review Comment:  Correcting the typo in the filename here as well. ```suggestion python examples/cubin_launcher/example_embedded_cubin.py ``` -- This is an automated message from the Apache Git Service. To respond to the message, please log on to GitHub and use the URL above to go to the specific comment. To unsubscribe, e-mail: [email protected] For queries about this service, please contact Infrastructure at: [email protected] --------------------------------------------------------------------- To unsubscribe, e-mail: [email protected] For additional commands, e-mail: [email protected]
