huawei
/
mindspore2022
0
0
Fork 53
Code Issues Pull Requests 29 Projects Releases Wiki Activity

!32107 [MSLITE] add compile for cuda op 

Merge pull request !32107 from Liu_Xuu/cu_0328_compile
This commit is contained in:
i-robot 2022-04-06 06:36:52 +00:00 committed by Gitee
parent edf77ace25 d1b74fa756
commit 80218e3ffa
No known key found for this signature in database
GPG Key ID: 173E9B9CA92EEF8F
18 changed files with 530 additions and 21 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

17
mindspore/lite/CMakeLists.txt
View File

@ -234,7 +234,7 @@ else()
set(MSLITE_GPU_BACKEND "off")
endif()
if((NOT MSLITE_GPU_BACKEND STREQUAL "tensorrt") AND (NOT MSLITE_GPU_BACKEND STREQUAL "off") AND
(NOT MSLITE_GPU_BACKEND STREQUAL "opencl"))
(NOT MSLITE_GPU_BACKEND STREQUAL "cuda") AND (NOT MSLITE_GPU_BACKEND STREQUAL "opencl"))
message("invalid MSLITE_GPU_BACKEND value ${MSLITE_GPU_BACKEND} for x86_64, MSLITE_GPU_BACKEND is set to off.")
set(MSLITE_GPU_BACKEND "off")
endif()
@ -344,6 +344,12 @@ endif()
if(MSVC)
set(MSLITE_ENABLE_CONVERTER off)
endif()
if(MSLITE_GPU_BACKEND STREQUAL cuda)
set(MSLITE_ENABLE_CONVERTER on)
set(MSLITE_ENABLE_RUNTIME_GLOG on)
endif()
if(MSLITE_ENABLE_CONVERTER)
add_compile_definitions(ENABLE_CONVERTER)
endif()
@ -562,9 +568,16 @@ if(MSLITE_GPU_BACKEND STREQUAL opencl)
endif()
if(MSLITE_GPU_BACKEND STREQUAL cuda)
add_definitions(-DGPU_CUDA)
# add_definitions(-DGPU_CUDA)
add_compile_definitions(SUPPORT_GPU)
set(SUPPORT_CUDA on)
if(DEFINED ENV{CUDA_HOME})
message("CUDA_HOME = $ENV{CUDA_HOME}")
else()
message(FATAL_ERROR "please set CUDA_HOME, example: export CUDA_HOME=/usr/local/cuda-10.1/")
endif()
endif()
if(MSLITE_GPU_BACKEND STREQUAL tensorrt)
add_compile_definitions(SUPPORT_GPU)
set(SUPPORT_TENSORRT on)

19
mindspore/lite/src/CMakeLists.txt
View File

@ -423,10 +423,6 @@ if(MSLITE_GPU_BACKEND STREQUAL opencl)
add_subdirectory(runtime/kernel/opencl)
target_link_libraries(mindspore-lite opencl_kernel_mid)
target_link_libraries(mindspore-lite_static opencl_kernel_mid)
elseif(MSLITE_GPU_BACKEND STREQUAL cuda)
add_subdirectory(runtime/kernel/cuda)
target_link_libraries(mindspore-lite cuda_kernel_mid)
target_link_libraries(mindspore-lite_static cuda_kernel_mid)
endif()
if(SUPPORT_NPU)
@ -498,6 +494,21 @@ if(MSLITE_ENABLE_RUNTIME_CONVERT)
endif()
endif()
if(SUPPORT_CUDA)
set(CUDA_PATH $ENV{CUDA_HOME})
include_directories(${CCSRC_DIR}/plugin/device/gpu/kernel)
set(ENABLE_GPU on)
add_definitions(-DENABLE_GPU)
set(CUDA_VERSION 11.1)
include_directories(${CUDA_PATH})
include_directories(${CUDA_PATH}/include)
find_package(CUDA)
add_subdirectory(runtime/kernel/cuda)
list(APPEND CUDA_NVCC_FLAGS -arch=sm_53 --expt-relaxed-constexpr)
add_subdirectory(${CCSRC_DIR}/plugin/device/gpu/kernel/cuda_impl/cuda_ops cuda_ops)
target_link_libraries(mindspore-lite cuda_lite_kernel_mid cuda_ops mindspore_core)
endif()
if(MSLITE_ENABLE_RUNTIME_GLOG)
target_link_libraries(mindspore-lite mindspore::glog)
target_link_libraries(mindspore-lite_static mindspore::glog)

1
mindspore/lite/src/delegate/tensorrt/CMakeLists.txt
View File

@ -82,6 +82,7 @@ find_package(CUDA)
file(GLOB_RECURSE CUDA_KERNEL_SRC
${CMAKE_CURRENT_SOURCE_DIR}/cuda_impl/*.cu
)
set_source_files_properties(${CUDA_KERNEL_SRC} PROPERTIES CUDA_SOURCE_PROPERTY_FORMAT OBJ)
SET(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGES} -std=c++14 -fPIC")
SET(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-std=c++14;)

1
mindspore/lite/src/delegate/tensorrt/cuda_impl/cublas_utils.h
View File

@ -19,6 +19,7 @@
#include <cublas_v2.h>
#include "src/delegate/tensorrt/cuda_impl/cuda_helper.h"
#include "src/common/log_util.h"
// cublas API error checking
#define CUBLAS_CHECK_VOID(err) \

1
mindspore/lite/src/delegate/tensorrt/cuda_impl/cuda_helper.cc
View File

@ -15,6 +15,7 @@
*/
#include "src/delegate/tensorrt/cuda_impl/cuda_helper.h"
#include "src/common/log_util.h"
CudaHelper &CudaHelper::GetInstance() {
static CudaHelper instance;

27
mindspore/lite/src/delegate/tensorrt/cuda_impl/cuda_helper.h
View File

@ -19,7 +19,6 @@
#include <cuda_runtime.h>
#include <algorithm>
#include "src/common/log_util.h"
class CudaHelper {
public:
@ -39,22 +38,20 @@ class CudaHelper {
#define GET_BLOCKS(total_threads) CudaHelper::GetInstance().GetBlocksNum(total_threads)
#define GET_THREADS CudaHelper::GetInstance().GetThreadNum()
#define CUDA_CHECK(ret) \
do { \
cudaError_t cuda_ret = (ret); \
if ((cuda_ret) != cudaSuccess) { \
MS_LOG(ERROR) << "cuda func call error: " << cudaGetErrorString(cuda_ret); \
return -1; \
} \
#define CUDA_CHECK(ret) \
do { \
cudaError_t cuda_ret = (ret); \
if ((cuda_ret) != cudaSuccess) { \
return -1; \
} \
} while (0)
#define CUDA_CHECK_VOID(ret) \
do { \
cudaError_t cuda_ret = (ret); \
if ((cuda_ret) != cudaSuccess) { \
MS_LOG(ERROR) << "cuda func call error: " << cudaGetErrorString(cuda_ret); \
return; \
} \
#define CUDA_CHECK_VOID(ret) \
do { \
cudaError_t cuda_ret = (ret); \
if ((cuda_ret) != cudaSuccess) { \
return; \
} \
} while (0)
#endif // MINDSPORE_LITE_SRC_DELEGATE_TENSORRT_CDUA_IMPL_CUDA_HELPER_H_

1
mindspore/lite/src/delegate/tensorrt/cuda_impl/cudnn_utils.h
View File

@ -20,6 +20,7 @@
#include <cudnn.h>
#include <NvInfer.h>
#include "src/delegate/tensorrt/cuda_impl/cuda_helper.h"
#include "src/common/log_util.h"
#define CUDNN_CHECK_VOID(err) \
do { \

6
mindspore/lite/src/runtime/kernel/cuda/CMakeLists.txt Normal file
View File

@ -0,0 +1,6 @@
file(GLOB CUDA_LITE_KERNEL_SRC LIST_DIRECTORIES false
${CMAKE_CURRENT_SOURCE_DIR}/*.cc
${CCSRC_DIR}/plugin/device/gpu/kernel/cuda_impl/cuda_class/*.cc
)
add_library(cuda_lite_kernel_mid OBJECT ${CUDA_LITE_KERNEL_SRC})

47
mindspore/lite/src/runtime/kernel/cuda/batchtospace.cc Normal file
View File

@ -0,0 +1,47 @@
/**
* Copyright 2022 Huawei Technologies Co., Ltd
*
* Licensed 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 "src/runtime/kernel/cuda/batchtospace.h"
#include <memory>
#include "nnacl/batch_to_space.h"
namespace mindspore::kernel {
int BatchtoSpaceCudaKernel::Prepare() {
CudaKernel::Prepare();
if (batch_to_space_helper_ == nullptr) {
batch_to_space_helper_ = std::make_shared<cukernel::BatchToSpaceHelperGpuKernel<float>>(type_name_);
helper_ = batch_to_space_helper_;
}
cukernel::BatchToSpaceAttr attr;
auto param = reinterpret_cast<BatchToSpaceParameter *>(op_parameter_);
attr.block_size = param->block_shape_[0];
attr.crops.push_back({param->crops_[0], param->crops_[1]});
attr.crops.push_back({param->crops_[2], param->crops_[3]});
int ret = batch_to_space_helper_->CheckKernelParam(&attr);
CHECK_NOT_EQUAL_RETURN(ret, RET_OK);
ret = ReSize();
CHECK_NOT_EQUAL_RETURN(ret, RET_OK);
return RET_OK;
}
int BatchtoSpaceCudaKernel::Run() {
int ret = batch_to_space_helper_->Process(input_device_ptrs_, output_device_ptrs_, work_device_ptrs_, stream_);
CHECK_NOT_EQUAL_RETURN(ret, RET_OK);
return RET_OK;
}
// REG_KERNEL(kGPU, kNumberTypeFloat32, PrimitiveType_BatchToSpace, CudaKernelCreator<BatchtoSpaceCudaKernel>)
} // namespace mindspore::kernel

39
mindspore/lite/src/runtime/kernel/cuda/batchtospace.h Normal file
View File

@ -0,0 +1,39 @@
/**
* Copyright 2022 Huawei Technologies Co., Ltd
*
* Licensed 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.
*/
#ifndef MINDSPORE_LITE_SRC_RUNTIME_KERNEL_CUDA_UNARY_H_
#define MINDSPORE_LITE_SRC_RUNTIME_KERNEL_CUDA_UNARY_H_
#include <memory>
#include <vector>
#include "src/runtime/kernel/cuda/cuda_kernel.h"
#include "cuda_impl/cuda_class/batchtospace_helper.h"
namespace mindspore::kernel {
class BatchtoSpaceCudaKernel : public CudaKernel {
public:
BatchtoSpaceCudaKernel(OpParameter *parameter, const std::vector<lite::Tensor *> &inputs,
const std::vector<lite::Tensor *> &outputs, const lite::InnerContext *ctx)
: CudaKernel(parameter, inputs, outputs, ctx) {}
~BatchtoSpaceCudaKernel() override = default;
int Prepare() override;
int Run() override;
private:
std::shared_ptr<cukernel::BatchToSpaceHelperGpuKernel<float>> batch_to_space_helper_{nullptr};
};
} // namespace mindspore::kernel
#endif

72
mindspore/lite/src/runtime/kernel/cuda/cuda_kernel.cc Normal file
View File

@ -0,0 +1,72 @@
/**
* Copyright 2022 Huawei Technologies Co., Ltd
*
* Licensed 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 "src/runtime/kernel/cuda/cuda_kernel.h"
namespace mindspore::kernel {
int CudaKernel::PreProcess() {
if (output_device_size_.size() == 0) {
for (size_t i = 0; i < out_tensors_.size(); i++) {
// allocator cudaMalloc mem_size: out_tensors_[i]->set_allocator(/*CudaAllocator*/)
output_device_size_.push_back(helper_->GetOutputSizeList()[i]);
output_device_ptrs_.push_back(out_tensors_[i]->MutableData());
}
} else {
for (size_t i = 0; i < out_tensors_.size(); i++) {
if (helper_->GetOutputSizeList()[i] > output_device_size_[i]) {
out_tensors_[i]->FreeData();
output_device_size_[i] = helper_->GetOutputSizeList()[i];
output_device_ptrs_[i] = out_tensors_[i]->MutableData();
}
}
}
for (size_t i = 0; i < in_tensors_.size(); i++) {
input_device_ptrs_[i] = in_tensors_[i]->data();
}
return RET_OK;
}
int CudaKernel::ReSize() {
// menory calculate
std::vector<std::vector<size_t>> input_shapes;
std::vector<std::vector<size_t>> output_shapes;
for (auto in : in_tensors_) {
std::vector<size_t> one_shape(in->shape().size());
for (size_t i = 0; i < in->shape().size(); i++) {
one_shape[i] = static_cast<size_t>(in->shape()[i]);
}
input_shapes.push_back(one_shape);
}
for (auto out : out_tensors_) {
std::vector<size_t> one_shape(out->shape().size());
for (size_t i = 0; i < out->shape().size(); i++) {
one_shape[i] = static_cast<size_t>(out->shape()[i]);
}
output_shapes.push_back(one_shape);
}
helper_->ResetResource();
auto ret = helper_->CalMemSize(input_shapes, output_shapes);
CHECK_NOT_EQUAL_RETURN(ret, RET_OK);
return RET_OK;
}
int CudaKernel::PostProcess() {
for (size_t i = 0; i < in_tensors_.size(); i++) {
in_tensors_[i]->DecRefCount();
}
return RET_OK;
}
CudaKernel::~CudaKernel() { helper_ = nullptr; }
} // namespace mindspore::kernel

62
mindspore/lite/src/runtime/kernel/cuda/cuda_kernel.h Normal file
View File

@ -0,0 +1,62 @@
/**
* Copyright 2022 Huawei Technologies Co., Ltd
*
* Licensed 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.
*/
#ifndef MINDSPORE_LITE_SRC_RUNTIME_KERNEL_CUDA_CUDA_KERNEL_H_
#define MINDSPORE_LITE_SRC_RUNTIME_KERNEL_CUDA_CUDA_KERNEL_H_
#include <cuda_runtime_api.h>
#include <vector>
#include <string>
#include <memory>
#include "src/inner_kernel.h"
#include "src/lite_kernel.h"
#include "cuda_impl/cuda_class/helper_base.h"
using mindspore::lite::RET_ERROR;
using mindspore::lite::RET_OK;
namespace mindspore::kernel {
class CudaKernel : public InnerKernel {
public:
CudaKernel(OpParameter *parameter, const std::vector<lite::Tensor *> &inputs,
const std::vector<lite::Tensor *> &outputs, const lite::InnerContext *ctx)
: InnerKernel(parameter, inputs, outputs, ctx) {}
~CudaKernel() override;
int Prepare() override {
type_name_ = std::string(EnumNamePrimitiveType(type()));
return RET_OK;
}
int PreProcess() override;
int PostProcess() override;
int ReSize() override;
int Run() override { return RET_ERROR; }
protected:
std::vector<size_t> output_device_size_;
std::vector<void *> input_device_ptrs_;
std::vector<void *> output_device_ptrs_;
std::vector<void *> work_device_ptrs_;
cudaStream_t stream_;
std::shared_ptr<cukernel::GpuKernelHelperBase> helper_{nullptr};
std::string type_name_;
};
template <class T>
kernel::InnerKernel *CudaKernelCreator(const std::vector<lite::Tensor *> &inputs,
const std::vector<lite::Tensor *> &outputs, OpParameter *opParameter,
const lite::Context *ctx, const kernel::KernelKey &desc) {
return new (std::nothrow) T(opParameter, inputs, outputs, static_cast<const lite::InnerContext *>(ctx));
}
} // namespace mindspore::kernel
#endif

37
mindspore/lite/src/runtime/kernel/cuda/unary.cc Normal file
View File

@ -0,0 +1,37 @@
/**
* Copyright 2022 Huawei Technologies Co., Ltd
*
* Licensed 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 "src/runtime/kernel/cuda/unary.h"
#include <memory>
namespace mindspore::kernel {
int UnaryCudaKernel::Prepare() {
CudaKernel::Prepare();
if (unary_helper_ == nullptr) {
unary_helper_ = std::make_shared<cukernel::UnaryHelperGpuKernel<float>>(type_name_);
helper_ = unary_helper_;
}
int ret = ReSize();
CHECK_NOT_EQUAL_RETURN(ret, RET_OK);
return RET_OK;
}
int UnaryCudaKernel::Run() {
int ret = unary_helper_->Process(input_device_ptrs_, output_device_ptrs_, work_device_ptrs_, stream_);
CHECK_NOT_EQUAL_RETURN(ret, RET_OK);
return RET_OK;
}
// REG_KERNEL(kGPU, kNumberTypeFloat32, PrimitiveType_Log, CudaKernelCreator<UnaryCudaKernel>)
} // namespace mindspore::kernel

39
mindspore/lite/src/runtime/kernel/cuda/unary.h Normal file
View File

@ -0,0 +1,39 @@
/**
* Copyright 2022 Huawei Technologies Co., Ltd
*
* Licensed 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.
*/
#ifndef MINDSPORE_LITE_SRC_RUNTIME_KERNEL_CUDA_UNARY_H_
#define MINDSPORE_LITE_SRC_RUNTIME_KERNEL_CUDA_UNARY_H_
#include <memory>
#include <vector>
#include "src/runtime/kernel/cuda/cuda_kernel.h"
#include "cuda_impl/cuda_class/unary_helper.h"
namespace mindspore::kernel {
class UnaryCudaKernel : public CudaKernel {
public:
UnaryCudaKernel(OpParameter *parameter, const std::vector<lite::Tensor *> &inputs,
const std::vector<lite::Tensor *> &outputs, const lite::InnerContext *ctx)
: CudaKernel(parameter, inputs, outputs, ctx) {}
~UnaryCudaKernel() override = default; // cudaFree
int Prepare() override;
int Run() override;
private:
std::shared_ptr<cukernel::UnaryHelperGpuKernel<float>> unary_helper_{nullptr};
};
} // namespace mindspore::kernel
#endif

47
mindspore/lite/src/runtime/kernel/cuda/unique.cc Normal file
View File

@ -0,0 +1,47 @@
/**
* Copyright 2022 Huawei Technologies Co., Ltd
*
* Licensed 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 "src/runtime/kernel/cuda/unique.h"
#include <memory>
#include <vector>
namespace mindspore::kernel {
int UniqueCudaKernel::Prepare() {
CudaKernel::Prepare();
if (unique_helper_ == nullptr) {
unique_helper_ = std::make_shared<cukernel::UniqueHelperGpuKernel<int, int>>(type_name_);
helper_ = unique_helper_;
}
return RET_OK;
}
int UniqueCudaKernel::PostProcess() {
auto ret = CudaKernel::PostProcess();
CHECK_NOT_EQUAL_RETURN(ret, RET_OK);
// set output tensor shape
std::vector<int> out_shape = out_tensors_[0]->shape();
out_shape[out_shape.size() - 1] = unique_helper_->GetOutSize();
out_tensors_[0]->set_shape(out_shape);
return RET_OK;
}
int UniqueCudaKernel::Run() {
int ret = unique_helper_->Process(input_device_ptrs_, output_device_ptrs_, work_device_ptrs_, stream_);
CHECK_NOT_EQUAL_RETURN(ret, RET_OK);
return RET_OK;
}
// REG_KERNEL(kGPU, kNumberTypeFloat32, PrimitiveType_Unique, CudaKernelCreator<UniqueCudaKernel>)
} // namespace mindspore::kernel

40
mindspore/lite/src/runtime/kernel/cuda/unique.h Normal file
View File

@ -0,0 +1,40 @@
/**
* Copyright 2022 Huawei Technologies Co., Ltd
*
* Licensed 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.
*/
#ifndef MINDSPORE_LITE_SRC_RUNTIME_KERNEL_CUDA_UNARY_H_
#define MINDSPORE_LITE_SRC_RUNTIME_KERNEL_CUDA_UNARY_H_
#include <memory>
#include <vector>
#include "src/runtime/kernel/cuda/cuda_kernel.h"
#include "cuda_impl/cuda_class/unique_helper.h"
namespace mindspore::kernel {
class UniqueCudaKernel : public CudaKernel {
public:
UniqueCudaKernel(OpParameter *parameter, const std::vector<lite::Tensor *> &inputs,
const std::vector<lite::Tensor *> &outputs, const lite::InnerContext *ctx)
: CudaKernel(parameter, inputs, outputs, ctx) {}
~UniqueCudaKernel() override = default;
int Prepare() override;
int PostProcess() override;
int Run() override;
private:
std::shared_ptr<cukernel::UniqueHelperGpuKernel<int, int>> unique_helper_{nullptr};
};
} // namespace mindspore::kernel
#endif

11
mindspore/lite/test/CMakeLists.txt
View File

@ -72,6 +72,17 @@ if(MSLITE_GPU_BACKEND STREQUAL opencl)
endif()
endif()
if(MSLITE_GPU_BACKEND STREQUAL cuda)
set(CUDA_PATH $ENV{CUDA_HOME})
include_directories(${CUDA_PATH}/include)
include_directories(${CCSRC_DIR}/plugin/device/gpu/kernel)
include_directories(${CUDA_PATH})
file(GLOB_RECURSE TEST_GPU_UT_SRC
${TEST_DIR}/ut/src/runtime/kernel/cuda/*.cc
)
list(APPEND TEST_UT_SRC ${TEST_GPU_UT_SRC})
endif()
if(MSLITE_ENABLE_INT8)
file(GLOB_RECURSE TEST_INT8_UT_SRC
${TEST_DIR}/ut/src/runtime/kernel/arm/int8/*.cc

84
mindspore/lite/test/ut/src/runtime/kernel/cuda/batchtospace_tests.cc Normal file
View File

@ -0,0 +1,84 @@
/**
* Copyright 2022 Huawei Technologies Co., Ltd
*
* Licensed 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 "common/common_test.h"
#include "schema/ops_generated.h"
#include "src/runtime/kernel/cuda/batchtospace.h"
#include "ut/src/runtime/kernel/cuda/common.h"
#include "nnacl/batch_to_space.h"
namespace mindspore {
class CudaTest_BatchToSpace : public CommonTest {
public:
CudaTest_BatchToSpace() {}
};
namespace {
// input: [batch*block_size*block_size, height_pad/block_size, width_pad/block_size, depth]
// crops = [[crop_top, crop_bottom], [crop_left, crop_right]]
// height = height_pad - crop_top - crop_bottom
// width = width_pad - crop_left - crop_right
OpParameter *CreateParameter() {
auto *param = mindspore::lite::cuda::test::CreateParameter<BatchToSpaceParameter>(schema::PrimitiveType_BatchToSpace);
param->block_shape_[0] = 2;
param->block_shape_[1] = 2;
param->crops_[0] = 0;
param->crops_[1] = 0;
param->crops_[THIRD_INPUT] = 0;
param->crops_[FOURTH_INPUT] = 0;
param->no_crop_ = true;
return reinterpret_cast<OpParameter *>(param);
}
} // namespace
TEST_F(CudaTest_BatchToSpace, basic) {
std::vector<int> input_shape{4, 1, 1, 1};
lite::Tensor *in_tensor = new (std::nothrow) lite::Tensor(TypeId::kNumberTypeFloat32, input_shape);
std::vector<lite::Tensor *> inputs{in_tensor};
std::vector<int> output_shape{1, 2, 2, 1};
lite::Tensor *out_tensor = new (std::nothrow) lite::Tensor(TypeId::kNumberTypeFloat32, output_shape);
std::vector<lite::Tensor *> outputs{out_tensor};
lite::InnerContext *ctx = new (std::nothrow) lite::InnerContext();
kernel::BatchtoSpaceCudaKernel *kernel =
new (std::nothrow) kernel::BatchtoSpaceCudaKernel(CreateParameter(), inputs, outputs, ctx);
void *input_device_ptr = nullptr;
cudaMalloc(&input_device_ptr, 4 * sizeof(float));
float input_host_ptr[4]{1, 2, 3, 4};
cudaMemcpy(input_device_ptr, input_host_ptr, 4 * sizeof(float), cudaMemcpyHostToDevice);
in_tensor->set_data(input_device_ptr);
void *output_device_ptr = nullptr;
cudaMalloc(&output_device_ptr, 4 * sizeof(float));
float output_host_ptr[4];
out_tensor->set_data(output_device_ptr);
kernel->Prepare();
kernel->Run();
cudaMemcpy(output_host_ptr, output_device_ptr, 4 * sizeof(float), cudaMemcpyDeviceToHost);
for (int i = 0; i < 4; i++) {
MS_LOG(ERROR) << "BatchtoSpaceCudaKernel out: " << output_host_ptr[i];
}
cudaFree(output_device_ptr);
cudaFree(input_device_ptr);
in_tensor->set_data(nullptr);
out_tensor->set_data(nullptr);
delete in_tensor;
delete out_tensor;
delete kernel;
}
} // namespace mindspore