tqchen commented on code in PR #170: URL: https://github.com/apache/tvm-ffi/pull/170#discussion_r2442267825
########## examples/quickstart/src/add_one_cuda.cu: ########## @@ -0,0 +1,48 @@ +/* + * 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.begin] +// File: src/add_one_cuda.cu +#include <tvm/ffi/container/tensor.h> +#include <tvm/ffi/extra/c_env_api.h> +#include <tvm/ffi/function.h> + +namespace tvm_ffi_example_cuda { + +__global__ void AddOneKernel(float* x, float* y, int n) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < n) { + y[idx] = x[idx] + 1; + } +} + +void AddOne(tvm::ffi::TensorView x, tvm::ffi::TensorView y) { + int64_t n = x.shape()[0]; Review Comment: n = x.size(0); ########## examples/quickstart/CMakeLists.txt: ########## @@ -0,0 +1,52 @@ +# 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. +cmake_minimum_required(VERSION 3.20) +project(tvm_ffi_example) + +option(EXAMPLE_NAME "Which example to build (cpp/cuda)" "cpp") +message(STATUS "Building example: ${EXAMPLE_NAME}") + +find_package(tvm_ffi CONFIG) +if (NOT tvm_ffi_FOUND) + message(FATAL_ERROR "Unable to find tvm_ffi. Set `tvm_ffi_DIR` with `tvm-ffi-config --cmakedir`") +endif () + +if (EXAMPLE_NAME STREQUAL "cpp") Review Comment: can we update, detect env and build both? ########## examples/quickstart/README.md: ########## @@ -0,0 +1,57 @@ +<!--- 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. --> + +# Quick Start Code Example + +This directory contains all the source code for [tutorial](https://tvm.apache.org/ffi/get_started/quickstart.html). + +To compile the C++ Example: + +```bash +cmake . -B build -DEXAMPLE_NAME="cpp" -DCMAKE_BUILD_TYPE=RelWithDebInfo Review Comment: DEXAMPLE_NAME="cpu" ########## examples/quick_start/src/add_one_c.c: ########## @@ -1,72 +0,0 @@ -/* - * 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. - */ -#include <tvm/ffi/c_api.h> -#include <tvm/ffi/extra/c_env_api.h> - -// This is a raw C variant of the add_one_cpu function -// it is used to demonstrate how low-level mechanism works -// to construct a tvm ffi compatible function -// -// This function can also serve as a reference for how to implement -// a compiler codegen to target tvm ffi -// -// if you are looking for a more high-level way to construct a tvm ffi compatible function, -// please refer to the add_one_cpu.cc instead -/*! - * \brief Helper code to read DLTensor from TVMFFIAny, can be inlined into generated code - * \param value The TVMFFIAny to read from - * \param out The DLTensor to read into - * \return 0 on success, -1 on error - */ -int ReadDLTensorPtr(const TVMFFIAny* value, DLTensor** out) { - if (value->type_index == kTVMFFIDLTensorPtr) { - *out = (DLTensor*)(value->v_ptr); - return 0; - } - if (value->type_index != kTVMFFITensor) { - // Use TVMFFIErrorSetRaisedFromCStr or TVMFFIErrorSetRaisedFromCStrParts to set an - // error which will be propagated to the caller - TVMFFIErrorSetRaisedFromCStr("ValueError", "Expects a Tensor input"); - return -1; - } - *out = (DLTensor*)((char*)(value->v_obj) + sizeof(TVMFFIObject)); - return 0; -} - -// FFI function implementing add_one operation -int __tvm_ffi_add_one_c( // - void* handle, const TVMFFIAny* args, int32_t num_args, TVMFFIAny* result // -) { - DLTensor *x, *y; - // Extract tensor arguments - // return -1 for error, error is set through TVMFFIErrorSetRaisedFromCStr - if (ReadDLTensorPtr(&args[0], &x) == -1) return -1; Review Comment: We also need run_example.py that comes with it ########## examples/quickstart/load_cpp.cc: ########## @@ -16,38 +16,52 @@ * specific language governing permissions and limitations * under the License. */ +// [example.begin] +// File: load_cpp.cc #include <tvm/ffi/container/tensor.h> #include <tvm/ffi/extra/module.h> -// This file shows how to load the same compiled module and interact with it in C++ +namespace { namespace ffi = tvm::ffi; -struct CPUNDAlloc { - void AllocData(DLTensor* tensor) { tensor->data = malloc(ffi::GetDataSize(*tensor)); } - void FreeData(DLTensor* tensor) { free(tensor->data); } -}; - -inline ffi::Tensor Empty(ffi::Shape shape, DLDataType dtype, DLDevice device) { - return ffi::Tensor::FromNDAlloc(CPUNDAlloc(), shape, dtype, device); -} - -int main() { - // load the module +int Run() { + // Step 1. Load `add_one_cpu.so` module and retrieve `add_one` function ffi::Module mod = ffi::Module::LoadFromFile("build/add_one_cpu.so"); + ffi::Function add_one_cpu = mod->GetFunction("add_one").value(); - // create an Tensor, alternatively, one can directly pass in a DLTensor* - ffi::Tensor x = Empty({5}, DLDataType({kDLFloat, 32, 1}), DLDevice({kDLCPU, 0})); - for (int i = 0; i < 5; ++i) { - reinterpret_cast<float*>(x.data_ptr())[i] = static_cast<float>(i); - } + // Step 2. Create input and output tensors + std::vector<int64_t> shape = {5}; + std::vector<float> x_data = {1, 2, 3, 4, 5}; + std::vector<float> y_data(5, 0); + DLTensor x{ + /*data=*/x_data.data(), + /*device=*/DLDevice{kDLCPU, 0}, + /*ndim=*/1, + /*dtype=*/DLDataType{kDLFloat, 32, 1}, + /*shape=*/shape.data(), + /*strides=*/nullptr, Review Comment: new standard must fill in strides ########## examples/quickstart/src/add_one_cpu.cc: ########## @@ -16,26 +16,24 @@ * specific language governing permissions and limitations * under the License. */ + +// [example.begin] +// File: src/add_one_cpu.cc #include <tvm/ffi/container/tensor.h> -#include <tvm/ffi/dtype.h> -#include <tvm/ffi/error.h> #include <tvm/ffi/function.h> -namespace tvm_ffi_example { +namespace tvm_ffi_example_cpp { +/*! \brief Perform vector add one: y = x + 1 (1-D float32) */ void AddOne(tvm::ffi::TensorView x, tvm::ffi::TensorView y) { - // implementation of a library function - TVM_FFI_ICHECK(x.ndim() == 1) << "x must be a 1D tensor"; - DLDataType f32_dtype{kDLFloat, 32, 1}; - TVM_FFI_ICHECK(x.dtype() == f32_dtype) << "x must be a float tensor"; - TVM_FFI_ICHECK(y.ndim() == 1) << "y must be a 1D tensor"; - TVM_FFI_ICHECK(y.dtype() == f32_dtype) << "y must be a float tensor"; - TVM_FFI_ICHECK(x.size(0) == y.size(0)) << "x and y must have the same shape"; - for (int i = 0; i < x.size(0); ++i) { - static_cast<float*>(y.data_ptr())[i] = static_cast<float*>(x.data_ptr())[i] + 1; + int64_t n = x.shape()[0]; Review Comment: n = x.size(0); ########## examples/quick_start/src/add_one_c.c: ########## @@ -1,72 +0,0 @@ -/* - * 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. - */ -#include <tvm/ffi/c_api.h> -#include <tvm/ffi/extra/c_env_api.h> - -// This is a raw C variant of the add_one_cpu function -// it is used to demonstrate how low-level mechanism works -// to construct a tvm ffi compatible function -// -// This function can also serve as a reference for how to implement -// a compiler codegen to target tvm ffi -// -// if you are looking for a more high-level way to construct a tvm ffi compatible function, -// please refer to the add_one_cpu.cc instead -/*! - * \brief Helper code to read DLTensor from TVMFFIAny, can be inlined into generated code - * \param value The TVMFFIAny to read from - * \param out The DLTensor to read into - * \return 0 on success, -1 on error - */ -int ReadDLTensorPtr(const TVMFFIAny* value, DLTensor** out) { - if (value->type_index == kTVMFFIDLTensorPtr) { - *out = (DLTensor*)(value->v_ptr); - return 0; - } - if (value->type_index != kTVMFFITensor) { - // Use TVMFFIErrorSetRaisedFromCStr or TVMFFIErrorSetRaisedFromCStrParts to set an - // error which will be propagated to the caller - TVMFFIErrorSetRaisedFromCStr("ValueError", "Expects a Tensor input"); - return -1; - } - *out = (DLTensor*)((char*)(value->v_obj) + sizeof(TVMFFIObject)); - return 0; -} - -// FFI function implementing add_one operation -int __tvm_ffi_add_one_c( // - void* handle, const TVMFFIAny* args, int32_t num_args, TVMFFIAny* result // -) { - DLTensor *x, *y; - // Extract tensor arguments - // return -1 for error, error is set through TVMFFIErrorSetRaisedFromCStr - if (ReadDLTensorPtr(&args[0], &x) == -1) return -1; Review Comment: addone c needs to be kept somewhere, perhaps examples/stable_c_bi ########## examples/quickstart/CMakeLists.txt: ########## @@ -0,0 +1,52 @@ +# 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. +cmake_minimum_required(VERSION 3.20) +project(tvm_ffi_example) + +option(EXAMPLE_NAME "Which example to build (cpp/cuda)" "cpp") +message(STATUS "Building example: ${EXAMPLE_NAME}") + +find_package(tvm_ffi CONFIG) +if (NOT tvm_ffi_FOUND) Review Comment: it is better for CMake to work out ot box, put in the standard line before ```cmake # Run `tvm_ffi.config --cmakedir` to find tvm-ffi targets execute_process( COMMAND "${Python_EXECUTABLE}" -m tvm_ffi.config --cmakedir OUTPUT_STRIP_TRAILING_WHITESPACE OUTPUT_VARIABLE tvm_ffi_ROOT ) ``` -- 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]
