一些序言：
- 正如 cmets 中已经提到的，1000 的总数组大小对于在 GPU 上运行是一个非常小问题。如果这是您在 GPU 上做的唯一事情，那么您可能不会从 GPU 代码中获得有趣的加速。
-
previous answer 建议使用一维模板。即使模板宽度为 100（或 200），仍然应该可以使用 1D 模板方法。这里唯一真正限制模板大小的因素是共享内存的大小，共享内存可以轻松保存等于每个块的线程数 + 模板（序列）“光环”的 200 的数据大小。但是该主题已经在此处介绍，因此我将介绍另一种方法，它应该能够处理 任意 数组长度和 任意 序列大小。
在这种情况下，看看我们是否可以利用prefix sum 来帮助我们可能是合适的。前缀和存在快速并行方法，因此这很有吸引力。前缀和只是一个数字的输出序列，表示输入序列中所有先前数字的总和。所以如果我有一个这样的数组：
1, 1, 3, 2, 0, 1
独占前缀总和将是：
0, 1, 2, 5, 7, 7, 8
排除这里的意思是当前求和位置包含所有之前的值，但排除当前位置的数据项。对于排他前缀和，请注意，我们可以（如果我们愿意）为比输入序列长度长 1 的序列生成“有用”数据。
前缀和可以很容易地适应您所问的问题。假设对于我上面的序列，我们想要计算 3 的子序列。我们可以取前缀和，并从中减去相同的前缀和序列，将 右 移动 3（序列长度），例如所以：
0, 1, 2, 5, 7, 7, 8
- 0, 1, 2, 5, 7, 7, 8
= 5, 6, 5, 3
在这种情况下，这个序列 (5, 6, 5, 3) 将是所需的答案。
下面是一个完整的例子，只是在推力上。我之所以这样做，是因为在 CUDA 中编写快速并行前缀和并不是一件容易的事，因此我更愿意使用（并建议其他人使用）library 实现。如果您想探索如何编写自己的并行前缀和，可以探索开源的推力，或阅读this 的介绍。
这是一个完整的例子：
$ cat t1279.cu
#include <thrust/device_vector.h>
#include <thrust/scan.h>
#include <thrust/transform.h>
#include <thrust/functional.h>
#include <thrust/copy.h>
#include <iostream>
const int arr_size = 1000;
const int seq_len = 100;
typedef int mytype;
int main(){
// test case
mytype t_arr[] = {80,12,14,5,70,9,26,30,8,12,16,15,60,12,38,32,17,67,19,11,0};
int t_arr_size = sizeof(t_arr)/sizeof(mytype);
int t_seq_len = 5;
thrust::device_vector<mytype> d_arr1(t_arr, t_arr+t_arr_size);
thrust::device_vector<mytype> d_res1(t_arr_size);
thrust::device_vector<mytype> d_out1(t_arr_size-t_seq_len);
thrust::exclusive_scan(d_arr1.begin(), d_arr1.end(), d_res1.begin());
thrust::transform(d_res1.begin()+t_seq_len, d_res1.end(), d_res1.begin(), d_out1.begin(), thrust::minus<mytype>());
thrust::copy_n(d_out1.begin(), t_arr_size-t_seq_len, std::ostream_iterator<mytype>(std::cout, ","));
std::cout << std::endl;
// case with larger array length and larger sequence length
thrust::device_vector<mytype> d_arr(arr_size+1, 1);
thrust::device_vector<mytype> d_res(arr_size+1);
thrust::device_vector<mytype> d_out(arr_size+1-seq_len);
thrust::inclusive_scan(d_arr.begin(), d_arr.end(), d_res.begin());
thrust::transform(d_res.begin()+seq_len, d_res.end(), d_res.begin(), d_out.begin(), thrust::minus<mytype>());
// validate
for (int i = 0; i < arr_size+1-seq_len; i++) {
mytype t = d_out[i];
if (t != seq_len) {std::cout << "mismatch at: " << i << "was: " << t << "should be: " << seq_len << std::endl; return 1;}
}
return 0;
}
$ nvcc -arch=sm_35 -o t1279 t1279.cu
$ ./t1279
181,110,124,140,143,85,92,81,111,115,141,157,159,166,173,146,
$