diff --git a/src/targets/gpu/CMakeLists.txt b/src/targets/gpu/CMakeLists.txt index 85765c6df47..09f4b7bdf39 100644 --- a/src/targets/gpu/CMakeLists.txt +++ b/src/targets/gpu/CMakeLists.txt @@ -126,7 +126,6 @@ add_library(migraphx_gpu fuse_ck.cpp fuse_mlir.cpp fuse_ops.cpp - gather.cpp gemm_impl.cpp hip.cpp kernel.cpp @@ -140,7 +139,6 @@ add_library(migraphx_gpu nonzero.cpp pack_args.cpp prefuse_ops.cpp - pad.cpp perfdb.cpp pooling.cpp reverse.cpp @@ -168,12 +166,10 @@ endfunction() register_migraphx_gpu_ops(hip_ argmax argmin - gather logsoftmax loop multinomial nonzero - pad prefix_scan_sum reverse scatter diff --git a/src/targets/gpu/device/gather.cpp b/src/targets/gpu/device/gather.cpp deleted file mode 100644 index 36f09245e99..00000000000 --- a/src/targets/gpu/device/gather.cpp +++ /dev/null @@ -1,67 +0,0 @@ -/* - * The MIT License (MIT) - * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to deal - * in the Software without restriction, including without limitation the rights - * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell - * copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in - * all copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN - * THE SOFTWARE. - */ -#include -#include -#include -#include -#include -#include - -namespace migraphx { -inline namespace MIGRAPHX_INLINE_NS { -namespace gpu { -namespace device { - -argument gather(hipStream_t stream, argument result, argument arg1, argument arg2, int64_t axis) -{ - const auto& input_shape = arg1.get_shape(); - auto lens = input_shape.lens(); - auto axis_dim_size = lens[axis]; - lens[axis] = arg2.get_shape().elements(); - shape out_comp_shape{result.get_shape().type(), lens}; - std::size_t nelements = result.get_shape().elements(); - - visit_all(result, arg1)([&](auto output, auto input_v) { - hip_visit_views(input_v, out_comp_shape)([&](auto input, auto out_comp) { - arg2.visit([&](auto indices) { - const auto* indices_ptr = device_cast(indices.data()); - auto* output_ptr = device_cast(output.data()); - gs_launch(stream, nelements, 256)([=](auto i) __device__ { - auto idx = out_comp.multi(i); - auto in_index = indices_ptr[idx[axis]]; - in_index = (in_index < 0) ? in_index + axis_dim_size : in_index; - idx[axis] = in_index; - output_ptr[i] = input[idx]; - }); - }); - }); - }); - - return result; -} - -} // namespace device -} // namespace gpu -} // namespace MIGRAPHX_INLINE_NS -} // namespace migraphx diff --git a/src/targets/gpu/device/pad.cpp b/src/targets/gpu/device/pad.cpp deleted file mode 100644 index ef2812f8651..00000000000 --- a/src/targets/gpu/device/pad.cpp +++ /dev/null @@ -1,66 +0,0 @@ -/* - * The MIT License (MIT) - * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to deal - * in the Software without restriction, including without limitation the rights - * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell - * copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in - * all copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN - * THE SOFTWARE. - */ -#include -#include -#include -#include -#include -#include -#include -#include - -namespace migraphx { -inline namespace MIGRAPHX_INLINE_NS { -namespace gpu { -namespace device { - -argument -pad(hipStream_t stream, argument result, argument arg1, float value, std::vector pads) -{ - std::size_t nelements = arg1.get_shape().elements(); - hip_visit_all(result, arg1)([&](auto output, auto input) { - using type = typename decltype(output)::value_type; - using hip_index = typename decltype(output)::hip_index; - type device_val = pad_clamp>(value); - gs_launch(stream, result.get_shape().elements())( - [=](auto i) __device__ { output.data()[i] = device_val; }); - - hip_index offsets; - std::copy(pads.begin(), pads.begin() + offsets.size(), offsets.begin()); - gs_launch(stream, nelements)([=](auto i) __device__ { - auto idx = input.get_shape().multi(i); - for(std::size_t j = 0; j < offsets.size(); j++) - { - idx[j] += offsets[j]; - } - output[idx] = input.data()[i]; - }); - }); - return result; -} - -} // namespace device -} // namespace gpu -} // namespace MIGRAPHX_INLINE_NS -} // namespace migraphx diff --git a/src/targets/gpu/gather.cpp b/src/targets/gpu/gather.cpp deleted file mode 100644 index 8160f47bab2..00000000000 --- a/src/targets/gpu/gather.cpp +++ /dev/null @@ -1,45 +0,0 @@ -/* - * The MIT License (MIT) - * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to deal - * in the Software without restriction, including without limitation the rights - * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell - * copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in - * all copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN - * THE SOFTWARE. - */ -#include -#include -#include - -namespace migraphx { -inline namespace MIGRAPHX_INLINE_NS { -namespace gpu { - -shape hip_gather::compute_shape(std::vector inputs) const -{ - inputs.pop_back(); - return op.normalize_compute_shape(inputs); -} - -argument hip_gather::compute(context& ctx, const shape&, const std::vector& args) const -{ - return device::gather(ctx.get_stream().get(), args.back(), args[0], args[1], op.axis); -} - -} // namespace gpu -} // namespace MIGRAPHX_INLINE_NS -} // namespace migraphx diff --git a/src/targets/gpu/include/migraphx/gpu/device/gather.hpp b/src/targets/gpu/include/migraphx/gpu/device/gather.hpp deleted file mode 100644 index 92912d8d2fd..00000000000 --- a/src/targets/gpu/include/migraphx/gpu/device/gather.hpp +++ /dev/null @@ -1,44 +0,0 @@ -/* - * The MIT License (MIT) - * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to deal - * in the Software without restriction, including without limitation the rights - * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell - * copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in - * all copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN - * THE SOFTWARE. - */ -#ifndef MIGRAPHX_GUARD_RTGLIB_DEVICE_GATHER_HPP -#define MIGRAPHX_GUARD_RTGLIB_DEVICE_GATHER_HPP - -#include -#include -#include - -namespace migraphx { -inline namespace MIGRAPHX_INLINE_NS { -namespace gpu { -namespace device { - -argument MIGRAPHX_DEVICE_EXPORT -gather(hipStream_t stream, argument result, argument arg1, argument arg2, int64_t axis); - -} // namespace device -} // namespace gpu -} // namespace MIGRAPHX_INLINE_NS -} // namespace migraphx - -#endif diff --git a/src/targets/gpu/include/migraphx/gpu/device/pad.hpp b/src/targets/gpu/include/migraphx/gpu/device/pad.hpp deleted file mode 100644 index 63e04daa473..00000000000 --- a/src/targets/gpu/include/migraphx/gpu/device/pad.hpp +++ /dev/null @@ -1,48 +0,0 @@ -/* - * The MIT License (MIT) - * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to deal - * in the Software without restriction, including without limitation the rights - * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell - * copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in - * all copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN - * THE SOFTWARE. - */ - -#ifndef MIGRAPHX_GUARD_RTGLIB_DEVICE_PAD_HPP -#define MIGRAPHX_GUARD_RTGLIB_DEVICE_PAD_HPP - -#include -#include -#include - -namespace migraphx { -inline namespace MIGRAPHX_INLINE_NS { -namespace gpu { -namespace device { - -argument MIGRAPHX_DEVICE_EXPORT pad(hipStream_t stream, - argument result, - argument arg1, - float value, - std::vector pads); - -} // namespace device -} // namespace gpu -} // namespace MIGRAPHX_INLINE_NS -} // namespace migraphx - -#endif diff --git a/src/targets/gpu/include/migraphx/gpu/gather.hpp b/src/targets/gpu/include/migraphx/gpu/gather.hpp deleted file mode 100644 index 4ca2019e69d..00000000000 --- a/src/targets/gpu/include/migraphx/gpu/gather.hpp +++ /dev/null @@ -1,62 +0,0 @@ -/* - * The MIT License (MIT) - * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to deal - * in the Software without restriction, including without limitation the rights - * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell - * copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in - * all copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN - * THE SOFTWARE. - */ -#ifndef MIGRAPHX_GUARD_RTGLIB_GATHER_HPP -#define MIGRAPHX_GUARD_RTGLIB_GATHER_HPP - -#include -#include -#include -#include - -namespace migraphx { -inline namespace MIGRAPHX_INLINE_NS { -namespace gpu { - -struct context; - -struct hip_gather -{ - op::gather op; - - template - static auto reflect(Self& self, F f) - { - return migraphx::reflect(self.op, f); - } - - std::string name() const { return "gpu::gather"; } - shape compute_shape(std::vector inputs) const; - argument - compute(context& ctx, const shape& output_shape, const std::vector& args) const; - std::ptrdiff_t output_alias(const std::vector& shapes) const - { - return shapes.size() - 1; - } -}; - -} // namespace gpu -} // namespace MIGRAPHX_INLINE_NS -} // namespace migraphx - -#endif diff --git a/src/targets/gpu/include/migraphx/gpu/pad.hpp b/src/targets/gpu/include/migraphx/gpu/pad.hpp deleted file mode 100644 index 9f31f132c74..00000000000 --- a/src/targets/gpu/include/migraphx/gpu/pad.hpp +++ /dev/null @@ -1,61 +0,0 @@ -/* - * The MIT License (MIT) - * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to deal - * in the Software without restriction, including without limitation the rights - * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell - * copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in - * all copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN - * THE SOFTWARE. - */ -#ifndef MIGRAPHX_GUARD_RTGLIB_PAD_HPP -#define MIGRAPHX_GUARD_RTGLIB_PAD_HPP - -#include -#include -#include - -namespace migraphx { -inline namespace MIGRAPHX_INLINE_NS { -namespace gpu { - -struct context; - -struct hip_pad -{ - op::pad op; - - template - static auto reflect(Self& self, F f) - { - return migraphx::reflect(self.op, f); - } - - std::string name() const { return "gpu::pad"; } - shape compute_shape(std::vector inputs) const; - argument - compute(context& ctx, const shape& output_shape, const std::vector& args) const; - std::ptrdiff_t output_alias(const std::vector& shapes) const - { - return shapes.size() - 1; - } -}; - -} // namespace gpu -} // namespace MIGRAPHX_INLINE_NS -} // namespace migraphx - -#endif diff --git a/src/targets/gpu/pad.cpp b/src/targets/gpu/pad.cpp deleted file mode 100644 index 5fb06af735f..00000000000 --- a/src/targets/gpu/pad.cpp +++ /dev/null @@ -1,46 +0,0 @@ -/* - * The MIT License (MIT) - * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to deal - * in the Software without restriction, including without limitation the rights - * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell - * copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in - * all copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN - * THE SOFTWARE. - */ -#include -#include -#include - -namespace migraphx { -inline namespace MIGRAPHX_INLINE_NS { -namespace gpu { - -shape hip_pad::compute_shape(std::vector inputs) const -{ - inputs.pop_back(); - check_shapes{inputs, *this}.has(1).standard(); - return op.compute_shape(inputs); -} - -argument hip_pad::compute(context& ctx, const shape&, const std::vector& args) const -{ - return device::pad(ctx.get_stream().get(), args.back(), args.front(), op.value, op.pads); -} - -} // namespace gpu -} // namespace MIGRAPHX_INLINE_NS -} // namespace migraphx