{"id":9603,"date":"2024-01-08T22:36:14","date_gmt":"2024-01-08T21:36:14","guid":{"rendered":"https:\/\/myoceane.fr\/?p=9603"},"modified":"2024-02-20T02:03:52","modified_gmt":"2024-02-20T01:03:52","slug":"gpu-implement-native-rapids-udf","status":"publish","type":"post","link":"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/","title":{"rendered":"[GPU] Native Rapids UDF – Create Custom UDF"},"content":{"rendered":"
<\/div>\n\n

\u5728\u4e0a\u4e00\u7bc7\u6211\u5011\u6210\u529f\u5be6\u4f5c\u4e86 HiveUDF<\/a>\uff0c\u70ba\u4e86\u8981\u9032\u4e00\u6b65\u5229\u7528 GPU \u52a0\u901f\uff0c\u6211\u5011\u9700\u8981\u53bb\u5be6\u4f5c evaluateColumnar \u9019\u4e00\u500b\u51fd\u6578\uff0c\u53c3\u8003 Spark Rapids \u7684 Github \u8207 ColumnView \/ ColumnVector<\/a> \u88e1\u9762\u7684\u7bc4\u4f8b\uff0c\u91dd\u5c0d\u6211\u5011\u60f3\u505a\u5230\u7684 UDF \u6211\u5011\u6c92\u6709\u767c\u73fe\u9069\u5408\u7684\u51fd\u6578\u53bb\u5be6\u4f5c\u91dd\u5c0d\u4e00\u500b Array[String] \u7684\u904e\u6ffe\u7a0b\u5f0f\uff0c\u6240\u4ee5\u6211\u5011\u9700\u8981\u81ea\u5df1\u53bb\u5be6\u4f5c Tutorial \u88e1\u9762\u6240\u8b02\u7684 Native Code Examples\uff0cTutorial \u88e1\u9762\u91dd\u5c0d HiveUDF \u53ea\u6709\u63d0\u4f9b\u4e00\u500b\u7bc4\u4f8b\u662f StringWordCount\uff0c\u672c\u7bc7\u6211\u5011\u7d00\u9304\u85c9\u7531\u9019\u500b\u7bc4\u4f8b\u53bb\u5be6\u4f5c\u4e00\u500b\u652f\u63f4 GPU \u7684 HiveUDF\u3002<\/p>\n

\u95dc\u65bc cuDF C++ \u4f60\u6240\u9700\u8981\u77e5\u9053\u7684\u4e8b\u60c5\uff0c\u7c21\u4ecb<\/a><\/h4>\n\n\n\n
\n

libcudf C++ Developer Guide <\/p>\nhttps:\/\/github.com\/rapidsai\/cudf\/blob\/branch-24.02\/cpp\/doxygen\/developer_guide\/DEVELOPER_GUIDE.md#variable-size-and-nested-data-types<\/a><\/cite><\/blockquote>\n\n\n

UDF \u67b6\u69cb<\/h4>\n

cudf \u7684 UDF \u4e3b\u8981\u5c31\u662f\u8f38\u5165\u4e00\u500b column_view \u8f38\u51fa\u4e00\u500b std::unique_ptr<cudf:column> \u4ee5\u4e0b\u7684\u7bc4\u4f8b\u90fd\u6703\u662f\u9019\u6a23\uff0c\u4e0b\u5716\u64f7\u53d6\u81ea\u7db2\u5740<\/a>\uff0c\u8aaa\u660e cudf\u00a0 \u91dd\u5c0d\u4e0d\u540c\u7684\u578b\u614b\u7684\u8cc7\u6599 column_view \u8981\u4f7f\u7528\u4e0d\u540c\u7684 type \u4f86\u5be6\u4f5c\u3002<\/p>\n\n\n

\"\"<\/figure>\n\n\n

Mastering String Transformation in RAPIDS libcudf<\/a><\/h4>\n

\u5728\u4ee5\u4e0a\u7684\u7bc4\u4f8b\u88e1\u9762\u6211\u5011\u89c0\u5bdf cudf \u662f\u600e\u9ebc\u5c0d strings_column \u505a\u8f49\u63db\u7684\uff0c\u9996\u5148\u4ed6\u6709\u5169\u500b column_view \u5206\u5225\u662f names \u8ddf visibilities\uff0c\u7136\u5f8c\u5229\u7528 cudf::string_scalar<\/a> \u5b9a\u7fa9\u5169\u500b cudf::scalar<\/a>\uff0c\u5206\u5225\u662f visible \u8ddf redaction\uff0c\u5229\u7528 cudf::strings::contains<\/a> \u6703\u56de\u50b3\u4e00\u500b cudf::column\uff0c\u6240\u4ee5 allowed \u8b8a\u6210\u4e00\u500b boolean type \u7684 column_view\uff0c\u518d\u900f\u904e cudf::copy_if_else<\/a> \u6703\u518d\u62ff\u5230\u4e00\u500b\u65b0\u7684 cudf::column\uff0c\u4ed6\u7684 type \u662f STRING\uff0c\u7d93\u904e cudf::strings::split<\/a> \u4e4b\u5f8c\u56de\u50b3\u4e00\u500b cudf::table<\/a>\uff0c\u7b2c\u4e00\u500b column \u662f first \u7b2c\u4e8c\u500b\u662f last\uff0c\u6700\u5f8c\u518d\u91cd\u7d44\u6210\u4e00\u500b table\uff0c\u5728\u7bc4\u4f8b\u4e2d\u5229\u7528 cudf::strings::concatenate<\/a> \u5c31\u53ef\u4ee5\u518d\u628a\u4e00\u500b column::table \u7d44\u56de\u53bb\u3002<\/p>\n\n\n

\/**\n * @brief Redacts each name per the corresponding visibility entry\n *\n * This implementation uses libcudf APIs to create the output result.\n *\n * @param names Column of names\n * @param visibilities Column of visibilities\n * @return Redacted column of names\n *\/\nstd::unique_ptr<cudf::column> redact_strings(cudf::column_view const& names,\n                                             cudf::column_view const& visibilities)\n{\n  auto const visible   = cudf::string_scalar(std::string(\"public\"));\n  auto const redaction = cudf::string_scalar(std::string(\"X X\"));\n\n  nvtxRangePushA(\"redact_strings\");\n\n  auto const allowed      = cudf::strings::contains(visibilities, visible);\n  auto const redacted     = cudf::copy_if_else(names, redaction, allowed->view());\n  auto const first_last   = cudf::strings::split(redacted->view());\n  auto const first        = first_last->view().column(0);\n  auto const last         = first_last->view().column(1);\n  auto const last_initial = cudf::strings::slice_strings(last, 0, 1);\n\n  auto const last_initial_first = cudf::table_view({last_initial->view(), first});\n\n  auto result = cudf::strings::concatenate(last_initial_first, std::string(\" \"));\n\n  cudaStreamSynchronize(0);\n\n  nvtxRangePop();\n  return result;\n}<\/code><\/pre>\n\n\n
\u53c3\u8003\u00a0string_word_count.cu<\/a> \/ string_word_count.hpp<\/a><\/h4>\n
spark-rapids-example \u63d0\u4f9b\u7684\u7bc4\u4f8b\u8207 Nvidia \u5b98\u7db2\u63d0\u4f9b\u7684\u7bc4\u4f8b\u4e0d\u592a\u4e00\u6a23\uff0c\u6709\u4ee5\u4e0b\u5e7e\u500b\u5dee\u5225\uff1a<\/p>\n
    \n
  1. \u4f7f\u7528 Thrust \u51fd\u5f0f\u5eab\u4f7f\u7528 GPU \u505a\u8fed\u4ee3\u5e73\u884c\u8a08\u7b97<\/li>\n
  2. \u4f7f\u7528 Rapids Memory Manager \u505a GPU memory allocation<\/li>\n<\/ol>\n\n\n
    \/*\n * Copyright (c) 2021-2022, NVIDIA CORPORATION.\n *\n * Licensed under the Apache License, Version 2.0 (the \"License\");\n * you may not use this file except in compliance with the License.\n * You may obtain a copy of the License at\n *\n *     http:\/\/www.apache.org\/licenses\/LICENSE-2.0\n *\n * Unless required by applicable law or agreed to in writing, software\n * distributed under the License is distributed on an \"AS IS\" BASIS,\n * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.\n * See the License for the specific language governing permissions and\n * limitations under the License.\n *\/\n\n#include \"string_word_count.hpp\"\n\n#include <cudf\/column\/column_factories.hpp>\n#include <cudf\/column\/column_device_view.cuh>\n#include <cudf\/strings\/string_view.cuh>\n\n#include <rmm\/cuda_stream_view.hpp>\n#include <rmm\/device_buffer.hpp>\n#include <rmm\/exec_policy.hpp>\n\n#include <thrust\/iterator\/counting_iterator.h>\n#include <thrust\/transform.h>\n\nnamespace {\n\n\/\/ count the words separated by whitespace characters\n__device__ cudf::size_type count_words(cudf::column_device_view const& d_strings,\n                                       cudf::size_type idx) {\n  if (d_strings.is_null(idx)) return 0;\n  cudf::string_view const d_str = d_strings.element<cudf::string_view>(idx);\n  cudf::size_type word_count    = 0;\n  \/\/ run of whitespace is considered a single delimiter\n  bool spaces = true;\n  auto itr    = d_str.begin();\n  while (itr != d_str.end()) {\n    cudf::char_utf8 ch = *itr;\n    if (spaces == (ch <= ' ')) {\n      itr++;\n    } else {\n      word_count += static_cast<cudf::size_type>(spaces);\n      spaces = !spaces;\n    }\n  }\n\n  return word_count;\n}\n\n\n} \/\/ anonymous namespace\n\n\/**\n * @brief Count the words in a string using whitespace as word boundaries\n *\n * @param strs The column containing the strings\n * @param stream The CUDA stream to use\n * @return The INT32 column containing the word count results per string\n *\/\nstd::unique_ptr<cudf::column> string_word_count(cudf::column_view const& strs) {\n  auto strings_count = strs.size();\n  if (strings_count == 0) {\n    return cudf::make_empty_column(cudf::data_type{cudf::type_id::INT32});\n  }\n\n  \/\/ the validity of the output matches the validity of the input\n  rmm::device_buffer null_mask = cudf::copy_bitmask(strs);\n\n  \/\/ allocate the column that will contain the word count results\n  std::unique_ptr<cudf::column> result =\n    cudf::make_numeric_column(\n      cudf::data_type{cudf::type_id::INT32},\n      strs.size(),\n      std::move(null_mask),\n      strs.null_count());\n\n  \/\/ compute the word counts, writing into the result column data buffer\n  auto stream = rmm::cuda_stream_default;\n  auto strs_device_view = cudf::column_device_view::create(strs, stream);\n  auto d_strs_view = *strs_device_view;\n  thrust::transform(\n    rmm::exec_policy(stream),\n    thrust::make_counting_iterator<cudf::size_type>(0),\n    thrust::make_counting_iterator<cudf::size_type>(strings_count),\n    result->mutable_view().data<cudf::size_type>(),\n    [d_strs_view] __device__(cudf::size_type idx) { return count_words(d_strs_view, idx); });\n\n  return result;\n}<\/code><\/pre>\n\n\n
    \u53c3\u8003 cosine_similarity.cu<\/a> \/ cosine_similarity.hpp<\/a><\/h4>\n
    \u5728\u8655\u7406 inner product \u7684\u6642\u5019\uff0c\u9700\u8981\u7528\u5230 Thrust<\/a> \u9019\u4e00\u500b\u7b97\u6cd5\u5eab\uff0c\u6839\u64da Nvidia \u5b98\u7db2\u7684\u4ecb\u7d39\uff0c<\/p>\n\n\n
    \n
    Thrust is a C++ template library for CUDA based on the Standard Template Library (STL). Thrust allows you to implement high performance parallel applications with minimal programming effort through a high-level interface that is fully interoperable with CUDA C.<\/p>\nhttps:\/\/docs.nvidia.com\/cuda\/thrust\/index.html<\/cite><\/blockquote>\n\n\n
    \u5728\u8a08\u7b97 cosine_similarity \u7684\u6642\u5019\uff0c\u6709\u7528\u5230\u4ee5\u4e0b\u9019\u5169\u500b thrust::all_of<\/a> \u8207 thrust::transform<\/a>\uff0c\u5f8c\u8005\u662f\u4ee5\u4e0a\u5169\u500b\u7bc4\u4f8b\u90fd\u6709\u7528\u5230\u7684\u51fd\u5f0f\uff0c<\/p>\n\n\n
    \/*\n * Copyright (c) 2021-2022, NVIDIA CORPORATION.\n *\n * Licensed under the Apache License, Version 2.0 (the \"License\");\n * you may not use this file except in compliance with the License.\n * You may obtain a copy of the License at\n *\n *     http:\/\/www.apache.org\/licenses\/LICENSE-2.0\n *\n * Unless required by applicable law or agreed to in writing, software\n * distributed under the License is distributed on an \"AS IS\" BASIS,\n * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.\n * See the License for the specific language governing permissions and\n * limitations under the License.\n *\/\n\n#include \"cosine_similarity.hpp\"\n\n#include <cudf\/column\/column_factories.hpp>\n#include <cudf\/column\/column_device_view.cuh>\n#include <cudf\/lists\/list_device_view.cuh>\n#include <cudf\/lists\/lists_column_device_view.cuh>\n#include <cudf\/lists\/lists_column_view.hpp>\n#include <cudf\/null_mask.hpp>\n#include <cudf\/table\/table_view.hpp>\n\n#include <rmm\/cuda_stream_view.hpp>\n#include <rmm\/device_buffer.hpp>\n#include <rmm\/exec_policy.hpp>\n\n#include <thrust\/iterator\/counting_iterator.h>\n#include <thrust\/logical.h>\n#include <thrust\/transform.h>\n\n#include <cmath>\n\nnamespace {\n\n\/**\n * @brief Functor for computing the cosine similarity between two list of float columns\n *\/\nstruct cosine_similarity_functor {\n  float const* const v1;\n  float const* const v2;\n  int32_t const* const v1_offsets;\n  int32_t const* const v2_offsets;\n\n  \/\/ This kernel executes thread-per-row which should be fine for relatively short lists\n  \/\/ but may need to be revisited for performance if operating on long lists.\n  __device__ float operator()(cudf::size_type row_idx) {\n    auto const v1_start_idx = v1_offsets[row_idx];\n    auto const v1_num_elems = v1_offsets[row_idx + 1] - v1_start_idx;\n    auto const v2_start_idx = v2_offsets[row_idx];\n    auto const v2_num_elems = v2_offsets[row_idx + 1] - v2_start_idx;\n    auto const num_elems = std::min(v1_num_elems, v2_num_elems);\n    double mag1 = 0;\n    double mag2 = 0;\n    double dot_product = 0;\n    for (auto i = 0; i < num_elems; i++) {\n      float const f1 = v1[v1_start_idx + i];\n      mag1 += f1 * f1;\n      float const f2 = v2[v2_start_idx + i];\n      mag2 += f2 * f2;\n      dot_product += f1 * f2;\n    }\n    mag1 = std::sqrt(mag1);\n    mag2 = std::sqrt(mag2);\n    return static_cast<float>(dot_product \/ (mag1 * mag2));\n  }\n};\n\n} \/\/ anonymous namespace\n\n\/**\n * @brief Compute the cosine similarity between two LIST of FLOAT32 columns\n *\n * The input vectors must have matching shapes, i.e.: same row count and same number of\n * list elements per row. A null list row is supported, but null float entries within a\n * list are not supported.\n *\n * @param lv1 The first LIST of FLOAT32 column view\n * @param lv2 The second LIST of FLOAT32 column view\n * @return A FLOAT32 column containing the cosine similarity corresponding to each input row\n *\/\nstd::unique_ptr<cudf::column> cosine_similarity(cudf::lists_column_view const& lv1,\n                                                cudf::lists_column_view const& lv2) {\n  \/\/ sanity-check the input types\n  if (lv1.child().type().id() != lv2.child().type().id() ||\n      lv1.child().type().id() != cudf::type_id::FLOAT32) {\n    throw std::invalid_argument(\"inputs are not lists of floats\");\n  }\n\n  \/\/ sanity check the input shape\n  auto const row_count = lv1.size();\n  if (row_count != lv2.size()) {\n    throw std::invalid_argument(\"input row counts do not match\");\n  }\n  if (row_count == 0) {\n    return cudf::make_empty_column(cudf::data_type{cudf::type_id::FLOAT32});\n  }\n  if (lv1.child().null_count() != 0 || lv2.child().null_count() != 0) {\n    throw std::invalid_argument(\"null floats are not supported\");\n  }\n\n  auto const stream = rmm::cuda_stream_default;\n  auto d_view1_ptr = cudf::column_device_view::create(lv1.parent());\n  auto d_lists1 = cudf::detail::lists_column_device_view(*d_view1_ptr);\n  auto d_view2_ptr = cudf::column_device_view::create(lv2.parent());\n  auto d_lists2 = cudf::detail::lists_column_device_view(*d_view2_ptr);\n  bool const are_offsets_equal =\n    thrust::all_of(rmm::exec_policy(stream),\n                   thrust::make_counting_iterator<cudf::size_type>(0),\n                   thrust::make_counting_iterator<cudf::size_type>(row_count),\n                   [d_lists1, d_lists2] __device__(cudf::size_type idx) {\n                     auto ldv1 = cudf::list_device_view(d_lists1, idx);\n                     auto ldv2 = cudf::list_device_view(d_lists2, idx);\n                     return ldv1.is_null() || ldv2.is_null() || ldv1.size() == ldv2.size();\n                   });\n  if (not are_offsets_equal) {\n    throw std::invalid_argument(\"input list lengths do not match for every row\");\n  }\n\n  \/\/ allocate the vector of float results\n  rmm::device_uvector<float> float_results(row_count, stream);\n\n  \/\/ compute the cosine similarity\n  auto const lv1_data = lv1.child().data<float>();\n  auto const lv2_data = lv2.child().data<float>();\n  auto const lv1_offsets = lv1.offsets().data<int32_t>();\n  auto const lv2_offsets = lv2.offsets().data<int32_t>();\n  thrust::transform(rmm::exec_policy(stream),\n                    thrust::make_counting_iterator<cudf::size_type>(0),\n                    thrust::make_counting_iterator<cudf::size_type>(row_count),\n                    float_results.data(),\n                    cosine_similarity_functor({lv1_data, lv2_data, lv1_offsets, lv2_offsets}));\n\n  \/\/ the validity of the output is the bitwise-and of the two input validity masks\n  auto [null_mask, null_count] = cudf::bitmask_and(cudf::table_view({lv1.parent(), lv2.parent()}));\n\n  return std::make_unique<cudf::column>(cudf::data_type{cudf::type_id::FLOAT32},\n                                        row_count,\n                                        float_results.release(),\n                                        std::move(null_mask),\n                                        null_count);\n}<\/code><\/pre>\n\n\n
    \u5be6\u4f5c Filtering List of String \u7684\u51fd\u5f0f<\/h4>\n
    \u4e0a\u65b9\u6709\u8aaa\u5230\u8981\u5be6\u4f5c\u91dd\u5c0d Array[String] \u7684\u904e\u6ffe\u6703\u9700\u8981\u540c\u6642\u5be6\u4f5c cudf \u88e1\u9762\u7684 cudf::lists<\/a> \u8ddf cudf::strings<\/a>\uff0c\u9996\u5148\u6211\u5011\u8981\u5148\u4e86\u89e3 List \u5728 Apache Arrow \u88e1\u9762\u662f\u600e\u9ebc\u505a\u5b58\u5132\u7684\uff1f\u5728 libcudf c++ Developer Guide \u88e1\u9762\u4ed6\u63d0\u4f9b\u4e00\u500b\u8907\u96dc\u7684\u7bc4\u4f8b\uff1aList<List<List<List<int>>>>\uff0c\u771f\u6b63\u7684\u6574\u6578\u503c\u6703\u5168\u90e8\u6524\u5e73\u653e\u5728\u4e00\u500b List<int> \u4e26\u4e14\u4f9d\u9760\u591a\u500b Offsets \u8207 Null array \u4f86\u9054\u5230\u5340\u5206\u4e0d\u540c\u5c64 List \u7684\u6548\u679c\u3002<\/p>\n\n\n
    lists_column = { {{{1, 2}, {3, 4}}, NULL}, {{{10, 20}, {30, 40}}, {{50, 60, 70}, {0}}} }\n\n   List<List<List<int>>>  (2 rows):\n   Length : 2\n   Offsets : 0, 2, 4\n   Children :\n      List<List<int>>:\n      Length : 4\n      Offsets : 0, 2, 2, 4, 6\n      Null count: 1\n        1101\n      Children :\n        List<int>:\n        Length : 6\n        Offsets : 0, 2, 4, 6, 8, 11, 12\n        Children :\n          Column of ints\n          1, 2, 3, 4, 10, 20, 30, 40, 50, 60, 70, 0<\/code><\/pre>\n\n\n
    \u90a3\u9ebc\u5728 Apache Arrow \u88e1\u9762\u662f\u600e\u9ebc\u8868\u793a List<String> \u7684\u5462\uff1f\u6211\u5011\u5229\u7528\u4ee5\u4e0b\u7684\u4f8b\u5b50\u505a\u8aaa\u660e\uff1a<\/p>\n\n\n
    lists_string = [[\"This\", \"is\", \"a\"], [\"column\", \"of\", \"strings\"]]\n\n   List<String> (2 rows):\n   Length : 2\n   Children :\n      Offsets : 0, 3, 6\n      String:\n      Length : 6\n      Children : \n      \t Offsets: 0, 4, 6, 7, 13, 15, 22\n         Characters: Thisisacolumnofstrings<\/code><\/pre>\n\n\n
    \u6240\u4ee5\u5728\u5be6\u4f5c\u91dd\u5c0d List<String> \u958b\u982d\u662f test \u7684 filtering \u53ef\u4ee5\u4f7f\u7528\u4ee5\u4e0b\u7684\u7a0b\u5f0f\u78bc\uff1a<\/p>\n\n\n
    std::unique_ptr<cudf::column> filtering_startswith_test(cudf::lists_column_view const& lv) {\n  \/\/ sanity-check the input types\n  if (lv.child().type().id() != cudf::type_id::STRING) {\n    throw std::invalid_argument(\"inputs are not lists of string\");\n  }\n\n  \/\/ sanity check the input shape\n  auto const row_count = lv.size();\n  if (row_count == 0) {\n    return cudf::make_empty_column(cudf::data_type{cudf::type_id::LIST});\n  }\n  if (lv.child().null_count() != 0) {\n    throw std::invalid_argument(\"null string are not supported\");\n  }\n\n  auto const scv = cudf::strings_column_view(lv.child());\n  auto const rs_start = cudf::string_scalar(std::string(\"test\"));\n  \/\/ build lists of bools column\n  auto const list_bools = cudf::make_lists_column(lv.size(),\n                                                  std::make_unique<cudf::column>(lv.offsets()),\n                                                  cudf::strings::starts_with(scv, rs_start),\n                                                  lv.null_count(),\n                                                  cudf::copy_bitmask(lv.parent()));\n  return cudf::lists::apply_boolean_mask(lv, list_bools->view());\n}\n                                                  <\/code><\/pre>\n\n\n
    \u53e6\u5916\u95dc\u65bc Struct \u5728 Apache Arrow \u4e2d\u7684 Struct \u7684\u5132\u5b58\u65b9\u6cd5\uff0c\u4ee5\u4e0b\u662f libcudf C++ Developer Guide \u63d0\u4f9b\u7684\u7bc4\u4f8b\uff0c<\/p>\n\n\n
    {\n  type = STRUCT\n  null_mask = [1, 1, 0, 1]\n  null_count = 1\n  children = {\n    {\n      type = FLOAT32\n      data =       [1.0, 4.0, X, 8.0]\n      null_mask  = [  1,   1, 0,   1]\n      null_count = 1\n    },\n    {\n      type = INT32\n      data =       [2, 5, X, X]\n      null_mask  = [1, 1, 0, 0]\n      null_count = 2\n    }\n  }\n}<\/code><\/pre>\n\n\n\n
    \n
    <\/p>\n<\/blockquote>\n","protected":false},"excerpt":{"rendered":"
    \u5728\u4e0a\u4e00\u7bc7\u6211\u5011\u6210\u529f\u5be6\u4f5c\u4e86 HiveUDF\uff0c\u70ba\u4e86\u8981\u9032\u4e00\u6b65\u5229\u7528 GPU \u52a0\u901f\uff0c\u6211\u5011\u9700\u8981\u53bb\u5be6\u4f5c evaluateColumnar \u9019\u4e00\u500b\u51fd\u6578\uff0c\u53c3\u8003 Spark Rapids \u7684 Github \u8207 ColumnView \u88e1\u9762\u7684\u7bc4\u4f8b\uff0c\u91dd\u5c0d\u6211\u5011\u60f3\u505a\u5230\u7684 UDF \u6211\u5011\u6c92\u6709\u767c\u73fe\u9069\u5408\u7684\u51fd\u6578\u53bb\u5be6\u4f5c\u91dd\u5c0d\u4e00\u500b Array[String] \u7684\u904e\u6ffe\u7a0b\u5f0f\uff0c\u6240\u4ee5\u6211\u5011\u9700\u8981\u81ea\u5df1\u53bb\u5be6\u4f5c Tutorial \u88e1\u9762\u6240\u8b02\u7684 Native Code Examples\uff0cTutorial \u88e1\u9762\u91dd\u5c0d HiveUDF \u53ea\u6709\u63d0\u4f9b\u4e00\u500b\u7bc4\u4f8b\u662f StringWordCount\uff0c\u672c\u7bc7\u6211\u5011\u7d00\u9304\u85c9\u7531\u9019\u500b\u7bc4\u4f8b\u53bb\u5be6\u4f5c\u4e00\u500b\u652f\u63f4 GPU \u7684 HiveUDF\u3002<\/p>\n","protected":false},"author":1,"featured_media":8700,"comment_status":"open","ping_status":"open","sticky":false,"template":"","format":"standard","meta":{"footnotes":""},"categories":[9,1761,14],"tags":[],"class_list":["post-9603","post","type-post","status-publish","format-standard","has-post-thumbnail","hentry","category-bigdata-ml","category-gpu","category-it-technology"],"yoast_head":"\n[GPU] Native Rapids UDF - Create Custom UDF - \u60f3\u65b9\u6d89\u6cd5 - \u91cf\u74f6\u5916\u7684\u5929\u7a7a M-Y-Oceane<\/title>\n<meta name=\"robots\" content=\"index, follow, max-snippet:-1, max-image-preview:large, max-video-preview:-1\" \/>\n<link rel=\"canonical\" href=\"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/\" \/>\n<meta property=\"og:locale\" content=\"en_US\" \/>\n<meta property=\"og:type\" content=\"article\" \/>\n<meta property=\"og:title\" content=\"[GPU] Native Rapids UDF - Create Custom UDF - \u60f3\u65b9\u6d89\u6cd5 - \u91cf\u74f6\u5916\u7684\u5929\u7a7a M-Y-Oceane\" \/>\n<meta property=\"og:description\" content=\"\u5728\u4e0a\u4e00\u7bc7\u6211\u5011\u6210\u529f\u5be6\u4f5c\u4e86 HiveUDF\uff0c\u70ba\u4e86\u8981\u9032\u4e00\u6b65\u5229\u7528 GPU \u52a0\u901f\uff0c\u6211\u5011\u9700\u8981\u53bb\u5be6\u4f5c evaluateColumnar \u9019\u4e00\u500b\u51fd\u6578\uff0c\u53c3\u8003 Spark Rapids \u7684 Github \u8207 ColumnView \u88e1\u9762\u7684\u7bc4\u4f8b\uff0c\u91dd\u5c0d\u6211\u5011\u60f3\u505a\u5230\u7684 UDF \u6211\u5011\u6c92\u6709\u767c\u73fe\u9069\u5408\u7684\u51fd\u6578\u53bb\u5be6\u4f5c\u91dd\u5c0d\u4e00\u500b Array[String] \u7684\u904e\u6ffe\u7a0b\u5f0f\uff0c\u6240\u4ee5\u6211\u5011\u9700\u8981\u81ea\u5df1\u53bb\u5be6\u4f5c Tutorial \u88e1\u9762\u6240\u8b02\u7684 Native Code Examples\uff0cTutorial \u88e1\u9762\u91dd\u5c0d HiveUDF \u53ea\u6709\u63d0\u4f9b\u4e00\u500b\u7bc4\u4f8b\u662f StringWordCount\uff0c\u672c\u7bc7\u6211\u5011\u7d00\u9304\u85c9\u7531\u9019\u500b\u7bc4\u4f8b\u53bb\u5be6\u4f5c\u4e00\u500b\u652f\u63f4 GPU \u7684 HiveUDF\u3002\" \/>\n<meta property=\"og:url\" content=\"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/\" \/>\n<meta property=\"og:site_name\" content=\"\u60f3\u65b9\u6d89\u6cd5 - \u91cf\u74f6\u5916\u7684\u5929\u7a7a M-Y-Oceane\" \/>\n<meta property=\"article:published_time\" content=\"2024-01-08T21:36:14+00:00\" \/>\n<meta property=\"article:modified_time\" content=\"2024-02-20T01:03:52+00:00\" \/>\n<meta property=\"og:image\" content=\"https:\/\/myoceane.fr\/wp-content\/uploads\/2022\/08\/RAPIDSSpark.png\" \/>\n\t<meta property=\"og:image:width\" content=\"936\" \/>\n\t<meta property=\"og:image:height\" content=\"248\" \/>\n\t<meta property=\"og:image:type\" content=\"image\/png\" \/>\n<meta name=\"author\" content=\"\u6ab8\u6aac\u7238\" \/>\n<meta name=\"twitter:card\" content=\"summary_large_image\" \/>\n<meta name=\"twitter:label1\" content=\"Written by\" \/>\n\t<meta name=\"twitter:data1\" content=\"\u6ab8\u6aac\u7238\" \/>\n\t<meta name=\"twitter:label2\" content=\"Est. reading time\" \/>\n\t<meta name=\"twitter:data2\" content=\"2 minutes\" \/>\n<script type=\"application\/ld+json\" class=\"yoast-schema-graph\">{\"@context\":\"https:\/\/schema.org\",\"@graph\":[{\"@type\":\"Article\",\"@id\":\"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/#article\",\"isPartOf\":{\"@id\":\"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/\"},\"author\":{\"name\":\"\u6ab8\u6aac\u7238\",\"@id\":\"https:\/\/myoceane.fr\/#\/schema\/person\/4a4552fb8c27693083d465e12db7658b\"},\"headline\":\"[GPU] Native Rapids UDF – Create Custom UDF\",\"datePublished\":\"2024-01-08T21:36:14+00:00\",\"dateModified\":\"2024-02-20T01:03:52+00:00\",\"mainEntityOfPage\":{\"@id\":\"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/\"},\"wordCount\":259,\"commentCount\":0,\"publisher\":{\"@id\":\"https:\/\/myoceane.fr\/#\/schema\/person\/4a4552fb8c27693083d465e12db7658b\"},\"image\":{\"@id\":\"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/#primaryimage\"},\"thumbnailUrl\":\"https:\/\/myoceane.fr\/wp-content\/uploads\/2022\/08\/RAPIDSSpark.png\",\"articleSection\":[\"Big Data & Machine Learning\",\"GPU\",\"IT Technology\"],\"inLanguage\":\"en-US\",\"potentialAction\":[{\"@type\":\"CommentAction\",\"name\":\"Comment\",\"target\":[\"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/#respond\"]}]},{\"@type\":\"WebPage\",\"@id\":\"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/\",\"url\":\"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/\",\"name\":\"[GPU] Native Rapids UDF - Create Custom UDF - \u60f3\u65b9\u6d89\u6cd5 - \u91cf\u74f6\u5916\u7684\u5929\u7a7a M-Y-Oceane\",\"isPartOf\":{\"@id\":\"https:\/\/myoceane.fr\/#website\"},\"primaryImageOfPage\":{\"@id\":\"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/#primaryimage\"},\"image\":{\"@id\":\"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/#primaryimage\"},\"thumbnailUrl\":\"https:\/\/myoceane.fr\/wp-content\/uploads\/2022\/08\/RAPIDSSpark.png\",\"datePublished\":\"2024-01-08T21:36:14+00:00\",\"dateModified\":\"2024-02-20T01:03:52+00:00\",\"breadcrumb\":{\"@id\":\"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/#breadcrumb\"},\"inLanguage\":\"en-US\",\"potentialAction\":[{\"@type\":\"ReadAction\",\"target\":[\"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/\"]}]},{\"@type\":\"ImageObject\",\"inLanguage\":\"en-US\",\"@id\":\"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/#primaryimage\",\"url\":\"https:\/\/myoceane.fr\/wp-content\/uploads\/2022\/08\/RAPIDSSpark.png\",\"contentUrl\":\"https:\/\/myoceane.fr\/wp-content\/uploads\/2022\/08\/RAPIDSSpark.png\",\"width\":936,\"height\":248},{\"@type\":\"BreadcrumbList\",\"@id\":\"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/#breadcrumb\",\"itemListElement\":[{\"@type\":\"ListItem\",\"position\":1,\"name\":\"Home\",\"item\":\"https:\/\/myoceane.fr\/\"},{\"@type\":\"ListItem\",\"position\":2,\"name\":\"[GPU] Native Rapids UDF – Create Custom UDF\"}]},{\"@type\":\"WebSite\",\"@id\":\"https:\/\/myoceane.fr\/#website\",\"url\":\"https:\/\/myoceane.fr\/\",\"name\":\"M-Y-Oceane \u60f3\u65b9\u6d89\u6cd5\u3002\u91cf\u74f6\u5916\u7684\u5929\u7a7a\",\"description\":\"\u60f3\u65b9\u6d89\u6cd5, France, Taiwan, Health, Information Technology\",\"publisher\":{\"@id\":\"https:\/\/myoceane.fr\/#\/schema\/person\/4a4552fb8c27693083d465e12db7658b\"},\"potentialAction\":[{\"@type\":\"SearchAction\",\"target\":{\"@type\":\"EntryPoint\",\"urlTemplate\":\"https:\/\/myoceane.fr\/?s={search_term_string}\"},\"query-input\":{\"@type\":\"PropertyValueSpecification\",\"valueRequired\":true,\"valueName\":\"search_term_string\"}}],\"inLanguage\":\"en-US\"},{\"@type\":[\"Person\",\"Organization\"],\"@id\":\"https:\/\/myoceane.fr\/#\/schema\/person\/4a4552fb8c27693083d465e12db7658b\",\"name\":\"\u6ab8\u6aac\u7238\",\"image\":{\"@type\":\"ImageObject\",\"inLanguage\":\"en-US\",\"@id\":\"https:\/\/myoceane.fr\/#\/schema\/person\/image\/\",\"url\":\"https:\/\/secure.gravatar.com\/avatar\/6cc678684664f8ad45a8d56a6630b183?s=96&d=mm&r=g\",\"contentUrl\":\"https:\/\/secure.gravatar.com\/avatar\/6cc678684664f8ad45a8d56a6630b183?s=96&d=mm&r=g\",\"caption\":\"\u6ab8\u6aac\u7238\"},\"logo\":{\"@id\":\"https:\/\/myoceane.fr\/#\/schema\/person\/image\/\"},\"url\":\"https:\/\/myoceane.fr\/index.php\/author\/johnny5584767gmail-com\/\"}]}<\/script>\n<!-- \/ Yoast SEO plugin. -->","yoast_head_json":{"title":"[GPU] Native Rapids UDF - Create Custom UDF - \u60f3\u65b9\u6d89\u6cd5 - \u91cf\u74f6\u5916\u7684\u5929\u7a7a M-Y-Oceane","robots":{"index":"index","follow":"follow","max-snippet":"max-snippet:-1","max-image-preview":"max-image-preview:large","max-video-preview":"max-video-preview:-1"},"canonical":"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/","og_locale":"en_US","og_type":"article","og_title":"[GPU] Native Rapids UDF - Create Custom UDF - \u60f3\u65b9\u6d89\u6cd5 - \u91cf\u74f6\u5916\u7684\u5929\u7a7a M-Y-Oceane","og_description":"\u5728\u4e0a\u4e00\u7bc7\u6211\u5011\u6210\u529f\u5be6\u4f5c\u4e86 HiveUDF\uff0c\u70ba\u4e86\u8981\u9032\u4e00\u6b65\u5229\u7528 GPU \u52a0\u901f\uff0c\u6211\u5011\u9700\u8981\u53bb\u5be6\u4f5c evaluateColumnar \u9019\u4e00\u500b\u51fd\u6578\uff0c\u53c3\u8003 Spark Rapids \u7684 Github \u8207 ColumnView \u88e1\u9762\u7684\u7bc4\u4f8b\uff0c\u91dd\u5c0d\u6211\u5011\u60f3\u505a\u5230\u7684 UDF \u6211\u5011\u6c92\u6709\u767c\u73fe\u9069\u5408\u7684\u51fd\u6578\u53bb\u5be6\u4f5c\u91dd\u5c0d\u4e00\u500b Array[String] \u7684\u904e\u6ffe\u7a0b\u5f0f\uff0c\u6240\u4ee5\u6211\u5011\u9700\u8981\u81ea\u5df1\u53bb\u5be6\u4f5c Tutorial \u88e1\u9762\u6240\u8b02\u7684 Native Code Examples\uff0cTutorial \u88e1\u9762\u91dd\u5c0d HiveUDF \u53ea\u6709\u63d0\u4f9b\u4e00\u500b\u7bc4\u4f8b\u662f StringWordCount\uff0c\u672c\u7bc7\u6211\u5011\u7d00\u9304\u85c9\u7531\u9019\u500b\u7bc4\u4f8b\u53bb\u5be6\u4f5c\u4e00\u500b\u652f\u63f4 GPU \u7684 HiveUDF\u3002","og_url":"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/","og_site_name":"\u60f3\u65b9\u6d89\u6cd5 - \u91cf\u74f6\u5916\u7684\u5929\u7a7a M-Y-Oceane","article_published_time":"2024-01-08T21:36:14+00:00","article_modified_time":"2024-02-20T01:03:52+00:00","og_image":[{"width":936,"height":248,"url":"https:\/\/myoceane.fr\/wp-content\/uploads\/2022\/08\/RAPIDSSpark.png","type":"image\/png"}],"author":"\u6ab8\u6aac\u7238","twitter_card":"summary_large_image","twitter_misc":{"Written by":"\u6ab8\u6aac\u7238","Est. reading time":"2 minutes"},"schema":{"@context":"https:\/\/schema.org","@graph":[{"@type":"Article","@id":"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/#article","isPartOf":{"@id":"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/"},"author":{"name":"\u6ab8\u6aac\u7238","@id":"https:\/\/myoceane.fr\/#\/schema\/person\/4a4552fb8c27693083d465e12db7658b"},"headline":"[GPU] Native Rapids UDF – Create Custom UDF","datePublished":"2024-01-08T21:36:14+00:00","dateModified":"2024-02-20T01:03:52+00:00","mainEntityOfPage":{"@id":"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/"},"wordCount":259,"commentCount":0,"publisher":{"@id":"https:\/\/myoceane.fr\/#\/schema\/person\/4a4552fb8c27693083d465e12db7658b"},"image":{"@id":"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/#primaryimage"},"thumbnailUrl":"https:\/\/myoceane.fr\/wp-content\/uploads\/2022\/08\/RAPIDSSpark.png","articleSection":["Big Data & Machine Learning","GPU","IT Technology"],"inLanguage":"en-US","potentialAction":[{"@type":"CommentAction","name":"Comment","target":["https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/#respond"]}]},{"@type":"WebPage","@id":"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/","url":"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/","name":"[GPU] Native Rapids UDF - Create Custom UDF - \u60f3\u65b9\u6d89\u6cd5 - \u91cf\u74f6\u5916\u7684\u5929\u7a7a M-Y-Oceane","isPartOf":{"@id":"https:\/\/myoceane.fr\/#website"},"primaryImageOfPage":{"@id":"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/#primaryimage"},"image":{"@id":"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/#primaryimage"},"thumbnailUrl":"https:\/\/myoceane.fr\/wp-content\/uploads\/2022\/08\/RAPIDSSpark.png","datePublished":"2024-01-08T21:36:14+00:00","dateModified":"2024-02-20T01:03:52+00:00","breadcrumb":{"@id":"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/#breadcrumb"},"inLanguage":"en-US","potentialAction":[{"@type":"ReadAction","target":["https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/"]}]},{"@type":"ImageObject","inLanguage":"en-US","@id":"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/#primaryimage","url":"https:\/\/myoceane.fr\/wp-content\/uploads\/2022\/08\/RAPIDSSpark.png","contentUrl":"https:\/\/myoceane.fr\/wp-content\/uploads\/2022\/08\/RAPIDSSpark.png","width":936,"height":248},{"@type":"BreadcrumbList","@id":"https:\/\/myoceane.fr\/index.php\/gpu-implement-native-rapids-udf\/#breadcrumb","itemListElement":[{"@type":"ListItem","position":1,"name":"Home","item":"https:\/\/myoceane.fr\/"},{"@type":"ListItem","position":2,"name":"[GPU] Native Rapids UDF – Create Custom UDF"}]},{"@type":"WebSite","@id":"https:\/\/myoceane.fr\/#website","url":"https:\/\/myoceane.fr\/","name":"M-Y-Oceane \u60f3\u65b9\u6d89\u6cd5\u3002\u91cf\u74f6\u5916\u7684\u5929\u7a7a","description":"\u60f3\u65b9\u6d89\u6cd5, France, Taiwan, Health, Information Technology","publisher":{"@id":"https:\/\/myoceane.fr\/#\/schema\/person\/4a4552fb8c27693083d465e12db7658b"},"potentialAction":[{"@type":"SearchAction","target":{"@type":"EntryPoint","urlTemplate":"https:\/\/myoceane.fr\/?s={search_term_string}"},"query-input":{"@type":"PropertyValueSpecification","valueRequired":true,"valueName":"search_term_string"}}],"inLanguage":"en-US"},{"@type":["Person","Organization"],"@id":"https:\/\/myoceane.fr\/#\/schema\/person\/4a4552fb8c27693083d465e12db7658b","name":"\u6ab8\u6aac\u7238","image":{"@type":"ImageObject","inLanguage":"en-US","@id":"https:\/\/myoceane.fr\/#\/schema\/person\/image\/","url":"https:\/\/secure.gravatar.com\/avatar\/6cc678684664f8ad45a8d56a6630b183?s=96&d=mm&r=g","contentUrl":"https:\/\/secure.gravatar.com\/avatar\/6cc678684664f8ad45a8d56a6630b183?s=96&d=mm&r=g","caption":"\u6ab8\u6aac\u7238"},"logo":{"@id":"https:\/\/myoceane.fr\/#\/schema\/person\/image\/"},"url":"https:\/\/myoceane.fr\/index.php\/author\/johnny5584767gmail-com\/"}]}},"amp_enabled":false,"_links":{"self":[{"href":"https:\/\/myoceane.fr\/index.php\/wp-json\/wp\/v2\/posts\/9603","targetHints":{"allow":["GET"]}}],"collection":[{"href":"https:\/\/myoceane.fr\/index.php\/wp-json\/wp\/v2\/posts"}],"about":[{"href":"https:\/\/myoceane.fr\/index.php\/wp-json\/wp\/v2\/types\/post"}],"author":[{"embeddable":true,"href":"https:\/\/myoceane.fr\/index.php\/wp-json\/wp\/v2\/users\/1"}],"replies":[{"embeddable":true,"href":"https:\/\/myoceane.fr\/index.php\/wp-json\/wp\/v2\/comments?post=9603"}],"version-history":[{"count":47,"href":"https:\/\/myoceane.fr\/index.php\/wp-json\/wp\/v2\/posts\/9603\/revisions"}],"predecessor-version":[{"id":9683,"href":"https:\/\/myoceane.fr\/index.php\/wp-json\/wp\/v2\/posts\/9603\/revisions\/9683"}],"wp:featuredmedia":[{"embeddable":true,"href":"https:\/\/myoceane.fr\/index.php\/wp-json\/wp\/v2\/media\/8700"}],"wp:attachment":[{"href":"https:\/\/myoceane.fr\/index.php\/wp-json\/wp\/v2\/media?parent=9603"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"https:\/\/myoceane.fr\/index.php\/wp-json\/wp\/v2\/categories?post=9603"},{"taxonomy":"post_tag","embeddable":true,"href":"https:\/\/myoceane.fr\/index.php\/wp-json\/wp\/v2\/tags?post=9603"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}