mindspore-Ecosystem
/
mindspore
21
0
Fork 27
Code Issues Pull Requests 6 Projects Releases Wiki Activity

!45895 回退 'Pull Request !44567 : [feat] [assistant] [I5EWKF] add new gpu operator Cauchy' 

Merge pull request !45895 from fengyixing/revert-merge-44567-master
This commit is contained in:
fengyixing 2022-11-22 13:40:57 +00:00 committed by Gitee
parent a044f6499a 9ce69c54da
commit 1fad6df286
No known key found for this signature in database
GPG Key ID: 173E9B9CA92EEF8F
6 changed files with 3 additions and 280 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

42
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/cauchy_impl.cu
View File

@ -1,42 +0,0 @@
/**
* 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 <curand_kernel.h>
#include "include/cuda_fp16.h"
#include "cauchy_impl.cuh"
__global__ void CauchyKernel(float *output, uint64_t seed, const float median, const float sigma, const size_t count) {
curandStatePhilox4_32_10_t state;
for (size_t i = blockIdx.x * blockDim.x + threadIdx.x; i < (count); i += blockDim.x * gridDim.x * 4) {
curand_init(seed, 0, i, &state);
auto rand = curand_uniform4(&state);
float randu[4];
randu[0] = static_cast<float>(rand.x);
randu[1] = static_cast<float>(rand.y);
randu[2] = static_cast<float>(rand.z);
randu[3] = static_cast<float>(rand.w);
#pragma unroll
for (int j = 0; j < 4; ++j, i += blockDim.x * gridDim.x) {
output[i] = median + sigma * tanf(static_cast<float>(M_PI) * (randu[j] - static_cast<float>(0.5)));
}
}
}
void Cauchy(float *output, uint64_t seed, const float median, const float sigma, const size_t count, uint32_t device_id,
cudaStream_t cuda_stream) {
CauchyKernel<<<CUDA_BLOCKS(device_id, count), CUDA_THREADS(device_id), 0, cuda_stream>>>(output, seed, median, sigma,
count);
return;
}

23
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/cauchy_impl.cuh
View File

@ -1,23 +0,0 @@
/**
* Copyright 2020 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_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_CAUCHY_IMPL_CUH_
#define MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_CAUCHY_IMPL_CUH_
#include <curand_kernel.h>
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/cuda_common.h"
CUDA_LIB_EXPORT void Cauchy(float *output, uint64_t seed, const float median, const float sigma, const size_t count,
uint32_t device_id, cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_CAUCHY_IMPL_CUH_

87
mindspore/ccsrc/plugin/device/gpu/kernel/math/cauchy_gpu_kernel.cc
View File

@ -1,87 +0,0 @@
/**
* 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 "plugin/device/gpu/kernel/math/cauchy_gpu_kernel.h"
#include <curand_kernel.h>
#include <map>
#include <vector>
#include <algorithm>
#include <utility>
#include "plugin/device/gpu/kernel/kernel_constants.h"
#include "mindspore/core/ops/cauchy.h"
#include "include/common/utils/convert_utils.h"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/cauchy_impl.cuh"
namespace mindspore {
namespace kernel {
std::vector<std::pair<KernelAttr, CauchyGpuKernelMod::CauchyFunc>> CauchyGpuKernelMod::func_list_ = {
{KernelAttr().AddOutputAttr(kNumberTypeFloat32), &CauchyGpuKernelMod::LaunchKernel<float>}};
bool CauchyGpuKernelMod::Init(const BaseOperatorPtr &base_operator, const std::vector<KernelTensorPtr> &inputs,
const std::vector<KernelTensorPtr> &outputs) {
kernel_name_ = base_operator->name();
size_ = GetValue<std::vector<int64_t>>(base_operator->GetAttr("size"));
median_ = GetValue<float>(base_operator->GetAttr("median"));
sigma_ = GetValue<float>(base_operator->GetAttr("sigma"));
auto kernel_attr = GetKernelAttrFromTensors(inputs, outputs);
auto [is_match, index] = MatchKernelAttr(kernel_attr, GetOpSupport());
if (!is_match) {
MS_LOG(ERROR) << "For 'Cauchy', it does not support this kernel type: " << kernel_attr;
return false;
}
kernel_func_ = func_list_[index].second;
return true;
}
int CauchyGpuKernelMod::Resize(const BaseOperatorPtr &base_operator, const std::vector<KernelTensorPtr> &inputs,
const std::vector<KernelTensorPtr> &outputs,
const std::map<uint32_t, tensor::TensorPtr> &) {
if (int ret = KernelMod::Resize(base_operator, inputs, outputs); ret != KRET_OK) {
return ret;
}
output_elements_ = std::accumulate(size_.begin(), size_.end(), int64_t(1), std::multiplies{});
workspace_size_list_.clear();
workspace_size_list_ = {
output_elements_ * sizeof(float),
};
return KRET_OK;
}
template <typename T>
bool CauchyGpuKernelMod::LaunchKernel(const std::vector<AddressPtr> &inputs, const std::vector<AddressPtr> &workspace,
const std::vector<AddressPtr> &outputs) {
auto output_addr = GetDeviceAddress<T>(outputs, kDim0);
auto seed = static_cast<uint64_t>(time(NULL));
Cauchy(output_addr, seed, median_, sigma_, output_elements_, device_id_,
reinterpret_cast<cudaStream_t>(cuda_stream_));
return true;
}
std::vector<KernelAttr> CauchyGpuKernelMod::GetOpSupport() {
std::vector<KernelAttr> support_list;
(void)std::transform(func_list_.begin(), func_list_.end(), std::back_inserter(support_list),
[](const std::pair<KernelAttr, CauchyGpuKernelMod::CauchyFunc> &pair) { return pair.first; });
return support_list;
}
MS_KERNEL_FACTORY_REG(NativeGpuKernelMod, Cauchy, CauchyGpuKernelMod);
} // namespace kernel
} // namespace mindspore

74
mindspore/ccsrc/plugin/device/gpu/kernel/math/cauchy_gpu_kernel.h
View File

@ -1,74 +0,0 @@
/**
* 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_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_MATH_CAUCHY_GPU_KERNEL_H_
#define MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_MATH_CAUCHY_GPU_KERNEL_H_
#include <utility>
#include <vector>
#include <map>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "include/common/utils/utils.h"
#include "plugin/device/gpu/kernel/kernel_constants.h"
namespace mindspore {
namespace kernel {
class CauchyGpuKernelMod : public NativeGpuKernelMod {
public:
CauchyGpuKernelMod() = default;
~CauchyGpuKernelMod() override = default;
bool Init(const BaseOperatorPtr &base_operator, const std::vector<KernelTensorPtr> &inputs,
const std::vector<KernelTensorPtr> &outputs) override;
int Resize(const BaseOperatorPtr &base_operator, const std::vector<KernelTensorPtr> &inputs,
const std::vector<KernelTensorPtr> &outputs, const std::map<uint32_t, tensor::TensorPtr> &) override;
bool Launch(const std::vector<AddressPtr> &inputs, const std::vector<AddressPtr> &workspace,
const std::vector<AddressPtr> &outputs, void *stream_ptr) override {
if (is_null_input_) {
return true;
}
cuda_stream_ = reinterpret_cast<cudaStream_t>(stream_ptr);
return kernel_func_(this, inputs, workspace, outputs);
}
std::vector<KernelAttr> GetOpSupport() override;
private:
using CauchyFunc =
std::function<bool(CauchyGpuKernelMod *, const std::vector<kernel::AddressPtr> &,
const std::vector<kernel::AddressPtr> &, const std::vector<kernel::AddressPtr> &)>;
template <typename T>
bool LaunchKernel(const std::vector<AddressPtr> &inputs, const std::vector<AddressPtr> &,
const std::vector<AddressPtr> &outputs);
CauchyFunc kernel_func_;
bool is_null_input_{false};
int64_t batch_num_{0};
std::vector<int64_t> size_;
float median_{0.0};
float sigma_{1.0};
int64_t output_elements_;
void *cuda_stream_{nullptr};
static std::vector<std::pair<KernelAttr, CauchyFunc>> func_list_;
};
} // namespace kernel
} // namespace mindspore
#endif // MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_MATH_CAUCHY_GPU_KERNEL_H_

6
mindspore/python/mindspore/ops/operations/math_ops.py
View File

@ -7724,14 +7724,14 @@ class Cauchy(Primitive):
ValueError: If data of `size` is not a positive integer.
Supported Platforms:
``Ascend`` ``GPU`` ``CPU``
``Ascend`` ``CPU``
Examples:
>>> size = [1]
>>> net = ops.Cauchy(size)
>>> y = net()
>>> print(y.shape)
(1,)
>>> print(y)
[0.03128606]
"""
@prim_attr_register

51
tests/st/ops/gpu/test_cauchy_op.py
View File

@ -1,51 +0,0 @@
# 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.
# ============================================================================
import pytest
from mindspore import nn
from mindspore import context
from mindspore.ops.operations.math_ops import Cauchy
class Net(nn.Cell):
"""a class used to test Cauchy gpu operator."""
def __init__(self, size, median=0.0, sigma=1.0):
super(Net, self).__init__()
self.cauchy = Cauchy(size=size, median=median, sigma=sigma)
def construct(self):
"""construct."""
return self.cauchy()
@pytest.mark.level0
@pytest.mark.platform_x86_gpu_training
@pytest.mark.env_onecard
def test_cauchy():
"""
Feature: CholeskySolve gpu TEST.
Description: test CholeskySolve operator
Expectation: the result match to numpy
"""
context.set_context(mode=context.GRAPH_MODE, device_target='GPU')
size = [1000]
net = Net(size)
result = net()
assert list(result.shape) == size
if __name__ == '__main__':
test_cauchy()