thrust 相关问题

在CUDA工具包中包含的Thrust库中找不到thrust/universal_vector.h

我目前正在使用 Thrust 在 GPU 和 CPU 之间传输数据。但是当我在代码中包含 并使用 CMake 配置项目时，“致命错误：没有这样的...

Thrust：使用device_ptr时如何获取copy_if函数复制的元素数量

我正在使用 Thrust 库的 Thrust::copy_if 函数以及计数迭代器来获取数组中非零元素的索引。我还需要获取复制元素的数量。 我...

推力收集/过滤

我想要做的是在向量上创建一个过滤器，以便它删除未通过谓词测试的元素；但不太确定我该怎么做。 我再次评估输入向量中的每个元素...

CUDA强制指令执行顺序

我正在尝试将一些数据操作从CPU传输到GPU（CUDA），但有一小部分需要指令以特定顺序运行。原则上我可以做前几页...

CUDA 二阶递归推力 Include_scan

我试图了解如何并行递归计算。连续地，计算采用以下形式： 对于 (int i = 2; i 我试图了解如何并行递归计算。连续地，计算采用以下形式： for (int i = 2; i<size; i++) { result[i] = oldArray[i] + k * result[i-2]; } 对于i-1索引，我之前的问题有一个解决方案：CUDA强制指令执行顺序 我想修改它以使用i-2，但我不明白如何将相同的过程应用于二阶计算。应该可以使用 thrust::inclusive_scan 函数，但我不知道如何使用。有谁知道解决办法吗 从上一个问题/答案中继续，我们将注意力转移到 Blelloch 的参考论文中的方程 1.11。我们观察到您的问题表述： for (int i = 2; i<size; i++) { result[i] = oldArray[i] + k * result[i-2]; } 如果我们设置 m=2， 似乎与方程 1.11 中的情况相匹配，在这种情况下，我们还可以观察到，对于您的公式，所有 ai,1 均为零（并且，如前所述，所有 ai,2 均为零） k). 根据该论文中的方程 1.12，我们的状态变量 si 现在变成了一个二元组： si = |xi xi-1| 记下这些事情，我们观察方程 1.13 的“正确性”： si = |xi-1 xi-2| 。 |0 1，k 0| + |bi 0| 重写： si,1 = xi = k*xi-2 + bi si,2 = xi-1 = xi-1 （在我看来，其他答案让你在这一点上。这种认识，即result.data[0] = right + k * left.data[1];足以进行串行扫描，但不适用于并行扫描。同样明显的是，函子/扫描操作不具有关联性.) 我们现在需要提出一个二元运算符bop，它是（1.7）中定义对这种情况的扩展。参考之前方程1.7中的定义，我们在1.13中的处理基础上将其扩展如下： Ci = |Ai , Bi| 地点： Ai = |0 1, k 0| 和： Bi = |bi 0| 然后我们有： Ci bop Cj = | Ai。 Aj，Bi。 Aj + Bj | 这将成为我们函子/扫描运算符的公式。我们需要始终携带 6 个标量“状态”量：2 个用于 B 向量，4 个用于 A 矩阵。 接下来就是上面的实现： $ cat t1930.cu #include <iostream> #include <thrust/device_vector.h> #include <thrust/host_vector.h> #include <thrust/scan.h> #include <thrust/copy.h> #include <thrust/iterator/zip_iterator.h> #include <thrust/iterator/constant_iterator.h> #include <cstdlib> #include <cstdio> template <typename T> void cpufunction(T *result, T *oldArray, size_t size, T k){ for (int i = 2; i<size; i++) { result[i] = oldArray[i] + k * result[i-2]; } } struct scan_op // as per blelloch (1.7) { template <typename T1, typename T2> __host__ __device__ T1 operator()(const T1 &t1, const T2 &t2){ T1 ret; thrust::get<0>(ret) = thrust::get<0>(t1)*thrust::get<2>(t2) + thrust::get<1>(t1)*thrust::get<4>(t2)+thrust::get<0>(t2); thrust::get<1>(ret) = thrust::get<0>(t1)*thrust::get<3>(t2) + thrust::get<1>(t1)*thrust::get<5>(t2)+thrust::get<1>(t2); thrust::get<2>(ret) = thrust::get<2>(t1)*thrust::get<2>(t2) + thrust::get<3>(t1)*thrust::get<4>(t2); thrust::get<3>(ret) = thrust::get<2>(t1)*thrust::get<3>(t2) + thrust::get<3>(t1)*thrust::get<5>(t2); thrust::get<4>(ret) = thrust::get<4>(t1)*thrust::get<2>(t2) + thrust::get<5>(t1)*thrust::get<4>(t2); thrust::get<5>(ret) = thrust::get<4>(t1)*thrust::get<3>(t2) + thrust::get<5>(t1)*thrust::get<5>(t2); return ret; } }; typedef float mt; const size_t ds = 512; const mt k = 1.01; const int snip = 10; int main(){ mt *b1 = new mt[ds]; // b as in blelloch (1.5) mt *cr = new mt[ds]; // cpu result for (int i = 0; i < ds; i++) { b1[i] = rand()/(float)RAND_MAX;} cr[0] = b1[0]; cr[1] = b1[1]; cpufunction(cr, b1, ds, k); for (int i = 0; i < snip; i++) std::cout << cr[i] << ","; for (int i = ds-snip; i < ds; i++) std::cout << cr[i] << ","; std::cout << std::endl; thrust::device_vector<mt> db(b1, b1+ds); auto b0 = thrust::constant_iterator<mt>(0); auto a0 = thrust::constant_iterator<mt>(0); auto a1 = thrust::constant_iterator<mt>(1); auto a2 = thrust::constant_iterator<mt>(k); auto a3 = thrust::constant_iterator<mt>(0); thrust::device_vector<mt> dx1(ds); thrust::device_vector<mt> dx0(ds); thrust::device_vector<mt> dy0(ds); thrust::device_vector<mt> dy1(ds); thrust::device_vector<mt> dy2(ds); thrust::device_vector<mt> dy3(ds); auto my_i_zip = thrust::make_zip_iterator(thrust::make_tuple(db.begin(), b0, a0, a1, a2, a3)); auto my_o_zip = thrust::make_zip_iterator(thrust::make_tuple(dx1.begin(), dx0.begin(), dy0.begin(), dy1.begin(), dy2.begin(), dy3.begin())); thrust::inclusive_scan(my_i_zip, my_i_zip+ds, my_o_zip, scan_op()); thrust::host_vector<mt> hx1 = dx1; thrust::copy_n(hx1.begin(), snip, std::ostream_iterator<mt>(std::cout, ",")); thrust::copy_n(hx1.begin()+ds-snip, snip, std::ostream_iterator<mt>(std::cout, ",")); std::cout << std::endl; } $ nvcc -std=c++14 t1930.cu -o t1930 $ cuda-memcheck ./t1930 ========= CUDA-MEMCHECK 0.840188,0.394383,1.63169,1.19677,2.55965,1.40629,2.92047,2.18858,3.22745,2.76443,570.218,601.275,576.315,607.993,582.947,614.621,589.516,621.699,595.644,628.843, 0.840188,0.394383,1.63169,1.19677,2.55965,1.40629,2.92047,2.18858,3.22745,2.76443,570.219,601.275,576.316,607.994,582.948,614.621,589.516,621.7,595.644,628.843, ========= ERROR SUMMARY: 0 errors $ 是的，上面有一些结果在第 6 位数字上有所不同。考虑到串行和并行方法之间非常不同的操作顺序，我将此归因于 float 分辨率的限制。如果您将 typedef 更改为 double，结果将看起来完全匹配。 既然您已经询问过它，这里有一个等效的实现，它是使用之前使用 cudaMalloc 分配的设备数据进行演示的： $ cat t1930.cu #include <iostream> #include <thrust/device_vector.h> #include <thrust/host_vector.h> #include <thrust/scan.h> #include <thrust/copy.h> #include <thrust/iterator/zip_iterator.h> #include <thrust/iterator/constant_iterator.h> #include <cstdlib> #include <cstdio> template <typename T> void cpufunction(T *result, T *oldArray, size_t size, T k){ for (int i = 2; i<size; i++) { result[i] = oldArray[i] + k * result[i-2]; } } struct scan_op // as per blelloch (1.7) { template <typename T1, typename T2> __host__ __device__ T1 operator()(const T1 &t1, const T2 &t2){ T1 ret; thrust::get<0>(ret) = thrust::get<0>(t1)*thrust::get<2>(t2) + thrust::get<1>(t1)*thrust::get<4>(t2)+thrust::get<0>(t2); thrust::get<1>(ret) = thrust::get<0>(t1)*thrust::get<3>(t2) + thrust::get<1>(t1)*thrust::get<5>(t2)+thrust::get<1>(t2); thrust::get<2>(ret) = thrust::get<2>(t1)*thrust::get<2>(t2) + thrust::get<3>(t1)*thrust::get<4>(t2); thrust::get<3>(ret) = thrust::get<2>(t1)*thrust::get<3>(t2) + thrust::get<3>(t1)*thrust::get<5>(t2); thrust::get<4>(ret) = thrust::get<4>(t1)*thrust::get<2>(t2) + thrust::get<5>(t1)*thrust::get<4>(t2); thrust::get<5>(ret) = thrust::get<4>(t1)*thrust::get<3>(t2) + thrust::get<5>(t1)*thrust::get<5>(t2); return ret; } }; typedef double mt; const size_t ds = 512; const mt k = 1.01; const int snip = 10; int main(){ mt *b1 = new mt[ds]; // b as in blelloch (1.5) mt *cr = new mt[ds]; // cpu result for (int i = 0; i < ds; i++) { b1[i] = rand()/(float)RAND_MAX;} cr[0] = b1[0]; cr[1] = b1[1]; cpufunction(cr, b1, ds, k); for (int i = 0; i < snip; i++) std::cout << cr[i] << ","; for (int i = ds-snip; i < ds; i++) std::cout << cr[i] << ","; std::cout << std::endl; mt *db; cudaMalloc(&db, ds*sizeof(db[0])); cudaMemcpy(db, b1, ds*sizeof(db[0]), cudaMemcpyHostToDevice); thrust::device_ptr<mt> dp_db = thrust::device_pointer_cast(db); auto b0 = thrust::constant_iterator<mt>(0); auto a0 = thrust::constant_iterator<mt>(0); auto a1 = thrust::constant_iterator<mt>(1); auto a2 = thrust::constant_iterator<mt>(k); auto a3 = thrust::constant_iterator<mt>(0); thrust::device_vector<mt> dx1(ds); thrust::device_vector<mt> dx0(ds); thrust::device_vector<mt> dy0(ds); thrust::device_vector<mt> dy1(ds); thrust::device_vector<mt> dy2(ds); thrust::device_vector<mt> dy3(ds); auto my_i_zip = thrust::make_zip_iterator(thrust::make_tuple(dp_db, b0, a0, a1, a2, a3)); auto my_o_zip = thrust::make_zip_iterator(thrust::make_tuple(dx1.begin(), dx0.begin(), dy0.begin(), dy1.begin(), dy2.begin(), dy3.begin())); thrust::inclusive_scan(my_i_zip, my_i_zip+ds, my_o_zip, scan_op()); cudaMemcpy(cr, thrust::raw_pointer_cast(dx1.data()), ds*sizeof(cr[0]), cudaMemcpyDeviceToHost); for (int i = 0; i < snip; i++) std::cout << cr[i] << ","; for (int i = ds-snip; i < ds; i++) std::cout << cr[i] << ","; std::cout << std::endl; } $ nvcc -std=c++14 t1930.cu -o t1930 $ cuda-memcheck ./t1930 ========= CUDA-MEMCHECK 0.840188,0.394383,1.63169,1.19677,2.55965,1.40629,2.92047,2.18858,3.22745,2.76443,570.219,601.275,576.316,607.994,582.948,614.622,589.516,621.7,595.645,628.844, 0.840188,0.394383,1.63169,1.19677,2.55965,1.40629,2.92047,2.18858,3.22745,2.76443,570.219,601.275,576.316,607.994,582.948,614.622,589.516,621.7,595.645,628.844, ========= ERROR SUMMARY: 0 errors 这两种方法之间应该没有显着的性能差异。 （但是，在本例中，我碰巧将 typedef 切换为 double，因此这会有所不同。）使用 cudaMalloc 作为各种状态向量（device_vector、dx0）的 dx1 的替代品、 dy0、dy1 ...）可能会稍微快一些，因为 device_vector 首先执行 cudaMalloc 样式分配，然后启动内核将分配归零。对于状态向量来说，这个归零步骤是不必要的。如果您有兴趣，这里给出的模式应该演示如何做到这一点。 这是一个完全消除使用 thrust::device_vector 和 thrust::host_vector 的版本： #include <iostream> #include <thrust/device_ptr.h> #include <thrust/scan.h> #include <thrust/iterator/zip_iterator.h> #include <thrust/iterator/constant_iterator.h> #include <cstdlib> template <typename T> void cpufunction(T *result, T *oldArray, size_t size, T k){ for (int i = 2; i<size; i++) { result[i] = oldArray[i] + k * result[i-2]; } } struct scan_op // as per blelloch (1.7) { template <typename T1, typename T2> __host__ __device__ T1 operator()(const T1 &t1, const T2 &t2){ T1 ret; thrust::get<0>(ret) = thrust::get<0>(t1)*thrust::get<2>(t2) + thrust::get<1>(t1)*thrust::get<4>(t2)+thrust::get<0>(t2); thrust::get<1>(ret) = thrust::get<0>(t1)*thrust::get<3>(t2) + thrust::get<1>(t1)*thrust::get<5>(t2)+thrust::get<1>(t2); thrust::get<2>(ret) = thrust::get<2>(t1)*thrust::get<2>(t2) + thrust::get<3>(t1)*thrust::get<4>(t2); thrust::get<3>(ret) = thrust::get<2>(t1)*thrust::get<3>(t2) + thrust::get<3>(t1)*thrust::get<5>(t2); thrust::get<4>(ret) = thrust::get<4>(t1)*thrust::get<2>(t2) + thrust::get<5>(t1)*thrust::get<4>(t2); thrust::get<5>(ret) = thrust::get<4>(t1)*thrust::get<3>(t2) + thrust::get<5>(t1)*thrust::get<5>(t2); return ret; } }; typedef float mt; const size_t ds = 32768*4; const mt k = 1.001; const int snip = 10; int main(){ mt *b1 = new mt[ds]; // b as in blelloch (1.5) mt *cr = new mt[ds]; // result for (int i = 0; i < ds; i++) { b1[i] = (rand()/(float)RAND_MAX)-0.5;} cr[0] = b1[0]; cr[1] = b1[1]; cpufunction(cr, b1, ds, k); for (int i = 0; i < snip; i++) std::cout << cr[i] << ","; for (int i = ds-snip; i < ds; i++) std::cout << cr[i] << ","; std::cout << std::endl; mt *db, *dstate; cudaMalloc(&db, ds*sizeof(db[0])); cudaMalloc(&dstate, 6*ds*sizeof(dstate[0])); cudaMemcpy(db, b1, ds*sizeof(db[0]), cudaMemcpyHostToDevice); thrust::device_ptr<mt> dp_db = thrust::device_pointer_cast(db); auto b0 = thrust::constant_iterator<mt>(0); auto a0 = thrust::constant_iterator<mt>(0); auto a1 = thrust::constant_iterator<mt>(1); auto a2 = thrust::constant_iterator<mt>(k); auto a3 = thrust::constant_iterator<mt>(0); thrust::device_ptr<mt> dx1 = thrust::device_pointer_cast(dstate); thrust::device_ptr<mt> dx0 = thrust::device_pointer_cast(dstate+ds); thrust::device_ptr<mt> dy0 = thrust::device_pointer_cast(dstate+2*ds); thrust::device_ptr<mt> dy1 = thrust::device_pointer_cast(dstate+3*ds); thrust::device_ptr<mt> dy2 = thrust::device_pointer_cast(dstate+4*ds); thrust::device_ptr<mt> dy3 = thrust::device_pointer_cast(dstate+5*ds); auto my_i_zip = thrust::make_zip_iterator(thrust::make_tuple(dp_db, b0, a0, a1, a2, a3)); auto my_o_zip = thrust::make_zip_iterator(thrust::make_tuple(dx1, dx0, dy0, dy1, dy2, dy3)); thrust::inclusive_scan(my_i_zip, my_i_zip+ds, my_o_zip, scan_op()); cudaMemcpy(cr, dstate, ds*sizeof(cr[0]), cudaMemcpyDeviceToHost); for (int i = 0; i < snip; i++) std::cout << cr[i] << ","; for (int i = ds-snip; i < ds; i++) std::cout << cr[i] << ","; std::cout << std::endl; } 这里是一些 CPU 代码，显示了源自 https://www.cs.cmu.edu/~guyb/papers/Ble93.pdf 的公式的可能实现，以将高阶递归表示为扫描操作。 关键思想是扫描结果的每个元素不是标量，而是包含前 n 个标量结果的向量。这样，扫描运算符中就可以使用所有必需的先前结果来计算下一个结果。 #include <iostream> #include <algorithm> #include <numeric> #include <array> void calculate1(std::vector<int> vec, int k){ std::vector<int> result(vec.size(), 0); for(int i = 2; i < vec.size(); i++){ result[i] = vec[i] + k * result[i-2]; } std::cerr << "calculate1 result: "; for(auto x : result){ std::cerr << x << ", "; } std::cerr << "\n"; } struct S{ //data[0] stores result of last iteration //data[1] stores result of second last iteration std::array<int, 2> data; }; std::ostream& operator<<(std::ostream& os, S s){ os << "(" << s.data[0] << "," << s.data[1] << ")"; } void calculate2(std::vector<int> vec, int k){ S initvalue{{0,0}}; std::vector<S> result(vec.size(), initvalue); std::exclusive_scan( vec.begin() + 2, vec.end(), result.begin(), initvalue, [k](S left, int right){ S result; /*A = ( 0 1 k 0 ) Compute result = left * A + (right 0) */ result.data[0] = right + k * left.data[1]; result.data[1] = left.data[0]; return result; } ); std::cerr << "calculate2 result: "; for(auto x : result){ std::cerr << x << ", "; } std::cerr << "\n"; } int main(){ const int k = 5; const std::vector<int> vec1{1,3,5,7,9,11,3,6,7,1,2,4}; calculate1(vec1, k); calculate2(vec1, k); } https://godbolt.org/z/cszzn8Ec8 输出： calculate1 result: 0, 0, 5, 7, 34, 46, 173, 236, 872, 1181, 4362, 5909, calculate2 result: (0,0), (5,0), (7,5), (34,7), (46,34), (173,46), (236,173), (872,236), (1181,872), (4362,1181), (0,0), (0,0), 某处仍然存在逐一错误，但人们可以理解其背后的想法。 我之前说过这种方法可以用于 CUDA 中的并行扫描。这是不正确的。对于并行扫描，扫描算子必须具有一个附加属性，即结合性，即 (a OP b) OP c == a OP (b OP c)。这种方法的情况并非如此。 Robert Crovella 的答案展示了如何导出可用于并行扫描的关联扫描算子。

thrust::transform() 导致从主机到设备的 cudaErrorIllegalAddress

以下test.cu程序 #包括 #包括 #包括 #包括 #包括 以下test.cu程序 #include <thrust/copy.h> #include <thrust/execution_policy.h> #include <thrust/transform.h> #include <thrust/device_vector.h> #include <thrust/host_vector.h> #include <iostream> using HOST_TYPE=int32_t; using DEVICE_TYPE=int; template <typename T> struct Cast { __host__ __device__ T operator()(HOST_TYPE i) const { return static_cast<T>(i); } }; int main() { // Initialize host data thrust::host_vector<HOST_TYPE> const h_vec{1, 2, 3, 4, 5}; // Allocate space on the device thrust::device_vector<DEVICE_TYPE> device_data(h_vec.size()); // Copy data from host to device //thrust::copy(h_vec.cbegin(), h_vec.cend(), device_data.begin()); // this works thrust::transform(h_vec.cbegin(), h_vec.cend(), device_data.begin(), Cast<DEVICE_TYPE>{}); // Copy back to host to check thrust::host_vector<DEVICE_TYPE> host_data_copy = device_data; for (DEVICE_TYPE val : host_data_copy) { std::cout << val << " "; } std::cout << std::endl; return 0; } 原因 $ nvcc test.cu $ ./a.out terminate called after throwing an instance of 'thrust::system::system_error' what(): parallel_for: failed to synchronize: cudaErrorIllegalAddress: an illegal memory access was encountered Aborted (core dumped) 这发生在生产线上 thrust::transform(h_vec.cbegin(), h_vec.cend(), device_data.begin(), Cast<DEVICE_TYPE>{}); 即使类似的thrust::copy()运行良好： thrust::copy(h_vec.cbegin(), h_vec.cend(), device_data.begin()); // this works 我在文档中找不到任何内容说明thrust::transform()不应在设备和主机之间转换数据。我是不是在某个地方错过了这个？ 使用 thrust::host 或 thrust::device 执行策略没有帮助。 版本： $ nvcc --version nvcc: NVIDIA (R) Cuda compiler driver Copyright (c) 2005-2023 NVIDIA Corporation Built on Fri_Nov__3_17:16:49_PDT_2023 Cuda compilation tools, release 12.3, V12.3.103 Build cuda_12.3.r12.3/compiler.33492891_0 注意：实际应用需要 HOST_TYPE=char，但为了调试/说明目的，将其更改为 HOST_TYPE=int32_t，并与 std::copy() 进行比较。 请参阅转换： 除了可以在主机和设备之间复制数据的 Thrust::copy 之外，Thrust 算法的所有迭代器参数都应该位于同一位置：要么全部位于主机上，要么全部位于设备上。当违反此要求时，编译器将产生错误消息。

获取 CUDA Thrust 以使用您选择的 CUDA 流

查看 CUDA Thrust 代码中的内核启动，似乎它们总是使用默认流。我可以让 Thrust 使用我选择的流吗？我是否遗漏了 API 中的某些内容？

在 CUDA 11.5 中使用推力的 nvcc 编译错误

我最初在尝试编译包含 的代码时询问了一个错误，但通过添加标志 -std=c++14 解决了这个问题，如 nvcc -std=c++14 test1.cu 中所示。然而，现在我明白了...

CUDA 上的配对重复数据删除

我已经在 CUDA 上运行了一个数据结构并收集数据如下： 结构SearchDataOnDevice { size_t npair； int * id1; int * id2; }; 我想删除重复的 id p...

如何在给定 Thrust 中的另一个向量的情况下查找向量的索引

我有两个推力装置矢量，假设a和b。我想找到向量 a 的索引，其中它小于/大于向量 b 的成对分量 ($a_{ij}>b_{ij}$)。 我发现...

循环旋转GPU矢量？

我有一个想要实现的算法，其中涉及 坐标加法， 坐标乘法，以及 坐标的循环旋转。 我的加法和乘法有点

使用 CUDA Thrust 在向量中进行替换/合并操作

我有两个使用 CUDA Thrust 操作设备向量中的元素的操作。哪些方法可以更高效地实现这两个功能？ 批量替换向量的部分值...

CUDA Thrust 如何在没有物化数据的情况下组合 copy_if 和转换

假设我们有两个输入，第一个是数组，第二个是位图 推力::device_vector点； 位组位； // 假设可以在内核中访问它。 什...

在 cuda Thrust 中融合两个归约操作

有没有办法在 Thrust 的一个内核调用中执行一个 reduce_by_key 操作和一个 reduce（或者理想情况下另一个 reduce_by_key）操作？除了获得计算速度，让我们说我想...

使用 thrust 时，在 __host__ __device__ 仿函数中创建 std::array 是否合法？

我写了一个玩具代码来测试一些想法 #包括 #包括 #包括 #包括 #包括<

进一步优化 CUDA 内核的 Thrust 操作的机会

我有一个 CUDA 内核，基本上如下所示。 __global__ void myOpKernel(double *vals, double *data, int *nums, double *crit, int N, int K) { int index = blockIdx.x*blockDim.x +

如何使用 CUDA thrust 进行分段缩减？

我想将部分归约结果存储在一个数组中。 假设我有数据 [8] = {10,20,30,40,50,60,70,80}。 如果我将数据划分为 chunk_size 2，则块将为 {10,20}, {30,40}, ... , ...

如何防止 thrust::reduce_by_key 写入可分页内存？

我正在编写一个使用多个并发 CUDA 流的应用程序。当我的 thrust::reduce_by_key 调用似乎写入可分页内存时，我的其他流正在阻塞。我认为返回的 v...

thrust::reduce_by_key() 返回重复的键

这是我的代码： //初始化设备向量 整数大小 = N; thrust::device_vector 值（大小）； thrust::device_vector key(size); //获取设备的设备指针_...

计算 CUDA 数组中数字的出现次数

我有一个无符号整数数组存储在带有 CUDA 的 GPU 上（通常有 1000000 个元素）。我想计算数组中每个数字的出现次数。只有几个不同的数字......