提交 666c6003 编写于 作者: Z zhiru

update


Former-commit-id: 1a3242b1ac87cad3274cc28d2e73697e55936e8b
......@@ -19,6 +19,7 @@ Please mark all change in change log and use the ticket from JIRA.
## Task
- MS-554 - Change license to Apache 2.0
- MS-561 - Add contributing guidelines, code of conduct and README docs
# Milvus 0.4.0 (2019-09-12)
......
# Milvus Code of Conduct
In the interest of fostering an open and welcoming environment, we as contributors and maintainers pledge to making participation in our project and our community a harassment-free experience for everyone, regardless of age, body size, disability, ethnicity, gender identity and expression, level of experience, nationality, personal appearance, race, religion, or sexual identity and orientation.
## Our Standards
Examples of behavior that contributes to creating a positive environment include:
- Using welcoming and inclusive language.
- Being respectful of differing viewpoints and experiences.
- Gracefully accepting constructive criticism.
- Focusing on what is best for the community.
- Showing empathy towards other community members.
Examples of unacceptable behavior by participants include:
- The use of sexualized language or imagery and unwelcome sexual attention or advances.
- Trolling, insulting/derogatory comments, and personal or political attacks.
- Public or private harassment.
- Publishing others' private information, such as a physical or electronic address, without explicit permission.
- Conduct which could reasonably be considered inappropriate for the forum in which it occurs.
All Milvus forums and spaces are meant for professional interactions, and any behavior which could reasonably be considered inappropriate in a professional setting is unacceptable.
## Our Responsibilities
Project maintainers are responsible for clarifying the standards of acceptable behavior and are expected to take appropriate and fair corrective action in response to any instances of unacceptable behavior.
Project maintainers have the right and responsibility to remove, edit, or reject comments, commits, code, wiki edits, issues, and other contributions that are not aligned to this Code of Conduct, or to ban temporarily or permanently any contributor for other behaviors that they deem inappropriate, threatening, offensive, or harmful.
## Scope
This Code of Conduct applies to all content on milvus.io, Milvus’s GitHub organization, or any other official Milvus web presence allowing for community interactions, as well as at all official Milvus events, whether offline or online.
The Code of Conduct also applies within all project spaces and in public spaces whenever an individual is representing Milvus or its community. Examples of representing a project or community include using an official project e-mail address, posting via an official social media account, or acting as an appointed or de facto representative at an online or offline event. Representation of Milvus may be further defined and clarified by project maintainers.
## Enforcement
Instances of abusive, harassing, or otherwise unacceptable behavior may be reported by contacting the project team at support@zilliz.com. All complaints will be reviewed and investigated and will result in a response that is deemed necessary and appropriate to the circumstances. The project team is obligated to maintain confidentiality with regard to the reporter of an incident. Further details of specific enforcement policies may be posted separately.
Project maintainers who do not follow or enforce the Code of Conduct in good faith may face temporary or permanent repercussions as determined by other members of the project's leadership.
## Attribution
This Code of Conduct is adapted from the [Contributor Covenant](https://www.contributor-covenant.org/), version 1.4, available at https://www.contributor-covenant.org/version/1/4/code-of-conduct.html
For answers to common questions about this code of conduct, see https://www.contributor-covenant.org/faq
# Contributing to Milvus
First of all, thanks for taking the time to contribute to Milvus! It's people like you that help Milvus come to fruition.
The following are a set of guidelines for contributing to Milvus. Following these guidelines helps contributing to this project easy and transparent. These are mostly guideline, not rules. Use your best judgment, and feel free to propose changes to this document in a pull request.
As for everything else in the project, the contributions to Milvus are governed by our [Code of Conduct](http://hood.ie/code-of-conduct/).
TOC
## Contribution Checklist
Before you make any contributions, make sure you follow this list.
- Read [Contributing to Milvus](CONTRIBUTING.md).
- Check if the changes are consistent with the [coding style](CONTRIBUTING.md#coding-style).
- Run [unit tests](CONTRIBUTING.md#run-unit-test).
## What contributions can I make?
Contributions to Milvus fall into the following categories.
1. To report a bug or a problem with documentation, please file an [issue](https://github.com/milvus-io/milvus/issues/new) providing the details of the problem. If you believe the issue needs priority attention, please comment on the issue to notify the team.
2. To propose a new feature, please file a new feature request [issue](https://github.com/milvus-io/milvus/issues/new). Describe the intended feature and discuss the design and implementation with the team and community. Once the team agrees that the plan looks good, go ahead and implement it, following the [Contributing code].
3. To implement a feature or bug-fix for an existing outstanding issue, follow the [Contributing code]. If you need more context on a particular issue, comment on the issue to let people know.
## How can I contribute?
### Contributing code
If you have improvements to Milvus, send us your pull requests! For those just getting started, GitHub has a [how-to](https://help.github.com/en/articles/about-pull-requests).
The Milvus team members will review your pull requests, and once it is accepted, it will be given a `ready to merge` label. This means we are working on submitting your pull request to the internal repository. After the change has been submitted internally, your pull request will be merged automatically on GitHub.
### General guidelines
Before sending your pull requests for review, make sure your changes are consistent with the guidelines and follow the Milvus coding style.
- Include unit tests when you contribute new features, as they help to a) prove that your code works correctly, and b) guard against future breaking changes to lower the maintenance cost.
- Bug fixes also generally require unit tests, because the presence of bugs usually indicates insufficient test coverage.
- Keep API compatibility in mind when you change code in Milvus. Reviewers of your pull request will comment on any API compatibility issues.
- When you contribute a new feature to Milvus, the maintenance burden is (by default) transferred to the Milvus team. This means that the benefit of the contribution must be compared against the cost of maintaining the feature.
## Coding Style
## Run unit test
```shell
$ ./build.sh -u
or
$ ./build.sh --unittest
```
### Compilation
#### Step 1: install necessery tools
# Welcome to Milvus
centos7 :
yum install gfortran qt4 flex bison mysql-devel mysql
ubuntu16.04 :
sudo apt-get install gfortran qt4-qmake flex bison libmysqlclient-dev mysql-client
cd scripts && sudo ./requirements.sh
Firstly, welcome, and thanks for your interest in [Milvus](https://milvus.io)! No matter who you are, what you do, we greatly appreciate your contribution to help us reinvent data science with Milvus.
If `libmysqlclient_r.so` does not exist after installing MySQL Development Files, you need to create a symbolic link:
## What is Milvus
Milvus is an open source vector search engine that makes incredibly fast querying speed enhancement over current solutions of massive vector processing. Built on optimized indexing algorithm, it is compatible with major AI/ML models.
Milvus was developed by researchers and engineers in ZILLIZ, a tech startup that intends to reinvent data science, with the purpose of providing enterprises with efficient and scalable similarity search and analysis of feature vectors and unstructured data.
Milvus provides stable Python and C++ APIs, as well as RESTful API.
Keep up-to-date with newest releases and latest updates by reading [Releases](https://www.milvus-io/docs/master/releases).
- GPU-accelerated search engine
Milvus is designed for the largest scale of vector index. CPU/GPU heterogeneous computing architecture allows you to process data at a speed 1000 times faster.
- Intelligent index
With a “Decide Your Own Algorithm” approach, you can embed machine learning and advanced algorithms into Milvus without the headache of complex data engineering or migrating data between disparate systems. Milvus is built on optimized indexing algorithm based on quantization indexing, tree-based and graph indexing methods.
- Strong scalability
The data is stored and computed on a distributed architecture. This lets you scale data sizes up and down without redesigning the system.
## Architecture
![Milvus_arch](https://www.milvus-io/docs/master/assets/milvus_arch.png)
## Get started
### Install and start Milvus server
#### Use Docker
Use Docker to install Milvus is a breeze. See the [Milvus install guide](https://www.milvus-io/docs/master/userguide/install_milvus.md) for details.
#### Use source code
##### Compilation
###### Step 1 Install necessary tools
```shell
# Install tools
Centos7 :
$ yum install gfortran qt4 flex bison mysql-devel mysql
Ubuntu16.04 :
$ sudo apt-get install gfortran qt4-qmake flex bison libmysqlclient-dev mysql-client
```
Verify the existence of `libmysqlclient_r.so`:
```shell
# Verify existence
$ locate libmysqlclient_r.so
```
sudo ln -s /path/to/libmysqlclient.so /path/to/libmysqlclient_r.so
If not, you need to create a symbolic link:
```shell
# Create symbolic link
$ sudo ln -s /path/to/libmysqlclient.so /path/to/libmysqlclient_r.so
```
#### Step 2: build(output to cmake_build folder)
###### Step 2 Build
cmake_build/src/milvus_server is the server
```shell
TBD
cd [Milvus sourcecode path]/cpp/thirdparty
git clone git@192.168.1.105:megasearch/knowhere.git
cd knowhere
./build.sh -t Debug
or ./build.sh -t Release
cd [sourcecode path]/cpp/thirdparty
git clone git@192.168.1.105:megasearch/knowhere.git
cd knowhere
./build.sh -t Debug
or ./build.sh -t Release
cd [sourcecode path]/cpp
./build.sh -t Debug
or ./build.sh -t Release
```
cd [sourcecode path]/cpp
./build.sh -t Debug
or ./build.sh -t Release
When the build is completed, all the stuff that you need in order to run Milvus will be installed under `[Milvus root path]/cpp/milvus`.
If you encounter the following error when building:
If you encounter the following error message,
`protocol https not supported or disabled in libcurl`
1. Install libcurl4-openssl-dev
2. Install cmake 3.14:
please reinstall CMake with curl:
1. Install curl development files:
```shell
CentOS 7:
$ yum install curl-devel
Ubuntu 16.04:
$ sudo apt-get install libcurl4-openssl-dev
```
./bootstrap --system-curl
make
sudo make install
2. Install [CMake 3.14](https://github.com/Kitware/CMake/releases/download/v3.14.6/cmake-3.14.6.tar.gz):
```shell
$ ./bootstrap --system-curl
$ make
$ sudo make install
```
#### To build unittest:
##### Run unit test
```shell
$ ./build.sh -u
or
$ ./build.sh --unittest
```
./build.sh -u
or
./build.sh --unittest
##### Run code coverage
#### To run code coverage
```shell
CentOS 7:
$ yum install lcov
Ubuntu 16.04:
$ sudo apt-get install lcov
$ ./build.sh -u -c
```
apt-get install lcov
./build.sh -u -c
##### Launch Milvus server
### Launch server
Set config in cpp/conf/server_config.yaml
```shell
$ cd [Milvus root path]/cpp/milvus
```
Add milvus/lib to LD_LIBRARY_PATH
Add `lib/` directory to `LD_LIBRARY_PATH`
```
export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/path/to/milvus/lib
$ export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/path/to/milvus/lib
```
Then start Milvus server:
```
$ cd scripts
$ ./start_server.sh
```
To stop Milvus server, run:
```shell
$ ./stop_server.sh
```
To edit Milvus settings in `conf/server_config.yaml` and `conf/log_config.conf`, please read [Milvus Configuration](https://www.milvus-io/docs/master/reference/milvus_config.md).
### Try your first Milvus program
#### Run Python example code
Make sure [Python 3.4](https://www.python.org/downloads/) or higher is already installed and in use.
Install Milvus Python SDK.
```shell
# Install Milvus Python SDK
$ pip install pymilvus==0.2.0
```
Create a new file `example.py`, and add [Python example code](https://github.com/milvus-io/pymilvus/blob/branch-0.3.1/examples/AdvancedExample.py) to it.
Run the example code.
```python
# Run Milvus Python example
$ python3 example.py
```
#### Run C++ example code
```shell
# Run Milvus C++ example
$ cd [Milvus root path]/cpp/milvus/bin
$ ./sdk_simple
```
## Contribution guidelines
Contributions are welcomed and greatly appreciated. If you want to contribute to Milvus, please read the [contribution guidelines](CONTRIBUTING.md). This project adheres to the [code of conduct](CODE OF CONDUCT.md) of Milvus. By participating, you are expected to uphold this code.
We use [GitHub issues] to track issues and bugs. For general questions and discussions, please go to [Milvus Forum].
## Join the Milvus community
For public discussion of Milvus, please join our [discussion group].
## Milvus Roadmap
Please read our [roadmap](milvus-io/milvus/docs/master/roadmap.md) to learn about upcoming features.
## Resources
[Milvus official website](https://milvus.io)
[Milvus docs](https://www.milvus.io/docs/en/QuickStart/)
[Milvus blog](https://www.milvus.io/blog/)
[Milvus CSDN](https://mp.csdn.net/mdeditor/100041006#)
[Milvus roadmap](https://www.milvus-io/docs/master/roadmap.md)
[Milvus white paper]
Then launch server with config:
cd [build output path]
start_server.sh
stop_server.sh
## License
### Launch test_client(only for debug)
If you want to test remote api, you can run sdk example.
[build output path]/sdk/examples/grpcsimple/sdk_simple
[Apache 2.0 license](milvus-io/milvus/LICENSE.md)
\ No newline at end of file
......@@ -73,7 +73,7 @@ aux_source_directory(${MILVUS_ENGINE_SRC}/server/grpc_impl grpc_server_files)
aux_source_directory(${MILVUS_ENGINE_SRC}/utils utils_files)
aux_source_directory(${MILVUS_ENGINE_SRC}/wrapper/knowhere wrapper_knowhere_files)
aux_source_directory(${MILVUS_ENGINE_SRC}/wrapper wrapper_files)
set(engine_files
${CMAKE_CURRENT_SOURCE_DIR}/main.cpp
......@@ -85,7 +85,7 @@ set(engine_files
${db_scheduler_files}
${metrics_files}
${utils_files}
${wrapper_knowhere_files}
${wrapper_files}
)
include_directories("${CUDA_TOOLKIT_ROOT_DIR}/include")
......
......@@ -18,7 +18,7 @@
#pragma once
#include "wrapper/knowhere/vec_index.h"
#include "src/wrapper/vec_index.h"
#include <memory>
......
......@@ -84,7 +84,7 @@ include(DefineOptionsCore)
include(BuildUtilsCore)
include(ThirdPartyPackagesCore)
add_subdirectory(src)
add_subdirectory(knowhere)
if (BUILD_COVERAGE STREQUAL "ON")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fprofile-arcs -ftest-coverage")
......
#pragma once
#include <memory>
#include <knowhere/common/array.h>
namespace zilliz {
namespace knowhere {
ArrayPtr
CopyArray(const ArrayPtr &origin);
SchemaPtr
CopySchema(const SchemaPtr &origin);
} // namespace knowhere
} // namespace zilliz
#pragma once
#include <memory>
#include <knowhere/common/dataset.h>
#include <SPTAG/AnnService/inc/Core/VectorIndex.h>
namespace zilliz {
namespace knowhere {
std::shared_ptr<SPTAG::VectorSet>
ConvertToVectorSet(const DatasetPtr &dataset);
std::shared_ptr<SPTAG::MetadataSet>
ConvertToMetadataSet(const DatasetPtr &dataset);
std::vector<SPTAG::QueryResult>
ConvertToQueryResult(const DatasetPtr &dataset, const Config &config);
DatasetPtr
ConvertToDataset(std::vector<SPTAG::QueryResult> query_results);
} // namespace knowhere
} // namespace zilliz
#pragma once
#include <jsoncons/json.hpp>
namespace zilliz {
namespace knowhere {
using Config = jsoncons::json;
} // namespace knowhere
} // namespace zilliz
#pragma once
namespace zilliz {
namespace sched {
namespace master {
} // namespace master
} // namespace sched
} // namespace zilliz
#pragma once
#include <cstdint>
#include "zlibrary/error/error.h"
namespace zilliz {
namespace knowhere {
using Error = zilliz::lib::ErrorCode;
constexpr Error STORE_SUCCESS = zilliz::lib::SUCCESS_CODE;
constexpr Error ERROR_CODE_BASE = 0x36000;
constexpr Error ERROR_CODE_END = 0x37000;
constexpr Error
ToGlobalErrorCode(const Error error_code) {
return zilliz::lib::ToGlobalErrorCode(error_code, ERROR_CODE_BASE);
}
class Exception : public zilliz::lib::Exception {
public:
Exception(const Error error_code,
const std::string &message = nullptr)
: zilliz::lib::Exception(error_code, "KNOWHERE", message) {}
};
constexpr Error UNEXPECTED = ToGlobalErrorCode(0x001);
constexpr Error UNSUPPORTED = ToGlobalErrorCode(0x002);
constexpr Error NULL_POINTER = ToGlobalErrorCode(0x003);
constexpr Error OVERFLOW = ToGlobalErrorCode(0x004);
constexpr Error INVALID_ARGUMENT = ToGlobalErrorCode(0x005);
constexpr Error UNSUPPORTED_TYPE = ToGlobalErrorCode(0x006);
} // namespace store
} // namespace zilliz
using Error = zilliz::store::Error;
#pragma once
//#include "zcommon/id/id.h"
//using ID = zilliz::common::ID;
#include <stdint.h>
#include <string>
namespace zilliz {
namespace knowhere {
class ID {
public:
constexpr static int64_t kIDSize = 20;
public:
const int32_t *
data() const { return content_; }
int32_t *
mutable_data() { return content_; }
bool
IsValid() const;
std::string
ToString() const;
bool
operator==(const ID &that) const;
bool
operator<(const ID &that) const;
protected:
int32_t content_[5] = {};
};
} // namespace knowhere
} // namespace zilliz
#pragma once
#include <memory>
#include "arrow/type.h"
namespace zilliz {
namespace knowhere {
using DataType = arrow::DataType;
using Field = arrow::Field;
using FieldPtr = std::shared_ptr<arrow::Field>;
using Schema = arrow::Schema;
using SchemaPtr = std::shared_ptr<Schema>;
using SchemaConstPtr = std::shared_ptr<const Schema>;
} // namespace knowhere
} // namespace zilliz
#pragma once
#include <memory>
#include "arrow/tensor.h"
namespace zilliz {
namespace knowhere {
using Tensor = arrow::Tensor;
using TensorPtr = std::shared_ptr<Tensor>;
} // namespace knowhere
} // namespace zilliz
#pragma once
#include <memory>
#include "knowhere/common/binary_set.h"
#include "knowhere/common/dataset.h"
#include "knowhere/index/index_type.h"
#include "knowhere/index/index_model.h"
#include "knowhere/index/preprocessor/preprocessor.h"
namespace zilliz {
namespace knowhere {
class Index {
public:
virtual BinarySet
Serialize() = 0;
virtual void
Load(const BinarySet &index_binary) = 0;
// @throw
virtual DatasetPtr
Search(const DatasetPtr &dataset, const Config &config) = 0;
public:
IndexType
idx_type() const { return idx_type_; }
void
set_idx_type(IndexType idx_type) { idx_type_ = idx_type; }
virtual void
set_preprocessor(PreprocessorPtr preprocessor) {}
virtual void
set_index_model(IndexModelPtr model) {}
private:
IndexType idx_type_;
};
using IndexPtr = std::shared_ptr<Index>;
} // namespace knowhere
} // namespace zilliz
#pragma once
#include <memory>
#include "knowhere/common/binary_set.h"
namespace zilliz {
namespace knowhere {
class IndexModel {
public:
virtual BinarySet
Serialize() = 0;
virtual void
Load(const BinarySet &binary) = 0;
};
using IndexModelPtr = std::shared_ptr<IndexModel>;
} // namespace knowhere
} // namespace zilliz
#pragma once
namespace zilliz {
namespace knowhere {
enum class IndexType {
kUnknown = 0,
kVecIdxBegin = 100,
kVecIVFFlat = kVecIdxBegin,
kVecIdxEnd,
};
} // namespace knowhere
} // namespace zilliz
#pragma once
#include <memory>
#include "knowhere/common/dataset.h"
namespace zilliz {
namespace knowhere {
class Preprocessor {
public:
virtual DatasetPtr
Preprocess(const DatasetPtr &input) = 0;
};
using PreprocessorPtr = std::shared_ptr<Preprocessor>;
} // namespace knowhere
} // namespace zilliz
#pragma once
#include <string>
namespace zilliz {
namespace knowhere {
#define META_ROWS ("rows")
#define META_DIM ("dimension")
#define META_K ("k")
} // namespace knowhere
} // namespace zilliz
#pragma once
#include <string>
#include <vector>
namespace zilliz {
namespace knowhere {
using KDTParameter = std::pair<std::string, std::string>;
class KDTParameterManagement {
public:
const std::vector<KDTParameter> &
GetKDTParameters();
public:
static KDTParameterManagement &
GetInstance() {
static KDTParameterManagement instance;
return instance;
}
KDTParameterManagement(const KDTParameterManagement &) = delete;
KDTParameterManagement &operator=(const KDTParameterManagement &) = delete;
private:
KDTParameterManagement();
private:
std::vector<KDTParameter> kdt_parameters_;
};
} // namespace knowhere
} // namespace zilliz
%module nsg
%{
#define SWIG_FILE_WITH_INIT
#include <numpy/arrayobject.h>
/* Include the header in the wrapper code */
#include "nsg.h"
%}
/* Parse the header file */
%include "index.h"
#pragma once
#include <memory>
#include "knowhere/common/config.h"
#include "knowhere/common/dataset.h"
#include "knowhere/index/index.h"
#include "knowhere/index/preprocessor/preprocessor.h"
namespace zilliz {
namespace knowhere {
class VectorIndex;
using VectorIndexPtr = std::shared_ptr<VectorIndex>;
class VectorIndex : public Index {
public:
virtual PreprocessorPtr
BuildPreprocessor(const DatasetPtr &dataset, const Config &config) { return nullptr; }
virtual IndexModelPtr
Train(const DatasetPtr &dataset, const Config &config) { return nullptr; }
virtual void
Add(const DatasetPtr &dataset, const Config &config) = 0;
virtual void
Seal() = 0;
virtual VectorIndexPtr
Clone() = 0;
virtual int64_t
Count() = 0;
virtual int64_t
Dimension() = 0;
};
} // namespace knowhere
} // namespace zilliz
......@@ -5,7 +5,7 @@ include_directories(${TBB_DIR}/include)
include_directories(${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES})
link_directories(${CUDA_TOOLKIT_ROOT_DIR}/lib64)
include_directories(${CORE_SOURCE_DIR}/include)
include_directories(${CORE_SOURCE_DIR}/knowhere)
include_directories(${CORE_SOURCE_DIR}/thirdparty)
include_directories(${CORE_SOURCE_DIR}/thirdparty/SPTAG/AnnService)
include_directories(${CORE_SOURCE_DIR}/thirdparty/jsoncons-0.126.0/include)
......@@ -107,7 +107,7 @@ INSTALL(FILES ${CORE_SOURCE_DIR}/thirdparty/tbb/libtbb.so
)
set(CORE_INCLUDE_DIRS
${CORE_SOURCE_DIR}/include
${CORE_SOURCE_DIR}/knowhere
${CORE_SOURCE_DIR}/thirdparty
${CORE_SOURCE_DIR}/thirdparty/SPTAG/AnnService
${CORE_SOURCE_DIR}/thirdparty/jsoncons-0.126.0/include
......@@ -120,18 +120,18 @@ set(CORE_INCLUDE_DIRS
set(CORE_INCLUDE_DIRS ${CORE_INCLUDE_DIRS} PARENT_SCOPE)
INSTALL(DIRECTORY
${CORE_SOURCE_DIR}/include/knowhere
${CORE_SOURCE_DIR}/thirdparty/jsoncons-0.126.0/include/jsoncons
${CORE_SOURCE_DIR}/thirdparty/jsoncons-0.126.0/include/jsoncons_ext
${ARROW_INCLUDE_DIR}/arrow
${FAISS_PREFIX}/include/faiss
${OPENBLAS_INCLUDE_DIR}/
${CORE_SOURCE_DIR}/thirdparty/tbb/include/tbb
DESTINATION
include)
INSTALL(DIRECTORY
${SPTAG_SOURCE_DIR}/AnnService/inc/
DESTINATION
include/SPTAG/AnnService/inc)
#INSTALL(DIRECTORY
# ${CORE_SOURCE_DIR}/include/knowhere
# ${CORE_SOURCE_DIR}/thirdparty/jsoncons-0.126.0/include/jsoncons
# ${CORE_SOURCE_DIR}/thirdparty/jsoncons-0.126.0/include/jsoncons_ext
# ${ARROW_INCLUDE_DIR}/arrow
# ${FAISS_PREFIX}/include/faiss
# ${OPENBLAS_INCLUDE_DIR}/
# ${CORE_SOURCE_DIR}/thirdparty/tbb/include/tbb
# DESTINATION
# include)
#
#INSTALL(DIRECTORY
# ${SPTAG_SOURCE_DIR}/AnnService/inc/
# DESTINATION
# include/SPTAG/AnnService/inc)
// 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 "knowhere/adapter/arrow.h"
#include "arrow.h"
namespace zilliz {
namespace knowhere {
......
// 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.
#pragma once
#include <memory>
#include "knowhere/common/array.h"
namespace zilliz {
namespace knowhere {
ArrayPtr
CopyArray(const ArrayPtr &origin);
SchemaPtr
CopySchema(const SchemaPtr &origin);
} // namespace knowhere
} // namespace zilliz
// 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 "knowhere/index/vector_index/definitions.h"
#include "knowhere/adapter/sptag.h"
#include "knowhere/adapter/structure.h"
#include "sptag.h"
#include "structure.h"
namespace zilliz {
......
// 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.
#pragma once
#include <memory>
#include <SPTAG/AnnService/inc/Core/VectorIndex.h>
#include "knowhere/common/dataset.h"
namespace zilliz {
namespace knowhere {
std::shared_ptr<SPTAG::VectorSet>
ConvertToVectorSet(const DatasetPtr &dataset);
std::shared_ptr<SPTAG::MetadataSet>
ConvertToMetadataSet(const DatasetPtr &dataset);
std::vector<SPTAG::QueryResult>
ConvertToQueryResult(const DatasetPtr &dataset, const Config &config);
DatasetPtr
ConvertToDataset(std::vector<SPTAG::QueryResult> query_results);
} // namespace knowhere
} // namespace zilliz
......@@ -16,7 +16,7 @@
// under the License.
#include "knowhere/adapter/structure.h"
#include "structure.h"
namespace zilliz {
......
......@@ -19,7 +19,7 @@
#pragma once
#include <memory>
#include <knowhere/common/dataset.h>
#include "knowhere/common/dataset.h"
namespace zilliz {
......
// 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.
#pragma once
#include "arrow/array.h"
#include "knowhere/common/schema.h"
#include "schema.h"
namespace zilliz {
......
// 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.
#pragma once
#include <map>
......@@ -5,7 +23,7 @@
#include <vector>
#include <memory>
#include "knowhere/common/id.h"
#include "id.h"
namespace zilliz {
......
// 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.
#pragma once
#include <memory>
#include "arrow/buffer.h"
......
// 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.
#pragma once
#include <jsoncons/json.hpp>
namespace zilliz {
namespace knowhere {
using Config = jsoncons::json;
} // namespace knowhere
} // namespace zilliz
// 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.
#pragma once
#include <vector>
#include <memory>
#include "knowhere/common/array.h"
#include "knowhere/common/buffer.h"
#include "knowhere/common/tensor.h"
#include "knowhere/common/schema.h"
#include "knowhere/common/config.h"
#include "array.h"
#include "buffer.h"
#include "tensor.h"
#include "schema.h"
#include "config.h"
#include "knowhere/adapter/arrow.h"
......
......@@ -16,9 +16,10 @@
// under the License.
#include "knowhere/common/exception.h"
#include <cstdio>
#include "exception.h"
namespace zilliz {
namespace knowhere {
......
// 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.
#pragma once
//#include "zcommon/id/id.h"
//using ID = zilliz::common::ID;
#include <stdint.h>
#include <string>
namespace zilliz {
namespace knowhere {
class ID {
public:
constexpr static int64_t kIDSize = 20;
public:
const int32_t *
data() const { return content_; }
int32_t *
mutable_data() { return content_; }
bool
IsValid() const;
std::string
ToString() const;
bool
operator==(const ID &that) const;
bool
operator<(const ID &that) const;
protected:
int32_t content_[5] = {};
};
} // namespace knowhere
} // namespace zilliz
// 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.
#pragma once
#include <memory>
#include "arrow/type.h"
namespace zilliz {
namespace knowhere {
using DataType = arrow::DataType;
using Field = arrow::Field;
using FieldPtr = std::shared_ptr<arrow::Field>;
using Schema = arrow::Schema;
using SchemaPtr = std::shared_ptr<Schema>;
using SchemaConstPtr = std::shared_ptr<const Schema>;
} // namespace knowhere
} // namespace zilliz
// 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.
#pragma once
#include <memory>
#include "arrow/tensor.h"
namespace zilliz {
namespace knowhere {
using Tensor = arrow::Tensor;
using TensorPtr = std::shared_ptr<Tensor>;
} // namespace knowhere
} // namespace zilliz
......@@ -18,7 +18,7 @@
#include <iostream> // TODO(linxj): using Log instead
#include "knowhere/common/timer.h"
#include "timer.h"
namespace zilliz {
namespace knowhere {
......
// 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.
#pragma once
#include <memory>
#include "knowhere/common/binary_set.h"
#include "knowhere/common/dataset.h"
#include "index_type.h"
#include "index_model.h"
#include "knowhere/index/preprocessor/preprocessor.h"
namespace zilliz {
namespace knowhere {
class Index {
public:
virtual BinarySet
Serialize() = 0;
virtual void
Load(const BinarySet &index_binary) = 0;
// @throw
virtual DatasetPtr
Search(const DatasetPtr &dataset, const Config &config) = 0;
public:
IndexType
idx_type() const { return idx_type_; }
void
set_idx_type(IndexType idx_type) { idx_type_ = idx_type; }
virtual void
set_preprocessor(PreprocessorPtr preprocessor) {}
virtual void
set_index_model(IndexModelPtr model) {}
private:
IndexType idx_type_;
};
using IndexPtr = std::shared_ptr<Index>;
} // namespace knowhere
} // namespace zilliz
// 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.
#pragma once
#include <memory>
#include "knowhere/common/binary_set.h"
namespace zilliz {
namespace knowhere {
class IndexModel {
public:
virtual BinarySet
Serialize() = 0;
virtual void
Load(const BinarySet &binary) = 0;
};
using IndexModelPtr = std::shared_ptr<IndexModel>;
} // namespace knowhere
} // namespace zilliz
// 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.
#pragma once
namespace zilliz {
namespace knowhere {
enum class IndexType {
kUnknown = 0,
kVecIdxBegin = 100,
kVecIVFFlat = kVecIdxBegin,
kVecIdxEnd,
};
} // namespace knowhere
} // namespace zilliz
// 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.
#pragma once
#include <memory>
#include "knowhere/common/dataset.h"
namespace zilliz {
namespace knowhere {
class Preprocessor {
public:
virtual DatasetPtr
Preprocess(const DatasetPtr &input) = 0;
};
using PreprocessorPtr = std::shared_ptr<Preprocessor>;
} // namespace knowhere
} // namespace zilliz
......@@ -17,10 +17,10 @@
#include "knowhere/common/exception.h"
#include "knowhere/index/vector_index/cloner.h"
#include "knowhere/index/vector_index/ivf.h"
#include "knowhere/index/vector_index/gpu_ivf.h"
#include "knowhere/index/vector_index/idmap.h"
#include "cloner.h"
#include "ivf.h"
#include "gpu_ivf.h"
#include "idmap.h"
namespace zilliz {
......
// 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 <sstream>
#include <SPTAG/AnnService/inc/Server/QueryParser.h>
......@@ -7,10 +24,10 @@
#undef mkdir
#include "knowhere/index/vector_index/cpu_kdt_rng.h"
#include "knowhere/index/vector_index/definitions.h"
#include "cpu_kdt_rng.h"
#include "definitions.h"
//#include "knowhere/index/preprocessor/normalize.h"
#include "knowhere/index/vector_index/kdt_parameters.h"
#include "kdt_parameters.h"
#include "knowhere/adapter/sptag.h"
#include "knowhere/common/exception.h"
......
// 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.
#pragma once
#include <cstdint>
#include <memory>
#include "knowhere/index/vector_index/vector_index.h"
#include "vector_index.h"
#include "knowhere/index/index_model.h"
#include <SPTAG/AnnService/inc/Core/VectorIndex.h>
......
// 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.
#pragma once
#include <string>
namespace zilliz {
namespace knowhere {
#define META_ROWS ("rows")
#define META_DIM ("dimension")
#define META_K ("k")
} // namespace knowhere
} // namespace zilliz
......@@ -27,9 +27,9 @@
#include "knowhere/common/exception.h"
#include "knowhere/index/vector_index/cloner.h"
#include "cloner.h"
#include "knowhere/adapter/faiss_adopt.h"
#include "knowhere/index/vector_index/gpu_ivf.h"
#include "gpu_ivf.h"
#include <algorithm>
......
// 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.
#pragma once
#include <faiss/gpu/StandardGpuResources.h>
......
......@@ -25,7 +25,7 @@
#include "knowhere/common/exception.h"
#include "knowhere/adapter/faiss_adopt.h"
#include "knowhere/index/vector_index/idmap.h"
#include "idmap.h"
namespace zilliz {
......
......@@ -25,13 +25,12 @@
#include <faiss/index_io.h>
#include <faiss/gpu/StandardGpuResources.h>
#include <faiss/gpu/GpuAutoTune.h>
#include <knowhere/index/vector_index/ivf.h>
#include "knowhere/common/exception.h"
#include "knowhere/index/vector_index/ivf.h"
#include "ivf.h"
#include "knowhere/adapter/faiss_adopt.h"
#include "knowhere/index/vector_index/gpu_ivf.h"
#include "gpu_ivf.h"
namespace zilliz {
......
......@@ -25,7 +25,7 @@
#include <faiss/AuxIndexStructures.h>
#include <faiss/Index.h>
#include "knowhere/index/vector_index/vector_index.h"
#include "vector_index.h"
namespace zilliz {
......
// 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 <mutex>
#include "knowhere/index/vector_index/kdt_parameters.h"
#include "kdt_parameters.h"
namespace zilliz {
......
// 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.
#pragma once
#include <string>
#include <vector>
namespace zilliz {
namespace knowhere {
using KDTParameter = std::pair<std::string, std::string>;
class KDTParameterManagement {
public:
const std::vector<KDTParameter> &
GetKDTParameters();
public:
static KDTParameterManagement &
GetInstance() {
static KDTParameterManagement instance;
return instance;
}
KDTParameterManagement(const KDTParameterManagement &) = delete;
KDTParameterManagement &operator=(const KDTParameterManagement &) = delete;
private:
KDTParameterManagement();
private:
std::vector<KDTParameter> kdt_parameters_;
};
} // namespace knowhere
} // namespace zilliz
......@@ -18,7 +18,7 @@
#pragma once
#include "knowhere/index/vector_index/nsg/nsg.h"
#include "nsg.h"
namespace zilliz {
namespace knowhere {
......
......@@ -22,7 +22,7 @@
#include <stack>
#include <omp.h>
#include "knowhere/index/vector_index/nsg/nsg.h"
#include "nsg.h"
#include "knowhere/common/exception.h"
#include "knowhere/common/timer.h"
#include "utils.h"
......
......@@ -18,7 +18,7 @@
#include <cstring>
#include "knowhere/index/vector_index/nsg/nsg_io.h"
#include "nsg_io.h"
namespace zilliz {
......
......@@ -23,7 +23,7 @@
#include <faiss/AutoTune.h>
#include "knowhere/index/vector_index/nsg/nsg.h"
#include "nsg.h"
#include "knowhere/common/config.h"
......
......@@ -16,12 +16,12 @@
// under the License.
#include "knowhere/index/vector_index/nsg_index.h"
#include "nsg_index.h"
#include "knowhere/index/vector_index/nsg/nsg.h"
#include "knowhere/index/vector_index/nsg/nsg_io.h"
#include "knowhere/index/vector_index/idmap.h"
#include "knowhere/index/vector_index/ivf.h"
#include "knowhere/index/vector_index/gpu_ivf.h"
#include "idmap.h"
#include "ivf.h"
#include "gpu_ivf.h"
#include "knowhere/adapter/faiss_adopt.h"
#include "knowhere/common/exception.h"
#include "knowhere/common/timer.h"
......
......@@ -18,7 +18,7 @@
#pragma once
#include "knowhere/index/vector_index/vector_index.h"
#include "vector_index.h"
namespace zilliz {
......
// 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.
#pragma once
#include <memory>
#include "knowhere/common/config.h"
#include "knowhere/common/dataset.h"
#include "knowhere/index/index.h"
#include "knowhere/index/preprocessor/preprocessor.h"
namespace zilliz {
namespace knowhere {
class VectorIndex;
using VectorIndexPtr = std::shared_ptr<VectorIndex>;
class VectorIndex : public Index {
public:
virtual PreprocessorPtr
BuildPreprocessor(const DatasetPtr &dataset, const Config &config) { return nullptr; }
virtual IndexModelPtr
Train(const DatasetPtr &dataset, const Config &config) { return nullptr; }
virtual void
Add(const DatasetPtr &dataset, const Config &config) = 0;
virtual void
Seal() = 0;
virtual VectorIndexPtr
Clone() = 0;
virtual int64_t
Count() = 0;
virtual int64_t
Dimension() = 0;
};
} // namespace knowhere
} // namespace zilliz
include_directories(${CORE_SOURCE_DIR}/thirdparty)
include_directories(${CORE_SOURCE_DIR}/thirdparty/SPTAG/AnnService)
include_directories(${CORE_SOURCE_DIR}/include)
include_directories(${CORE_SOURCE_DIR}/knowhere)
include_directories(${CORE_SOURCE_DIR}/thirdparty/jsoncons-0.126.0/include)
include_directories(/usr/local/cuda/include)
link_directories(/usr/local/cuda/lib64)
......@@ -25,13 +25,13 @@ set(basic_libs
#<IVF-TEST>
set(ivf_srcs
${CORE_SOURCE_DIR}/src/knowhere/index/vector_index/ivf.cpp
${CORE_SOURCE_DIR}/src/knowhere/index/vector_index/gpu_ivf.cpp
${CORE_SOURCE_DIR}/src/knowhere/index/vector_index/cloner.cpp
${CORE_SOURCE_DIR}/src/knowhere/index/vector_index/idmap.cpp
${CORE_SOURCE_DIR}/src/knowhere/adapter/structure.cpp
${CORE_SOURCE_DIR}/src/knowhere/common/exception.cpp
${CORE_SOURCE_DIR}/src/knowhere/common/timer.cpp
${CORE_SOURCE_DIR}/knowhere/knowhere/index/vector_index/ivf.cpp
${CORE_SOURCE_DIR}/knowhere/knowhere/index/vector_index/gpu_ivf.cpp
${CORE_SOURCE_DIR}/knowhere/knowhere/index/vector_index/cloner.cpp
${CORE_SOURCE_DIR}/knowhere/knowhere/index/vector_index/idmap.cpp
${CORE_SOURCE_DIR}/knowhere/knowhere/adapter/structure.cpp
${CORE_SOURCE_DIR}/knowhere/knowhere/common/exception.cpp
${CORE_SOURCE_DIR}/knowhere/knowhere/common/timer.cpp
utils.cpp
)
if(NOT TARGET test_ivf)
......@@ -41,13 +41,13 @@ target_link_libraries(test_ivf ${depend_libs} ${unittest_libs} ${basic_libs})
#<IDMAP-TEST>
set(idmap_srcs
${CORE_SOURCE_DIR}/src/knowhere/index/vector_index/idmap.cpp
${CORE_SOURCE_DIR}/src/knowhere/index/vector_index/ivf.cpp
${CORE_SOURCE_DIR}/src/knowhere/index/vector_index/cloner.cpp
${CORE_SOURCE_DIR}/src/knowhere/index/vector_index/gpu_ivf.cpp
${CORE_SOURCE_DIR}/src/knowhere/adapter/structure.cpp
${CORE_SOURCE_DIR}/src/knowhere/common/exception.cpp
${CORE_SOURCE_DIR}/src/knowhere/common/timer.cpp
${CORE_SOURCE_DIR}/knowhere/knowhere/index/vector_index/idmap.cpp
${CORE_SOURCE_DIR}/knowhere/knowhere/index/vector_index/ivf.cpp
${CORE_SOURCE_DIR}/knowhere/knowhere/index/vector_index/cloner.cpp
${CORE_SOURCE_DIR}/knowhere/knowhere/index/vector_index/gpu_ivf.cpp
${CORE_SOURCE_DIR}/knowhere/knowhere/adapter/structure.cpp
${CORE_SOURCE_DIR}/knowhere/knowhere/common/exception.cpp
${CORE_SOURCE_DIR}/knowhere/knowhere/common/timer.cpp
utils.cpp
)
if(NOT TARGET test_idmap)
......@@ -57,14 +57,14 @@ target_link_libraries(test_idmap ${depend_libs} ${unittest_libs} ${basic_libs})
#<KDT-TEST>
set(kdt_srcs
${CORE_SOURCE_DIR}/src/knowhere/index/preprocessor/normalize.cpp
${CORE_SOURCE_DIR}/src/knowhere/index/vector_index/kdt_parameters.cpp
${CORE_SOURCE_DIR}/src/knowhere/index/vector_index/cpu_kdt_rng.cpp
${CORE_SOURCE_DIR}/src/knowhere/adapter/structure.cpp
${CORE_SOURCE_DIR}/src/knowhere/adapter/sptag.cpp
${CORE_SOURCE_DIR}/src/knowhere/common/exception.cpp
${CORE_SOURCE_DIR}/src/knowhere/adapter/arrow.cpp
${CORE_SOURCE_DIR}/src/knowhere/common/timer.cpp
${CORE_SOURCE_DIR}/knowhere/knowhere/index/preprocessor/normalize.cpp
${CORE_SOURCE_DIR}/knowhere/knowhere/index/vector_index/kdt_parameters.cpp
${CORE_SOURCE_DIR}/knowhere/knowhere/index/vector_index/cpu_kdt_rng.cpp
${CORE_SOURCE_DIR}/knowhere/knowhere/adapter/structure.cpp
${CORE_SOURCE_DIR}/knowhere/knowhere/adapter/sptag.cpp
${CORE_SOURCE_DIR}/knowhere/knowhere/common/exception.cpp
${CORE_SOURCE_DIR}/knowhere/knowhere/adapter/arrow.cpp
${CORE_SOURCE_DIR}/knowhere/knowhere/common/timer.cpp
utils.cpp
)
if(NOT TARGET test_kdt)
......
......@@ -148,7 +148,9 @@ Status GetTableFilePath(const DBMetaOptions& options, meta::TableFileSchema& tab
}
std::string msg = "Table file doesn't exist: " + file_path;
ENGINE_LOG_ERROR << msg;
ENGINE_LOG_ERROR << msg << " in path: " << options.path
<< " for table: " << table_file.table_id_;
return Status(DB_ERROR, msg);
}
......
......@@ -23,8 +23,8 @@
#include "utils/CommonUtil.h"
#include "utils/Exception.h"
#include "wrapper/knowhere/vec_index.h"
#include "wrapper/knowhere/vec_impl.h"
#include "src/wrapper/vec_index.h"
#include "src/wrapper/vec_impl.h"
#include "knowhere/common/exception.h"
#include <stdexcept>
......
......@@ -18,7 +18,7 @@
#pragma once
#include "ExecutionEngine.h"
#include "wrapper/knowhere/vec_index.h"
#include "src/wrapper/vec_index.h"
#include <memory>
#include <string>
......
......@@ -32,7 +32,7 @@
#include <unistd.h>
#include <string.h>
#include <src/scheduler/SchedInst.h>
#include "wrapper/knowhere/KnowhereResource.h"
#include "src/wrapper/KnowhereResource.h"
#include "metrics/Metrics.h"
#include "DBWrapper.h"
......
// 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.
#pragma once
#include <cstdint>
#include <cstddef>
#include <limits>
#include <cstddef>
namespace zilliz {
namespace milvus {
namespace engine {
using Bool = int8_t;
using Byte = uint8_t;
using Word = unsigned long;
using EnumType = uint64_t;
using Float32 = float;
using Float64 = double;
constexpr bool kBoolMax = std::numeric_limits<bool>::max();
constexpr bool kBoolMin = std::numeric_limits<bool>::lowest();
constexpr int8_t kInt8Max = std::numeric_limits<int8_t>::max();
constexpr int8_t kInt8Min = std::numeric_limits<int8_t>::lowest();
constexpr int16_t kInt16Max = std::numeric_limits<int16_t>::max();
constexpr int16_t kInt16Min = std::numeric_limits<int16_t>::lowest();
constexpr int32_t kInt32Max = std::numeric_limits<int32_t>::max();
constexpr int32_t kInt32Min = std::numeric_limits<int32_t>::lowest();
constexpr int64_t kInt64Max = std::numeric_limits<int64_t>::max();
constexpr int64_t kInt64Min = std::numeric_limits<int64_t>::lowest();
constexpr float kFloatMax = std::numeric_limits<float>::max();
constexpr float kFloatMin = std::numeric_limits<float>::lowest();
constexpr double kDoubleMax = std::numeric_limits<double>::max();
constexpr double kDoubleMin = std::numeric_limits<double>::lowest();
constexpr uint32_t kFloat32DecimalPrecision = std::numeric_limits<Float32>::digits10;
constexpr uint32_t kFloat64DecimalPrecision = std::numeric_limits<Float64>::digits10;
constexpr uint8_t kByteWidth = 8;
constexpr uint8_t kCharWidth = kByteWidth;
constexpr uint8_t kWordWidth = sizeof(Word) * kByteWidth;
constexpr uint8_t kEnumTypeWidth = sizeof(EnumType) * kByteWidth;
template<typename T>
inline size_t
WidthOf() { return sizeof(T) << 3; }
template<typename T>
inline size_t
WidthOf(const T &) { return sizeof(T) << 3; }
}
} // namespace lib
} // namespace zilliz
// 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 "faiss/FaissAssert.h"
#include "faiss/gpu/utils/Limits.cuh"
#include "Arithmetic.h"
namespace faiss {
namespace gpu {
constexpr bool kBoolMax = zilliz::milvus::engine::kBoolMax;
constexpr bool kBoolMin = zilliz::milvus::engine::kBoolMin;
template<>
struct Limits<bool> {
static __device__ __host__
inline bool getMin() {
return kBoolMin;
}
static __device__ __host__
inline bool getMax() {
return kBoolMax;
}
};
constexpr int8_t kInt8Max = zilliz::milvus::engine::kInt8Max;
constexpr int8_t kInt8Min = zilliz::milvus::engine::kInt8Min;
template<>
struct Limits<int8_t> {
static __device__ __host__
inline int8_t getMin() {
return kInt8Min;
}
static __device__ __host__
inline int8_t getMax() {
return kInt8Max;
}
};
constexpr int16_t kInt16Max = zilliz::milvus::engine::kInt16Max;
constexpr int16_t kInt16Min = zilliz::milvus::engine::kInt16Min;
template<>
struct Limits<int16_t> {
static __device__ __host__
inline int16_t getMin() {
return kInt16Min;
}
static __device__ __host__
inline int16_t getMax() {
return kInt16Max;
}
};
constexpr int64_t kInt64Max = zilliz::milvus::engine::kInt64Max;
constexpr int64_t kInt64Min = zilliz::milvus::engine::kInt64Min;
template<>
struct Limits<int64_t> {
static __device__ __host__
inline int64_t getMin() {
return kInt64Min;
}
static __device__ __host__
inline int64_t getMax() {
return kInt64Max;
}
};
constexpr double kDoubleMax = zilliz::milvus::engine::kDoubleMax;
constexpr double kDoubleMin = zilliz::milvus::engine::kDoubleMin;
template<>
struct Limits<double> {
static __device__ __host__
inline double getMin() {
return kDoubleMin;
}
static __device__ __host__
inline double getMax() {
return kDoubleMax;
}
};
}
}
#include "faiss/gpu/utils/DeviceUtils.h"
#include "faiss/gpu/utils/MathOperators.cuh"
#include "faiss/gpu/utils/Pair.cuh"
#include "faiss/gpu/utils/Reductions.cuh"
#include "faiss/gpu/utils/Select.cuh"
#include "faiss/gpu/utils/Tensor.cuh"
#include "faiss/gpu/utils/StaticUtils.h"
#include "Topk.h"
namespace zilliz {
namespace milvus {
namespace engine {
namespace gpu {
constexpr int kWarpSize = 32;
template<typename T, int Dim, bool InnerContig>
using Tensor = faiss::gpu::Tensor<T, Dim, InnerContig>;
template<typename T, typename U>
using Pair = faiss::gpu::Pair<T, U>;
// select kernel for k == 1
template<typename T, int kRowsPerBlock, int kBlockSize>
__global__ void topkSelectMin1(Tensor<T, 2, true> productDistances,
Tensor<T, 2, true> outDistances,
Tensor<int64_t, 2, true> outIndices) {
// Each block handles kRowsPerBlock rows of the distances (results)
Pair<T, int64_t> threadMin[kRowsPerBlock];
__shared__
Pair<T, int64_t> blockMin[kRowsPerBlock * (kBlockSize / kWarpSize)];
T distance[kRowsPerBlock];
#pragma unroll
for (int i = 0; i < kRowsPerBlock; ++i) {
threadMin[i].k = faiss::gpu::Limits<T>::getMax();
threadMin[i].v = -1;
}
// blockIdx.x: which chunk of rows we are responsible for updating
int rowStart = blockIdx.x * kRowsPerBlock;
// FIXME: if we have exact multiples, don't need this
bool endRow = (blockIdx.x == gridDim.x - 1);
if (endRow) {
if (productDistances.getSize(0) % kRowsPerBlock == 0) {
endRow = false;
}
}
if (endRow) {
for (int row = rowStart; row < productDistances.getSize(0); ++row) {
for (int col = threadIdx.x; col < productDistances.getSize(1);
col += blockDim.x) {
distance[0] = productDistances[row][col];
if (faiss::gpu::Math<T>::lt(distance[0], threadMin[0].k)) {
threadMin[0].k = distance[0];
threadMin[0].v = col;
}
}
// Reduce within the block
threadMin[0] =
faiss::gpu::blockReduceAll<Pair<T, int64_t>, faiss::gpu::Min<Pair<T, int64_t> >, false, false>(
threadMin[0], faiss::gpu::Min<Pair<T, int64_t> >(), blockMin);
if (threadIdx.x == 0) {
outDistances[row][0] = threadMin[0].k;
outIndices[row][0] = threadMin[0].v;
}
// so we can use the shared memory again
__syncthreads();
threadMin[0].k = faiss::gpu::Limits<T>::getMax();
threadMin[0].v = -1;
}
} else {
for (int col = threadIdx.x; col < productDistances.getSize(1);
col += blockDim.x) {
#pragma unroll
for (int row = 0; row < kRowsPerBlock; ++row) {
distance[row] = productDistances[rowStart + row][col];
}
#pragma unroll
for (int row = 0; row < kRowsPerBlock; ++row) {
if (faiss::gpu::Math<T>::lt(distance[row], threadMin[row].k)) {
threadMin[row].k = distance[row];
threadMin[row].v = col;
}
}
}
// Reduce within the block
faiss::gpu::blockReduceAll<kRowsPerBlock, Pair<T, int64_t>, faiss::gpu::Min<Pair<T, int64_t> >, false, false>(
threadMin, faiss::gpu::Min<Pair<T, int64_t> >(), blockMin);
if (threadIdx.x == 0) {
#pragma unroll
for (int row = 0; row < kRowsPerBlock; ++row) {
outDistances[rowStart + row][0] = threadMin[row].k;
outIndices[rowStart + row][0] = threadMin[row].v;
}
}
}
}
// L2 + select kernel for k > 1, no re-use of ||c||^2
template<typename T, int NumWarpQ, int NumThreadQ, int ThreadsPerBlock>
__global__ void topkSelectMinK(Tensor<T, 2, true> productDistances,
Tensor<T, 2, true> outDistances,
Tensor<int64_t, 2, true> outIndices,
int k, T initK) {
// Each block handles a single row of the distances (results)
constexpr int kNumWarps = ThreadsPerBlock / kWarpSize;
__shared__
T smemK[kNumWarps * NumWarpQ];
__shared__
int64_t smemV[kNumWarps * NumWarpQ];
faiss::gpu::BlockSelect<T, int64_t, false, faiss::gpu::Comparator<T>,
NumWarpQ, NumThreadQ, ThreadsPerBlock>
heap(initK, -1, smemK, smemV, k);
int row = blockIdx.x;
// Whole warps must participate in the selection
int limit = faiss::gpu::utils::roundDown(productDistances.getSize(1), kWarpSize);
int i = threadIdx.x;
for (; i < limit; i += blockDim.x) {
T v = productDistances[row][i];
heap.add(v, i);
}
if (i < productDistances.getSize(1)) {
T v = productDistances[row][i];
heap.addThreadQ(v, i);
}
heap.reduce();
for (int i = threadIdx.x; i < k; i += blockDim.x) {
outDistances[row][i] = smemK[i];
outIndices[row][i] = smemV[i];
}
}
// FIXME: no TVec specialization
template<typename T>
void runTopKSelectMin(Tensor<T, 2, true> &productDistances,
Tensor<T, 2, true> &outDistances,
Tensor<int64_t, 2, true> &outIndices,
int k,
cudaStream_t stream) {
FAISS_ASSERT(productDistances.getSize(0) == outDistances.getSize(0));
FAISS_ASSERT(productDistances.getSize(0) == outIndices.getSize(0));
FAISS_ASSERT(outDistances.getSize(1) == k);
FAISS_ASSERT(outIndices.getSize(1) == k);
FAISS_ASSERT(k <= 1024);
if (k == 1) {
constexpr int kThreadsPerBlock = 256;
constexpr int kRowsPerBlock = 8;
auto block = dim3(kThreadsPerBlock);
auto grid = dim3(faiss::gpu::utils::divUp(outDistances.getSize(0), kRowsPerBlock));
topkSelectMin1<T, kRowsPerBlock, kThreadsPerBlock>
<< < grid, block, 0, stream >> > (productDistances, outDistances, outIndices);
} else {
constexpr int kThreadsPerBlock = 128;
auto block = dim3(kThreadsPerBlock);
auto grid = dim3(outDistances.getSize(0));
#define RUN_TOPK_SELECT_MIN(NUM_WARP_Q, NUM_THREAD_Q) \
do { \
topkSelectMinK<T, NUM_WARP_Q, NUM_THREAD_Q, kThreadsPerBlock> \
<<<grid, block, 0, stream>>>(productDistances, \
outDistances, outIndices, \
k, faiss::gpu::Limits<T>::getMax()); \
} while (0)
if (k <= 32) {
RUN_TOPK_SELECT_MIN(32, 2);
} else if (k <= 64) {
RUN_TOPK_SELECT_MIN(64, 3);
} else if (k <= 128) {
RUN_TOPK_SELECT_MIN(128, 3);
} else if (k <= 256) {
RUN_TOPK_SELECT_MIN(256, 4);
} else if (k <= 512) {
RUN_TOPK_SELECT_MIN(512, 8);
} else if (k <= 1024) {
RUN_TOPK_SELECT_MIN(1024, 8);
} else {
FAISS_ASSERT(false);
}
}
CUDA_TEST_ERROR();
}
////////////////////////////////////////////////////////////
// select kernel for k == 1
template<typename T, int kRowsPerBlock, int kBlockSize>
__global__ void topkSelectMax1(Tensor<T, 2, true> productDistances,
Tensor<T, 2, true> outDistances,
Tensor<int64_t, 2, true> outIndices) {
// Each block handles kRowsPerBlock rows of the distances (results)
Pair<T, int64_t> threadMax[kRowsPerBlock];
__shared__
Pair<T, int64_t> blockMax[kRowsPerBlock * (kBlockSize / kWarpSize)];
T distance[kRowsPerBlock];
#pragma unroll
for (int i = 0; i < kRowsPerBlock; ++i) {
threadMax[i].k = faiss::gpu::Limits<T>::getMin();
threadMax[i].v = -1;
}
// blockIdx.x: which chunk of rows we are responsible for updating
int rowStart = blockIdx.x * kRowsPerBlock;
// FIXME: if we have exact multiples, don't need this
bool endRow = (blockIdx.x == gridDim.x - 1);
if (endRow) {
if (productDistances.getSize(0) % kRowsPerBlock == 0) {
endRow = false;
}
}
if (endRow) {
for (int row = rowStart; row < productDistances.getSize(0); ++row) {
for (int col = threadIdx.x; col < productDistances.getSize(1);
col += blockDim.x) {
distance[0] = productDistances[row][col];
if (faiss::gpu::Math<T>::gt(distance[0], threadMax[0].k)) {
threadMax[0].k = distance[0];
threadMax[0].v = col;
}
}
// Reduce within the block
threadMax[0] =
faiss::gpu::blockReduceAll<Pair<T, int64_t>, faiss::gpu::Max<Pair<T, int64_t> >, false, false>(
threadMax[0], faiss::gpu::Max<Pair<T, int64_t> >(), blockMax);
if (threadIdx.x == 0) {
outDistances[row][0] = threadMax[0].k;
outIndices[row][0] = threadMax[0].v;
}
// so we can use the shared memory again
__syncthreads();
threadMax[0].k = faiss::gpu::Limits<T>::getMin();
threadMax[0].v = -1;
}
} else {
for (int col = threadIdx.x; col < productDistances.getSize(1);
col += blockDim.x) {
#pragma unroll
for (int row = 0; row < kRowsPerBlock; ++row) {
distance[row] = productDistances[rowStart + row][col];
}
#pragma unroll
for (int row = 0; row < kRowsPerBlock; ++row) {
if (faiss::gpu::Math<T>::gt(distance[row], threadMax[row].k)) {
threadMax[row].k = distance[row];
threadMax[row].v = col;
}
}
}
// Reduce within the block
faiss::gpu::blockReduceAll<kRowsPerBlock, Pair<T, int64_t>, faiss::gpu::Max<Pair<T, int64_t> >, false, false>(
threadMax, faiss::gpu::Max<Pair<T, int64_t> >(), blockMax);
if (threadIdx.x == 0) {
#pragma unroll
for (int row = 0; row < kRowsPerBlock; ++row) {
outDistances[rowStart + row][0] = threadMax[row].k;
outIndices[rowStart + row][0] = threadMax[row].v;
}
}
}
}
// L2 + select kernel for k > 1, no re-use of ||c||^2
template<typename T, int NumWarpQ, int NumThreadQ, int ThreadsPerBlock>
__global__ void topkSelectMaxK(Tensor<T, 2, true> productDistances,
Tensor<T, 2, true> outDistances,
Tensor<int64_t, 2, true> outIndices,
int k, T initK) {
// Each block handles a single row of the distances (results)
constexpr int kNumWarps = ThreadsPerBlock / kWarpSize;
__shared__
T smemK[kNumWarps * NumWarpQ];
__shared__
int64_t smemV[kNumWarps * NumWarpQ];
faiss::gpu::BlockSelect<T, int64_t, true, faiss::gpu::Comparator<T>,
NumWarpQ, NumThreadQ, ThreadsPerBlock>
heap(initK, -1, smemK, smemV, k);
int row = blockIdx.x;
// Whole warps must participate in the selection
int limit = faiss::gpu::utils::roundDown(productDistances.getSize(1), kWarpSize);
int i = threadIdx.x;
for (; i < limit; i += blockDim.x) {
T v = productDistances[row][i];
heap.add(v, i);
}
if (i < productDistances.getSize(1)) {
T v = productDistances[row][i];
heap.addThreadQ(v, i);
}
heap.reduce();
for (int i = threadIdx.x; i < k; i += blockDim.x) {
outDistances[row][i] = smemK[i];
outIndices[row][i] = smemV[i];
}
}
// FIXME: no TVec specialization
template<typename T>
void runTopKSelectMax(Tensor<T, 2, true> &productDistances,
Tensor<T, 2, true> &outDistances,
Tensor<int64_t, 2, true> &outIndices,
int k,
cudaStream_t stream) {
FAISS_ASSERT(productDistances.getSize(0) == outDistances.getSize(0));
FAISS_ASSERT(productDistances.getSize(0) == outIndices.getSize(0));
FAISS_ASSERT(outDistances.getSize(1) == k);
FAISS_ASSERT(outIndices.getSize(1) == k);
FAISS_ASSERT(k <= 1024);
if (k == 1) {
constexpr int kThreadsPerBlock = 256;
constexpr int kRowsPerBlock = 8;
auto block = dim3(kThreadsPerBlock);
auto grid = dim3(faiss::gpu::utils::divUp(outDistances.getSize(0), kRowsPerBlock));
topkSelectMax1<T, kRowsPerBlock, kThreadsPerBlock>
<< < grid, block, 0, stream >> > (productDistances, outDistances, outIndices);
} else {
constexpr int kThreadsPerBlock = 128;
auto block = dim3(kThreadsPerBlock);
auto grid = dim3(outDistances.getSize(0));
#define RUN_TOPK_SELECT_MAX(NUM_WARP_Q, NUM_THREAD_Q) \
do { \
topkSelectMaxK<T, NUM_WARP_Q, NUM_THREAD_Q, kThreadsPerBlock> \
<<<grid, block, 0, stream>>>(productDistances, \
outDistances, outIndices, \
k, faiss::gpu::Limits<T>::getMin()); \
} while (0)
if (k <= 32) {
RUN_TOPK_SELECT_MAX(32, 2);
} else if (k <= 64) {
RUN_TOPK_SELECT_MAX(64, 3);
} else if (k <= 128) {
RUN_TOPK_SELECT_MAX(128, 3);
} else if (k <= 256) {
RUN_TOPK_SELECT_MAX(256, 4);
} else if (k <= 512) {
RUN_TOPK_SELECT_MAX(512, 8);
} else if (k <= 1024) {
RUN_TOPK_SELECT_MAX(1024, 8);
} else {
FAISS_ASSERT(false);
}
}
CUDA_TEST_ERROR();
}
//////////////////////////////////////////////////////////////
template<typename T>
void runTopKSelect(Tensor<T, 2, true> &productDistances,
Tensor<T, 2, true> &outDistances,
Tensor<int64_t, 2, true> &outIndices,
bool dir,
int k,
cudaStream_t stream) {
if (dir) {
runTopKSelectMax<T>(productDistances,
outDistances,
outIndices,
k,
stream);
} else {
runTopKSelectMin<T>(productDistances,
outDistances,
outIndices,
k,
stream);
}
}
template<typename T>
void TopK(T *input,
int length,
int k,
T *output,
int64_t *idx,
// Ordering order_flag,
cudaStream_t stream) {
// bool dir = (order_flag == Ordering::kAscending ? false : true);
bool dir = 0;
Tensor<T, 2, true> t_input(input, {1, length});
Tensor<T, 2, true> t_output(output, {1, k});
Tensor<int64_t, 2, true> t_idx(idx, {1, k});
runTopKSelect<T>(t_input, t_output, t_idx, dir, k, stream);
}
//INSTANTIATION_TOPK_2(bool);
//INSTANTIATION_TOPK_2(int8_t);
//INSTANTIATION_TOPK_2(int16_t);
INSTANTIATION_TOPK_2(int32_t);
//INSTANTIATION_TOPK_2(int64_t);
INSTANTIATION_TOPK_2(float);
//INSTANTIATION_TOPK_2(double);
//INSTANTIATION_TOPK(TimeInterval);
//INSTANTIATION_TOPK(Float128);
//INSTANTIATION_TOPK(char);
}
void TopK(float *host_input,
int length,
int k,
float *output,
int64_t *indices) {
float *device_input, *device_output;
int64_t *ids;
cudaMalloc((void **) &device_input, sizeof(float) * length);
cudaMalloc((void **) &device_output, sizeof(float) * k);
cudaMalloc((void **) &ids, sizeof(int64_t) * k);
cudaMemcpy(device_input, host_input, sizeof(float) * length, cudaMemcpyHostToDevice);
gpu::TopK<float>(device_input, length, k, device_output, ids, nullptr);
cudaMemcpy(output, device_output, sizeof(float) * k, cudaMemcpyDeviceToHost);
cudaMemcpy(indices, ids, sizeof(int64_t) * k, cudaMemcpyDeviceToHost);
cudaFree(device_input);
cudaFree(device_output);
cudaFree(ids);
}
}
}
}
// 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.
#pragma once
#include <cuda.h>
#include <cuda_runtime.h>
namespace zilliz {
namespace milvus {
namespace engine {
namespace gpu {
template<typename T>
void
TopK(T *input,
int length,
int k,
T *output,
int64_t *indices,
// Ordering order_flag,
cudaStream_t stream = nullptr);
#define INSTANTIATION_TOPK_2(T) \
template void \
TopK<T>(T *input, \
int length, \
int k, \
T *output, \
int64_t *indices, \
cudaStream_t stream)
// Ordering order_flag, \
// cudaStream_t stream)
//extern INSTANTIATION_TOPK_2(int8_t);
//extern INSTANTIATION_TOPK_2(int16_t);
extern INSTANTIATION_TOPK_2(int32_t);
//extern INSTANTIATION_TOPK_2(int64_t);
extern INSTANTIATION_TOPK_2(float);
//extern INSTANTIATION_TOPK_2(double);
//extern INSTANTIATION_TOPK(TimeInterval);
//extern INSTANTIATION_TOPK(Float128);
}
// User Interface.
void TopK(float *input,
int length,
int k,
float *output,
int64_t *indices);
}
}
}
......@@ -66,7 +66,7 @@ aux_source_directory(${MILVUS_ENGINE_SRC}/server/grpc_impl grpc_server_files)
aux_source_directory(${MILVUS_ENGINE_SRC}/utils utils_files)
aux_source_directory(${MILVUS_ENGINE_SRC}/wrapper/knowhere wrapper_knowhere_files)
aux_source_directory(${MILVUS_ENGINE_SRC}/wrapper wrapper_files)
set(unittest_files
${CMAKE_CURRENT_SOURCE_DIR}/main.cpp)
......@@ -90,7 +90,7 @@ set(common_files
${db_scheduler_files}
${metrics_files}
${scheduler_files}
${wrapper_knowhere_files}
${wrapper_files}
${helper_files}
)
......
......@@ -25,7 +25,7 @@
#include "scheduler/ResourceFactory.h"
#include "scheduler/resource/Resource.h"
#include "utils/Error.h"
#include "wrapper/knowhere/vec_index.h"
#include "src/wrapper/vec_index.h"
#include "scheduler/tasklabel/SpecResLabel.h"
......
......@@ -21,7 +21,7 @@
#include "server/ServerConfig.h"
#include "utils/Error.h"
#include "wrapper/knowhere/vec_index.h"
#include "src/wrapper/vec_index.h"
using namespace zilliz::milvus;
......
......@@ -20,10 +20,10 @@
include_directories("${CUDA_TOOLKIT_ROOT_DIR}/include")
link_directories("${CUDA_TOOLKIT_ROOT_DIR}/lib64")
set(wrapper_knowhere_files
${MILVUS_ENGINE_SRC}/wrapper/knowhere/data_transfer.cpp
${MILVUS_ENGINE_SRC}/wrapper/knowhere/vec_impl.cpp
${MILVUS_ENGINE_SRC}/wrapper/knowhere/vec_index.cpp)
set(wrapper_files
${MILVUS_ENGINE_SRC}/wrapper/data_transfer.cpp
${MILVUS_ENGINE_SRC}/wrapper/vec_impl.cpp
${MILVUS_ENGINE_SRC}/wrapper/vec_index.cpp)
set(util_files
utils.cpp
......@@ -36,7 +36,7 @@ set(knowhere_libs
cublas
)
add_executable(wrapper_test wrapper_test.cpp ${wrapper_knowhere_files} ${util_files})
add_executable(wrapper_test wrapper_test.cpp ${wrapper_files} ${util_files})
target_link_libraries(wrapper_test ${knowhere_libs} ${unittest_libs})
install(TARGETS wrapper_test DESTINATION unittest)
\ No newline at end of file
......@@ -19,7 +19,7 @@
#include <gtest/gtest.h>
#include "utils/easylogging++.h"
#include <wrapper/knowhere/vec_index.h>
#include "src/wrapper/vec_index.h"
#include "knowhere/index/vector_index/gpu_ivf.h"
#include "utils.h"
......
Markdown is supported
0% .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑!
想要评论请 注册