NVIDIA / NVIDIA/cccl

[BUG]: Why does it slow down after multiple iterations

Open
#2,228 0 comments 0 reactions 0 assignees View on GitHub
Dominant language
C++
Stars
2.5k
Forks
486
Avg merge
2d 6h
Merged PRs (30d)
295

Description

### Is this a duplicate?

- [X] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this bug and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)

### Type of Bug

Performance

### Component

Thrust

### Describe the bug

```c++
#include
#include
#include
#include
#include
#include
#include
#include
#include
#include
#include
#include
using namespace std;
using namespace chrono;
class Timer {
private:
std::chrono::time_point startTime, endTime;
std::chrono::duration duration_double;
std::chrono::milliseconds duration_milliseconds;
std::chrono::microseconds duration_microseconds;
std::string m_infoMessage;
public:
Timer(const std::string& infoMessage) {
m_infoMessage = infoMessage;
}
void start() {
startTime = std::chrono::system_clock::now();
}
void ends() {
endTime = std::chrono::system_clock::now();
duration_double = std::chrono::duration_cast>(endTime - startTime);
std::cout << m_infoMessage << " cost time: " << duration_double.count() << " 秒\n";
}
void endms() {
endTime = std::chrono::system_clock::now();
duration_milliseconds = std::chrono::duration_cast(endTime - startTime);
std::cout << m_infoMessage << " cost time: " << duration_milliseconds.count() << " 毫秒\n";
}
void endns() {
endTime = std::chrono::system_clock::now();
duration_microseconds = std::chrono::duration_cast(endTime - startTime);
std::cout << m_infoMessage << " cost time: " << duration_microseconds.count() << " 微秒\n";
}
};

void torch_unique(const thrust::device_vector& d_raw_value, thrust::device_vector& d_unique_value,
thrust::device_vector& d_unique_counts, thrust::device_vector& d_inverse_indices,
thrust::device_vector& sorted_values,
thrust::pair::iterator, thrust::device_vector::iterator> &new_end, int& num_unique
) {
thrust::copy(d_raw_value.begin(), d_raw_value.end(), sorted_values.begin());
thrust::sort(sorted_values.begin(), sorted_values.end());
new_end = thrust::reduce_by_key(sorted_values.begin(), sorted_values.end(),
thrust::make_constant_iterator(1),
d_unique_value.begin(), d_unique_counts.begin());
num_unique = new_end.first - d_unique_value.begin();
thrust::lower_bound(d_unique_value.begin(), d_unique_value.begin() + num_unique,
d_raw_value.begin(), d_raw_value.end(),
d_inverse_indices.begin());
}

int main() {
int size = 100000; //100000
std::default_random_engine generator(0);
std::uniform_int_distribution distribution(0, 500); // 500
int* raw_value = new int[size];
int* unique_value = new int[size];
int* inverse_indices = new int[size]; int* unique_counts = new int[size];

for (size_t i = 0; i < size; i++) {
raw_value[i] = distribution(generator);
}

thrust::device_vector d_raw_value(size);
thrust::device_vector d_unique_value(size);
thrust::device_vector d_unique_counts(size);
thrust::device_vector d_inverse_indices(size);
thrust::device_vector sorted_values(size);
thrust::pair::iterator, thrust::device_vector::iterator> new_end;
thrust::copy(raw_value, raw_value + size, d_raw_value.begin());
int num_unique;
Timer t1("func");
for (size_t i = 0; i < 100000; i++) {
num_unique = 0;
t1.start();
torch_unique(d_raw_value, d_unique_value, d_unique_counts, d_inverse_indices, sorted_values, new_end, num_unique);
t1.endms();
}
//thrust::copy(d_unique_value.begin(), d_unique_value.begin() + num_unique, unique_value);
//thrust::copy(d_inverse_indices.begin(), d_inverse_indices.end(), inverse_indices);
//thrust::copy(d_unique_counts.begin(), d_unique_counts.begin() + num_unique, unique_counts);

//auto printArray = [](const char* label, const int* arr, int size) {
// std::cout << label << ": ";
// for (int i = 0; i < size; ++i) {
// std::cout << arr[i] << " ";
// }
// std::cout << std::endl;
// };
//printArray("Input", raw_value, size);
//printArray("Unique Value (CUDA)", unique_value, num_unique);
//printArray("Inverse Indices (CUDA)", inverse_indices, size);
//printArray("Unique Counts (CUDA)", unique_counts, num_unique);
//delete[] raw_value; delete[] unique_value; delete[] inverse_indices; delete[] unique_counts;
return 0;
}
```

After about 9000 cycles, it will slow down.

### How to Reproduce

just run

### Expected behavior

Running unchanged after multiple cycles

### Reproduction link

_No response_

### Operating System

windows 11

### nvidia-smi output

Tue Aug 13 15:10:46 2024
+---------------------------------------------------------------------------------------+
| NVIDIA-SMI 546.33 Driver Version: 546.33 CUDA Version: 12.3 |
|-----------------------------------------+----------------------+----------------------+
| GPU Name TCC/WDDM | Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap | Memory-Usage | GPU-Util Compute M. |
| | | MIG M. |
|=========================================+======================+======================|
| 0 NVIDIA GeForce RTX 4060 ... WDDM | 00000000:01:00.0 Off | N/A |
| N/A 41C P8 1W / 105W | 0MiB / 8188MiB | 0% Default |
| | | N/A |
+-----------------------------------------+----------------------+----------------------+

+---------------------------------------------------------------------------------------+
| Processes: |
| GPU GI CI PID Type Process name GPU Memory |
| ID ID Usage |
|=======================================================================================|
| 0 N/A N/A 672 C+G C:\anaconda3\python.exe N/A |
| 0 N/A N/A 6572 C+G C:\anaconda3\python.exe N/A |
| 0 N/A N/A 6616 C+G C:\anaconda3\python.exe N/A |
| 0 N/A N/A 17188 C C:\anaconda3\python.exe N/A |
| 0 N/A N/A 24008 C+G ...\Huorong\Sysdiag\bin\HipsDaemon.exe N/A |
+---------------------------------------------------------------------------------------+

### NVCC version

nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2022 NVIDIA Corporation
Built on Wed_Sep_21_10:41:10_Pacific_Daylight_Time_2022
Cuda compilation tools, release 11.8, V11.8.89
Build cuda_11.8.r11.8/compiler.31833905_0

Contributor guide

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.