You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by GitBox <gi...@apache.org> on 2021/02/15 08:20:03 UTC
[GitHub] [tvm] masahi opened a new pull request #7458: [ROCM] Add Thrust support
masahi opened a new pull request #7458:
URL: https://github.com/apache/tvm/pull/7458
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] tkonolige commented on a change in pull request #7458: [ROCM] Add Thrust support
Posted by GitBox <gi...@apache.org>.
tkonolige commented on a change in pull request #7458:
URL: https://github.com/apache/tvm/pull/7458#discussion_r576376750
##########
File path: src/runtime/contrib/rocthrust/thrust.cc
##########
@@ -0,0 +1,378 @@
+/*
+ * 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.
+ */
+
+/*!
+ * \file Use external Thrust library call
+ */
+
+#include <dlpack/dlpack.h>
+#include <thrust/device_ptr.h>
+#include <thrust/device_vector.h>
+#include <thrust/gather.h>
+#include <thrust/scan.h>
+#include <thrust/sequence.h>
+#include <thrust/sort.h>
+#include <tvm/runtime/registry.h>
+
+#include <algorithm>
+#include <functional>
+#include <vector>
+
+namespace tvm {
+namespace contrib {
+
+using namespace runtime;
+
+// Performs sorting along axis -1 and returns both sorted values and indices.
+template <typename DataType, typename IndicesType>
+void thrust_sort(DLTensor* input, DLTensor* out_values, DLTensor* out_indices, bool is_ascend,
+ int n_values) {
+ thrust::device_ptr<DataType> data_ptr(static_cast<DataType*>(input->data));
+ thrust::device_ptr<DataType> values_ptr(static_cast<DataType*>(out_values->data));
+ thrust::device_ptr<IndicesType> indices_ptr(static_cast<IndicesType*>(out_indices->data));
+
+ size_t size = 1;
+ for (int i = 0; i < input->ndim; ++i) {
+ size *= input->shape[i];
+ }
+ thrust::copy(data_ptr, data_ptr + size, values_ptr);
+
+ if (size == static_cast<size_t>(input->shape[input->ndim - 1])) {
+ // A fast path for single segment case
+ thrust::sequence(indices_ptr, indices_ptr + n_values);
+ if (is_ascend) {
+ thrust::sort_by_key(values_ptr, values_ptr + n_values, indices_ptr);
+ } else {
+ thrust::sort_by_key(values_ptr, values_ptr + n_values, indices_ptr,
+ thrust::greater<DataType>());
+ }
+ } else {
+ // segmented sort by key
+ // Follow the back-to-back stable_sort_by_key strategy explained below
+ // https://groups.google.com/g/thrust-users/c/BoLsxO6b4FY
+ thrust::device_vector<int64_t> argsort_order(size);
+ thrust::sequence(argsort_order.begin(), argsort_order.end());
+
+ // First, sort values and store the sorted order in argsort_order.
+ if (is_ascend) {
+ thrust::stable_sort_by_key(values_ptr, values_ptr + size, argsort_order.begin());
+ } else {
+ thrust::stable_sort_by_key(values_ptr, values_ptr + size, argsort_order.begin(),
+ thrust::greater<DataType>());
+ }
+
+ // The following is to create the indices array 0, 1, 2, 0, 1, 2 ... 0, 1, 2
+ // without materializing it
+ auto counting_iter = thrust::counting_iterator<int64_t>(0);
+ auto linear_index_to_sort_axis_index = [n_values] __host__ __device__(int64_t i) {
+ return i % n_values;
+ }; // NOLINT(*)
+ auto init_indices_iter =
+ thrust::make_transform_iterator(counting_iter, linear_index_to_sort_axis_index);
+
+ // This will reorder indices 0, 1, 2 ... in the sorted order of values_ptr
+ thrust::gather(argsort_order.begin(), argsort_order.end(), init_indices_iter, indices_ptr);
+
+ thrust::device_vector<int> segment_ids(size);
+ auto linear_index_to_segment_id = [n_values] __host__ __device__(int64_t i) {
+ return i / n_values;
+ }; // NOLINT(*)
+ // We also reorder segment indices 0, 0, 0, 1, 1, 1 ... in the order of values_ptr
+ thrust::transform(argsort_order.begin(), argsort_order.end(), segment_ids.begin(),
+ linear_index_to_segment_id);
+
+ // The second sort key-ed by segment_ids would bring segment_ids back to 0, 0, 0, 1, 1, 1 ...
+ // values_ptr and indices_ptr will also be sorted in the order of segmend_ids above
+ // Since sorting has been done in a stable way, relative orderings of values and indices
+ // in the segment do not change and hence they remain sorted.
+ auto key_val_zip = thrust::make_zip_iterator(thrust::make_tuple(values_ptr, indices_ptr));
+ thrust::stable_sort_by_key(segment_ids.begin(), segment_ids.end(), key_val_zip);
+ }
+}
+
+void thrust_sort_common(DLTensor* input, DLTensor* values_out, DLTensor* indices_out,
+ bool is_ascend, int sort_len, std::string data_dtype,
+ std::string out_dtype) {
+ if (data_dtype == "float32") {
+ if (out_dtype == "int32") {
+ thrust_sort<float, int32_t>(input, values_out, indices_out, is_ascend, sort_len);
+ } else if (out_dtype == "int64") {
+ thrust_sort<float, int64_t>(input, values_out, indices_out, is_ascend, sort_len);
+ } else if (out_dtype == "float32") {
+ thrust_sort<float, float>(input, values_out, indices_out, is_ascend, sort_len);
+ } else if (out_dtype == "float64") {
+ thrust_sort<float, double>(input, values_out, indices_out, is_ascend, sort_len);
+ } else {
+ LOG(FATAL) << "Unsupported output dtype: " << out_dtype;
+ }
+ } else if (data_dtype == "float64") {
+ if (out_dtype == "int32") {
+ thrust_sort<double, int32_t>(input, values_out, indices_out, is_ascend, sort_len);
+ } else if (out_dtype == "int64") {
+ thrust_sort<double, int64_t>(input, values_out, indices_out, is_ascend, sort_len);
+ } else if (out_dtype == "float32") {
+ thrust_sort<double, float>(input, values_out, indices_out, is_ascend, sort_len);
+ } else if (out_dtype == "float64") {
+ thrust_sort<double, double>(input, values_out, indices_out, is_ascend, sort_len);
+ } else {
+ LOG(FATAL) << "Unsupported output dtype: " << out_dtype;
+ }
+ } else if (data_dtype == "int32") {
+ if (out_dtype == "int32") {
+ thrust_sort<int32_t, int32_t>(input, values_out, indices_out, is_ascend, sort_len);
+ } else if (out_dtype == "int64") {
+ thrust_sort<int32_t, int64_t>(input, values_out, indices_out, is_ascend, sort_len);
+ } else if (out_dtype == "float32") {
+ thrust_sort<int32_t, float>(input, values_out, indices_out, is_ascend, sort_len);
+ } else if (out_dtype == "float64") {
+ thrust_sort<int32_t, double>(input, values_out, indices_out, is_ascend, sort_len);
+ } else {
+ LOG(FATAL) << "Unsupported output dtype: " << out_dtype;
+ }
+ } else if (data_dtype == "int64") {
+ if (out_dtype == "int32") {
+ thrust_sort<int64_t, int32_t>(input, values_out, indices_out, is_ascend, sort_len);
+ } else if (out_dtype == "int64") {
+ thrust_sort<int64_t, int64_t>(input, values_out, indices_out, is_ascend, sort_len);
+ } else if (out_dtype == "float32") {
+ thrust_sort<int64_t, float>(input, values_out, indices_out, is_ascend, sort_len);
+ } else if (out_dtype == "float64") {
+ thrust_sort<int64_t, double>(input, values_out, indices_out, is_ascend, sort_len);
+ } else {
+ LOG(FATAL) << "Unsupported output dtype: " << out_dtype;
+ }
+ } else {
+ LOG(FATAL) << "Unsupported input dtype: " << data_dtype;
+ }
+}
+
+TVM_REGISTER_GLOBAL("tvm.contrib.thrust.sort").set_body([](TVMArgs args, TVMRetValue* ret) {
+ ICHECK_GE(args.num_args, 4);
+ DLTensor* input = args[0];
+ DLTensor* values_out = args[1];
+ DLTensor* indices_out = args[2];
+ bool is_ascend = args[3];
+
+ auto data_dtype = DLDataType2String(input->dtype);
+ auto out_dtype = DLDataType2String(indices_out->dtype);
+
+ int n_values = input->shape[input->ndim - 1];
+ thrust_sort_common(input, values_out, indices_out, is_ascend, n_values, data_dtype, out_dtype);
+});
+
+template <typename KeyType, typename ValueType>
+void thrust_stable_sort_by_key(DLTensor* keys_in, DLTensor* values_in, DLTensor* keys_out,
+ DLTensor* values_out, bool for_scatter) {
+ const auto size = keys_in->shape[0];
+ thrust::device_ptr<KeyType> keys_in_ptr(static_cast<KeyType*>(keys_in->data));
+ thrust::device_ptr<ValueType> values_in_ptr(static_cast<ValueType*>(values_in->data));
+ thrust::device_ptr<KeyType> keys_out_ptr(static_cast<KeyType*>(keys_out->data));
+ thrust::device_ptr<ValueType> values_out_ptr(static_cast<ValueType*>(values_out->data));
+
+ if (for_scatter) {
+ thrust::transform(keys_in_ptr, keys_in_ptr + size, keys_out_ptr, [size] __device__(KeyType k) {
+ if (k < 0) return k + static_cast<KeyType>(size);
+ return k;
+ });
+ } else {
+ thrust::copy(keys_in_ptr, keys_in_ptr + size, keys_out_ptr);
+ }
+ thrust::copy(values_in_ptr, values_in_ptr + size, values_out_ptr);
+
+ thrust::stable_sort_by_key(keys_out_ptr, keys_out_ptr + size, values_out_ptr);
+}
+
+TVM_REGISTER_GLOBAL("tvm.contrib.thrust.stable_sort_by_key")
+ .set_body([](TVMArgs args, TVMRetValue* ret) {
+ ICHECK_GE(args.num_args, 5);
+ DLTensor* keys_in = args[0];
+ DLTensor* values_in = args[1];
+ DLTensor* keys_out = args[2];
+ DLTensor* values_out = args[3];
+ bool for_scatter = args[4];
+
+ auto key_dtype = DLDataType2String(keys_in->dtype);
+ auto value_dtype = DLDataType2String(values_in->dtype);
+
+ if (key_dtype == "int32") {
+ if (value_dtype == "int32") {
+ thrust_stable_sort_by_key<int, int>(keys_in, values_in, keys_out, values_out,
+ for_scatter);
+ } else if (value_dtype == "int64") {
+ thrust_stable_sort_by_key<int, int64_t>(keys_in, values_in, keys_out, values_out,
+ for_scatter);
+ } else if (value_dtype == "float32") {
+ thrust_stable_sort_by_key<int, float>(keys_in, values_in, keys_out, values_out,
+ for_scatter);
+ } else {
+ LOG(FATAL) << "Unsupported value dtype: " << value_dtype;
Review comment:
For all these errors, could you list the supported types?
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] tkonolige commented on pull request #7458: [ROCM] Add Thrust support
Posted by GitBox <gi...@apache.org>.
tkonolige commented on pull request #7458:
URL: https://github.com/apache/tvm/pull/7458#issuecomment-780192238
@masahi You can use `set_source_files_properties(src/runtime/contrib/thrust/thrust.cu PROPERTIES LANGUAGE CXX)`.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on a change in pull request #7458: [ROCM] Add Thrust support
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #7458:
URL: https://github.com/apache/tvm/pull/7458#discussion_r577182951
##########
File path: python/tvm/relay/op/strategy/rocm.py
##########
@@ -219,3 +221,93 @@ def batch_matmul_strategy_rocm(attrs, inputs, out_type, target):
plevel=12,
)
return strategy
+
+
+def can_use_thrust(target, func_name):
+ return (
+ target.kind.name == "rocm"
+ and "thrust" in target.libs
+ and get_global_func(func_name, allow_missing=True)
+ )
Review comment:
Yes, currently thrust dispatch is done in an ad hoc way, unlike cudnn/cublas etc where `-libs` option is required in addition to `USE_CUDNN=ON` etc.
@tkonolige and I discussed this issue a bit and we agreed that we should enforce `-libs=thrust` to enable thrust offload too. For rocm, this requirement is already enforced in this PR, but for cuda it is not. That's what `TODO(masahi)` will going to fix in the next PR.
We can discuss what the common helper should be / where it should be in the next PR, would you agree?
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on a change in pull request #7458: [ROCM] Add Thrust support
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #7458:
URL: https://github.com/apache/tvm/pull/7458#discussion_r577182951
##########
File path: python/tvm/relay/op/strategy/rocm.py
##########
@@ -219,3 +221,93 @@ def batch_matmul_strategy_rocm(attrs, inputs, out_type, target):
plevel=12,
)
return strategy
+
+
+def can_use_thrust(target, func_name):
+ return (
+ target.kind.name == "rocm"
+ and "thrust" in target.libs
+ and get_global_func(func_name, allow_missing=True)
+ )
Review comment:
Yes, currently thrust dispatch is done in an ad hoc way, unlike cudnn/cublas etc where `-libs` option is required in addition to `USE_CUDNN=ON` etc.
@tkonolige and I discussed this issue a bit and we agreed that we should enforce `-libs=thrust` to enable thrust offload too. For rocm, this requirement is already enforced in this PR, but for cuda it is not. That's what `TODO(masahi)` will going to fix in the next PR.
We can discuss what the common helper should be / where it should be placed in the next PR, would you agree?
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on pull request #7458: [ROCM] Add Thrust support
Posted by GitBox <gi...@apache.org>.
masahi commented on pull request #7458:
URL: https://github.com/apache/tvm/pull/7458#issuecomment-780290281
Thanks @mbrookhart @tkonolige @csullivan
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] csullivan commented on a change in pull request #7458: [ROCM] Add Thrust support
Posted by GitBox <gi...@apache.org>.
csullivan commented on a change in pull request #7458:
URL: https://github.com/apache/tvm/pull/7458#discussion_r577185869
##########
File path: python/tvm/relay/op/strategy/rocm.py
##########
@@ -219,3 +221,93 @@ def batch_matmul_strategy_rocm(attrs, inputs, out_type, target):
plevel=12,
)
return strategy
+
+
+def can_use_thrust(target, func_name):
+ return (
+ target.kind.name == "rocm"
+ and "thrust" in target.libs
+ and get_global_func(func_name, allow_missing=True)
+ )
Review comment:
Works for me. :shipit:
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] csullivan commented on a change in pull request #7458: [ROCM] Add Thrust support
Posted by GitBox <gi...@apache.org>.
csullivan commented on a change in pull request #7458:
URL: https://github.com/apache/tvm/pull/7458#discussion_r577148555
##########
File path: python/tvm/relay/op/strategy/rocm.py
##########
@@ -219,3 +221,93 @@ def batch_matmul_strategy_rocm(attrs, inputs, out_type, target):
plevel=12,
)
return strategy
+
+
+def can_use_thrust(target, func_name):
+ return (
+ target.kind.name == "rocm"
+ and "thrust" in target.libs
+ and get_global_func(func_name, allow_missing=True)
+ )
Review comment:
Thoughts on perhaps moving this helper to `python/tvm/contrib/rocm.py`? There are a few examples of helpers in python/tvm/contrib that help with op strategy dispatch based on conditions like these.
##########
File path: python/tvm/relay/op/strategy/rocm.py
##########
@@ -219,3 +221,93 @@ def batch_matmul_strategy_rocm(attrs, inputs, out_type, target):
plevel=12,
)
return strategy
+
+
+def can_use_thrust(target, func_name):
+ return (
+ target.kind.name == "rocm"
+ and "thrust" in target.libs
+ and get_global_func(func_name, allow_missing=True)
+ )
Review comment:
Or given a few of your TODOs in this PR, like
```
# TODO(masahi): Check -libs=thrust option
```
maybe we want a common helper for use with both cuda and rocm?
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on pull request #7458: [ROCM] Add Thrust support
Posted by GitBox <gi...@apache.org>.
masahi commented on pull request #7458:
URL: https://github.com/apache/tvm/pull/7458#issuecomment-780153529
@mbrookhart Yes, I wanted to reuse `thrust.cu` as is, by adding the following to `ROCM.cmake`:
```
file(GLOB CONTRIB_THRUST_SRC src/runtime/contrib/thrust/*.cu)
```
But unlike for CUDA, cmake doesn't recognize that this .cu file needs to be compiled by the same CXX compiler (`hipcc`), so it will not be compiled. I think CMake has a special support for NVCC and .cu file.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on pull request #7458: [ROCM] Add Thrust support
Posted by GitBox <gi...@apache.org>.
masahi commented on pull request #7458:
URL: https://github.com/apache/tvm/pull/7458#issuecomment-780193738
> @masahi You can use `set_source_files_properties(src/runtime/contrib/thrust/thrust.cu PROPERTIES LANGUAGE CXX)`.
Thanks I'll try this ASAP
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on a change in pull request #7458: [ROCM] Add Thrust support
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #7458:
URL: https://github.com/apache/tvm/pull/7458#discussion_r577182951
##########
File path: python/tvm/relay/op/strategy/rocm.py
##########
@@ -219,3 +221,93 @@ def batch_matmul_strategy_rocm(attrs, inputs, out_type, target):
plevel=12,
)
return strategy
+
+
+def can_use_thrust(target, func_name):
+ return (
+ target.kind.name == "rocm"
+ and "thrust" in target.libs
+ and get_global_func(func_name, allow_missing=True)
+ )
Review comment:
Yes, currently thrust dispatch is done in a ad hoc way, unlike cudnn/cublas etc where `-libs` option is required in addition to `USE_CUDNN=ON` etc.
@tkonolige and I discussed this issue a bit and we agreed that we should enforce `-libs=thrust` to enable thrust offload too. For rocm, this requirement is already enforced in this PR, but for cuda it is not. That's what `TODO(masahi)` will going to fix in the next PR.
We can discuss what the common helper should be / where it should be in the next PR, would you agree?
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi edited a comment on pull request #7458: [ROCM] Add Thrust support
Posted by GitBox <gi...@apache.org>.
masahi edited a comment on pull request #7458:
URL: https://github.com/apache/tvm/pull/7458#issuecomment-780153529
@mbrookhart Yes, I wanted to reuse `thrust.cu` as is, by adding the following to `ROCM.cmake`:
```
file(GLOB CONTRIB_THRUST_SRC src/runtime/contrib/thrust/*.cu)
```
But unlike for CUDA, cmake doesn't recognize that this .cu file needs to be compiled by the same CXX compiler (`hipcc`) together with other .cc files, so it will not be compiled (it is ignored during compilation).
I think CMake has a special support for NVCC and .cu file, to correctly compile .cu with NVCC and others with gcc/clang. On rocm, everything needs to be compiled by the same compiler, `hipcc` (a wrapper around clang).
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi edited a comment on pull request #7458: [ROCM] Add Thrust support
Posted by GitBox <gi...@apache.org>.
masahi edited a comment on pull request #7458:
URL: https://github.com/apache/tvm/pull/7458#issuecomment-780153529
@mbrookhart Yes, I wanted to reuse `thrust.cu` as is, by adding the following to `ROCM.cmake`:
```
file(GLOB CONTRIB_THRUST_SRC src/runtime/contrib/thrust/*.cu)
```
But unlike for CUDA, cmake doesn't recognize that this .cu file needs to be compiled by the same CXX compiler (`hipcc`) together with other .cc files, so it will not be compiled.
I think CMake has a special support for NVCC and .cu file, to correctly compile .cu with NVCC and others with gcc/clang. On rocm, everything needs to be compiled by the same compiler, `hipcc` (a wrapper around clang).
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on pull request #7458: [ROCM] Add Thrust support
Posted by GitBox <gi...@apache.org>.
masahi commented on pull request #7458:
URL: https://github.com/apache/tvm/pull/7458#issuecomment-779816625
@tkonolige @mbrookhart Thanks for review, it should be ready to go. This can be used for benchmarking the new TIR sort @mbrookhart is cooking.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] mbrookhart commented on pull request #7458: [ROCM] Add Thrust support
Posted by GitBox <gi...@apache.org>.
mbrookhart commented on pull request #7458:
URL: https://github.com/apache/tvm/pull/7458#issuecomment-780191065
@tkonolige @csullivan Any other change requests or are you happy?
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on pull request #7458: [ROCM] Add Thrust support
Posted by GitBox <gi...@apache.org>.
masahi commented on pull request #7458:
URL: https://github.com/apache/tvm/pull/7458#issuecomment-780165445
Maybe pytorch people have a solution to the duplication problem. But my cmake-fu is too low and I have no interest in learning about cmake :)
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi edited a comment on pull request #7458: [ROCM] Add Thrust support
Posted by GitBox <gi...@apache.org>.
masahi edited a comment on pull request #7458:
URL: https://github.com/apache/tvm/pull/7458#issuecomment-780153529
@mbrookhart Yes, I wanted to reuse `thrust.cu` as is, by adding the following to `ROCM.cmake`:
```
file(GLOB CONTRIB_THRUST_SRC src/runtime/contrib/thrust/*.cu)
```
But unlike for CUDA, cmake doesn't recognize that this .cu file needs to be compiled by the same CXX compiler (`hipcc`) together with other .cc files, so it will not be compiled. I think CMake has a special support for NVCC and .cu file.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi edited a comment on pull request #7458: [ROCM] Add Thrust support
Posted by GitBox <gi...@apache.org>.
masahi edited a comment on pull request #7458:
URL: https://github.com/apache/tvm/pull/7458#issuecomment-779816625
@tkonolige @mbrookhart Thanks for review, it should be ready to go. This can be used for benchmarking the new TIR sort @mbrookhart is cooking.
I'll follow up with another PR to enforce `-libs=thrust` option for other thrust usages.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on pull request #7458: [ROCM] Add Thrust support
Posted by GitBox <gi...@apache.org>.
masahi commented on pull request #7458:
URL: https://github.com/apache/tvm/pull/7458#issuecomment-780200850
Great `set_source_files_properties` worked! I removed `rocthrust/thrust.cc`, the same `thrust.cu` is used by CUDA and rocm now cc @mbrookhart
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi merged pull request #7458: [ROCM] Add Thrust support
Posted by GitBox <gi...@apache.org>.
masahi merged pull request #7458:
URL: https://github.com/apache/tvm/pull/7458
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org