为什么 Thrust 库的交集返回意外结果？
Why intersection of Thrust Library is returning unexpected result?
我正在使用 Thrust 库来获取两个较大的整数集的交集。在使用 2 个小输入进行测试时，我得到了正确的结果，但是当我使用两个包含 10^8 和 65535*1024 元素的集合时，我得到了一个空集合。谁能解释一下这个问题？将前两个变量更改为较小的值推力 returns 预期的交集。我的代码如下。
#include <thrust/set_operations.h>
#include <thrust/device_vector.h>
#include <thrust/device_ptr.h>
#include <iostream>
#include <stdio.h>
int main() {
int sizeArrayLonger = 100*1000*1000;
int sizeArraySmaller = 65535*1024;
int length_result = sizeArraySmaller;
int* list = (int*) malloc(4*sizeArrayLonger);
int* list_smaller = (int*) malloc(4*sizeArraySmaller);
int* result = (int*) malloc(4*length_result);
int* list_gpu;
int* list_smaller_gpu;
int* result_gpu;
// THE NEXT TWO FORS TRANSFORMS THE SMALLER ARRAY IN A SUBSET OF THE LARGER ARRAY
for (int i=0; i < sizeArraySmaller; i++) {
list_smaller[i] = i+1;
list[i] = i+1;
}
for (int i=sizeArraySmaller; i < sizeArrayLonger; i++) {
list[i] = i+1;
}
cudaMalloc(&list_gpu, sizeof(int) * sizeArrayLonger);
cudaMalloc(&list_smaller_gpu, sizeof(int) * sizeArraySmaller);
cudaMalloc(&result_gpu, sizeof(int) * length_result);
cudaMemcpy(list_gpu, list, sizeof(int) * sizeArrayLonger, cudaMemcpyHostToDevice);
cudaMemcpy(list_smaller_gpu, list_smaller, sizeof(int) * sizeArraySmaller, cudaMemcpyHostToDevice);
cudaMemset(result_gpu, 0, sizeof(int) * length_result);
typedef thrust::device_ptr<int> device_ptr;
thrust::set_intersection(device_ptr(list_gpu), device_ptr(list_gpu + sizeArrayLonger), device_ptr(list_smaller_gpu),
device_ptr(list_smaller_gpu + sizeArraySmaller), device_ptr(result_gpu), thrust::less<int>() );
// MOVING TO CPU THE MARKER ARRAY OF ELEMENTS OF INTERSECTION SET
cudaMemcpy(result, result_gpu, sizeof(int)*length_result, cudaMemcpyDeviceToHost);
cudaDeviceSynchronize();
// THIS LOOP ITERATES ALL ARRAY NAMED "result" WHERE THE POSITION ARE MARKED WITH 1
int counter = 0;
for (int i=0; i < length_result; i++)
if (result[i]) {
printf("\n-> %d", result[i]);
counter++;
}
printf("\nTHRUST -> Total of elements: %d\n", counter);
cudaDeviceReset();
return 0;
}
看来OP最近没有来过，所以我会为其他读者扩展我的评论。 （我希望得到一些确认，即在编译期间指定正在使用的设备的计算目标也会修复 OP 的观察结果。）
根据我的测试，OP 的代码将：
- 如果为 cc2.0 设备编译通过，并且 运行 在 cc2.0 设备上编译。
- 如果为 cc3.0 设备和 运行 在 cc3.x 设备上编译，则通过。
- 如果为 cc2.0 设备和 运行 在 cc3.x 设备上编译，则失败。
最后这个结果有点不直观。由于 the runtime JIT mechanism.
，通常我们喜欢将使用 PTX（例如 nvcc -arch=sm_20 ... 或类似）编译的 CUDA 代码视为与未来架构向前兼容
然而，有一个陷阱（以及一个与推力相关的问题。）CUDA 代码查询它们实际 运行 正在使用的设备的情况并不少见（例如通过 cudaGetDeviceProperties) 并根据正在使用的设备做出决定（例如内核配置决定）。具体来说，在这种情况下，thrust 在底层启动内核，并根据实际使用的设备决定要为该内核选择的网格 x 维度的大小。对于此参数，CC 2.x 设备限制为 65535，但 CC 3.x 和更高的设备 have a much higher limit。因此，在这种情况下，对于足够大的数据集，如果 thrust 检测到它在 cc3.0 设备上 运行ning，它将为这个特定的内核配置一个大于 65535 的网格 x 维度。（对于一个足够小的数据集，它不会这样做，所以这个可能的错误不会浮出水面。因此，问题与问题大小松散相关。）
如果我们将 cc 2.x 和 cc 3.x PTX（或适当的 SASS）嵌入到二进制文件中，那么仍然不会有问题。但是，如果我们只有 cc2.x PTX 嵌入到二进制文件中，那么 JIT 进程将使用它来创建适合于 运行 在 cc 3.x 设备上的机器代码，如果那是正在使用的设备。 但是这个正向JIT编译的SASS仍然受到CC2.x的限制，包括65535的网格X维度限制。但是cudaGetDevicePropertiesreturn事实是设备是 cc3.x 设备，因此如果将此信息用于此特定决策（可接受的网格 X 尺寸），则此信息将产生误导。
由于此序列，内核配置不正确，内核启动失败并出现特定类型的非粘性 CUDA 运行time API 错误。这种类型的非粘性错误不会破坏 CUDA 上下文，因此仍然允许进一步的 CUDA 操作，并且未来的 CUDA API 调用将不会 return 此错误。为了在 CUDA 内核启动后捕获此类错误，有必要在内核启动后发出 cudaGetLastError() 或 cudaPeekAtLastError() 调用，如 proper cuda error checking 所建议的那样。未能执行此操作意味着错误是 "lost" 并且无法从未来的 CUDA API 调用中发现（cudaGetLastError() 或 cudaPeekAtLastError() 除外），因为它们不会指示此错误的存在或在状态 return 值中启动内核失败。
以上大部分内容都可以通过仔细使用 cuda 分析工具来发现，例如nvprof，在通过和失败的情况下，以及cuda-memcheck。在过去的案例中，cuda-memcheck 没有报告任何错误，分析器显示了对 cudaLaunch 的 8 次调用以及在 GPU 上实际执行的 8 个内核。在失败的案例中，cuda-memcheck 报告了 2 个上述类型的内核启动失败，并且分析器显示了 8 次 cudaLaunch 调用，但实际上只有 6 个内核在 GPU 上执行。当在 cc2.x GPU 上 运行 时，失败的内核配置了 65535 的网格 X 维度，而在 cc3.x 上 运行 时配置了更大的数字GPU.
因此，通过适当的 cuda 错误检查，上述序列虽然不一定合乎需要，但至少会因显式错误而失败。但是 OP 的代码无声地失败了——它 return 在失败的情况下是一个不正确的结果，但是 thrust 不会抛出任何类型的错误。
事实证明，在幕后，对从集合操作（至少是这个，特别是）启动的内核的推力错误检查有这个特殊的错误检查差距。
通过仔细研究分析器输出，我们可以发现哪些文件包含 thrust 在这种情况下用于启动关闭的代码（即内核启动实际来自哪里）。 （您也可以通过仔细跟踪模板序列来解决这个问题。）在特定的失败案例中，我相信内核启动是由 here. If we look at one of the kernel launches there, we see something like this:
引起的
#ifndef __CUDA_ARCH__
kernel<<<(unsigned int) num_blocks, (unsigned int) block_size, (unsigned int) smem_size, stream(thrust::detail::derived_cast(exec))>>>(f);
#else
...
#endif // __CUDA_ARCH__
synchronize_if_enabled("launch_closure_by_value");
synchronize_if_enabled（在此特定代码路径中）将在内核启动后立即调用。可以找到该函数 here:
inline __host__ __device__
void synchronize_if_enabled(const char *message)
{
// XXX this could potentially be a runtime decision
// note we always have to synchronize in __device__ code
#if __THRUST_SYNCHRONOUS || defined(__CUDA_ARCH__)
synchronize(message);
#else
// WAR "unused parameter" warning
(void) message;
#endif
调用 synchronize():
inline __host__ __device__
void synchronize(const char *message)
{
throw_on_error(cudaDeviceSynchronize(), message);
} // end synchronize()
我们在 synchronize() 中看到 throw_on_error 调用 cudaDeviceSynchronize() 消除了先前的非粘性错误 11，表示错误配置的内核启动尝试，事实上 return s cudaSuccess（因为 cudaDeviceSynchronize() 操作本身实际上是成功的。）
所以总结是存在 2 个问题：
Thrust（在这种情况下）对内核启动配置做出运行时间决定，如果执行设备是 cc3.0 或更高版本并且代码是为 cc 编译的，这将是不正确的2.x（仅）。
对这个特定 set_intersection 调用的推力错误检查是有缺陷的，因为它没有适当的机制来捕获非粘性 CUDA 运行时间 API 与错误配置的内核启动相关的错误（错误 11）。
那么，如果您打算 运行 cc3.0 或更高版本的设备。 （当然，您可以指定两者 a cc 2.x和a cc3.x目标，并选择nvcc命令行开关。）引擎盖下的各种启动机制，并非所有（也许大多数）不受此特定缺陷的影响（#2），但（对我而言）似乎这个特定的 set_intersection 调用 受此不足，此时（推力v1.8）。
（对我而言）不清楚是否有系统地解决上述第一个问题 (#1) 的方法。我已将上面的第二个问题 (#2) 提请推力开发人员注意（通过 RFE 或错误报告。）
作为解决方法，推力开发人员可以在他们的推力应用程序中插入对 cudaGetLastError() 的调用（可能在最后），以防止此类错误成为 "silent".
有一件事让我感到困惑，并导致了奇怪的行为（我的 Cuda 程序的某些部分在调用 Thrust 函数后根本没有执行，也没有打印任何错误消息），是许多 Thrust 函数被可选的重载第一个参数 const thrust::detail::execution_policy_base< DerivedPolicy > & exec。在使用 Cuda 和 Thrust 处理 single-source 文件示例时，省略这个可选的第一个参数似乎没问题。但是，当使用多个源文件并在不同源文件的不同函数之间传递指向 GPU 内存中向量的指针时，我发现 Thrust 会无声地失败，除非我将 thrust::device 指定为这些 Thrust 函数的第一个参数。
TL;DR：只要您有这个选项，请将 thrust::device 指定为任何预期使用 GPU 内存的 Thrust 函数的第一个参数，并且 运行 在 GPU 上。
我正在使用 Thrust 库来获取两个较大的整数集的交集。在使用 2 个小输入进行测试时，我得到了正确的结果，但是当我使用两个包含 10^8 和 65535*1024 元素的集合时，我得到了一个空集合。谁能解释一下这个问题？将前两个变量更改为较小的值推力 returns 预期的交集。我的代码如下。
#include <thrust/set_operations.h>
#include <thrust/device_vector.h>
#include <thrust/device_ptr.h>
#include <iostream>
#include <stdio.h>
int main() {
int sizeArrayLonger = 100*1000*1000;
int sizeArraySmaller = 65535*1024;
int length_result = sizeArraySmaller;
int* list = (int*) malloc(4*sizeArrayLonger);
int* list_smaller = (int*) malloc(4*sizeArraySmaller);
int* result = (int*) malloc(4*length_result);
int* list_gpu;
int* list_smaller_gpu;
int* result_gpu;
// THE NEXT TWO FORS TRANSFORMS THE SMALLER ARRAY IN A SUBSET OF THE LARGER ARRAY
for (int i=0; i < sizeArraySmaller; i++) {
list_smaller[i] = i+1;
list[i] = i+1;
}
for (int i=sizeArraySmaller; i < sizeArrayLonger; i++) {
list[i] = i+1;
}
cudaMalloc(&list_gpu, sizeof(int) * sizeArrayLonger);
cudaMalloc(&list_smaller_gpu, sizeof(int) * sizeArraySmaller);
cudaMalloc(&result_gpu, sizeof(int) * length_result);
cudaMemcpy(list_gpu, list, sizeof(int) * sizeArrayLonger, cudaMemcpyHostToDevice);
cudaMemcpy(list_smaller_gpu, list_smaller, sizeof(int) * sizeArraySmaller, cudaMemcpyHostToDevice);
cudaMemset(result_gpu, 0, sizeof(int) * length_result);
typedef thrust::device_ptr<int> device_ptr;
thrust::set_intersection(device_ptr(list_gpu), device_ptr(list_gpu + sizeArrayLonger), device_ptr(list_smaller_gpu),
device_ptr(list_smaller_gpu + sizeArraySmaller), device_ptr(result_gpu), thrust::less<int>() );
// MOVING TO CPU THE MARKER ARRAY OF ELEMENTS OF INTERSECTION SET
cudaMemcpy(result, result_gpu, sizeof(int)*length_result, cudaMemcpyDeviceToHost);
cudaDeviceSynchronize();
// THIS LOOP ITERATES ALL ARRAY NAMED "result" WHERE THE POSITION ARE MARKED WITH 1
int counter = 0;
for (int i=0; i < length_result; i++)
if (result[i]) {
printf("\n-> %d", result[i]);
counter++;
}
printf("\nTHRUST -> Total of elements: %d\n", counter);
cudaDeviceReset();
return 0;
}
看来OP最近没有来过，所以我会为其他读者扩展我的评论。 （我希望得到一些确认，即在编译期间指定正在使用的设备的计算目标也会修复 OP 的观察结果。）
根据我的测试，OP 的代码将：
- 如果为 cc2.0 设备编译通过，并且 运行 在 cc2.0 设备上编译。
- 如果为 cc3.0 设备和 运行 在 cc3.x 设备上编译，则通过。
- 如果为 cc2.0 设备和 运行 在 cc3.x 设备上编译，则失败。
最后这个结果有点不直观。由于 the runtime JIT mechanism.
，通常我们喜欢将使用 PTX（例如nvcc -arch=sm_20 ... 或类似）编译的 CUDA 代码视为与未来架构向前兼容
然而，有一个陷阱（以及一个与推力相关的问题。）CUDA 代码查询它们实际 运行 正在使用的设备的情况并不少见（例如通过 cudaGetDeviceProperties) 并根据正在使用的设备做出决定（例如内核配置决定）。具体来说，在这种情况下，thrust 在底层启动内核，并根据实际使用的设备决定要为该内核选择的网格 x 维度的大小。对于此参数，CC 2.x 设备限制为 65535，但 CC 3.x 和更高的设备 have a much higher limit。因此，在这种情况下，对于足够大的数据集，如果 thrust 检测到它在 cc3.0 设备上 运行ning，它将为这个特定的内核配置一个大于 65535 的网格 x 维度。（对于一个足够小的数据集，它不会这样做，所以这个可能的错误不会浮出水面。因此，问题与问题大小松散相关。）
如果我们将 cc 2.x 和 cc 3.x PTX（或适当的 SASS）嵌入到二进制文件中，那么仍然不会有问题。但是，如果我们只有 cc2.x PTX 嵌入到二进制文件中，那么 JIT 进程将使用它来创建适合于 运行 在 cc 3.x 设备上的机器代码，如果那是正在使用的设备。 但是这个正向JIT编译的SASS仍然受到CC2.x的限制，包括65535的网格X维度限制。但是cudaGetDevicePropertiesreturn事实是设备是 cc3.x 设备，因此如果将此信息用于此特定决策（可接受的网格 X 尺寸），则此信息将产生误导。
由于此序列，内核配置不正确，内核启动失败并出现特定类型的非粘性 CUDA 运行time API 错误。这种类型的非粘性错误不会破坏 CUDA 上下文，因此仍然允许进一步的 CUDA 操作，并且未来的 CUDA API 调用将不会 return 此错误。为了在 CUDA 内核启动后捕获此类错误，有必要在内核启动后发出 cudaGetLastError() 或 cudaPeekAtLastError() 调用，如 proper cuda error checking 所建议的那样。未能执行此操作意味着错误是 "lost" 并且无法从未来的 CUDA API 调用中发现（cudaGetLastError() 或 cudaPeekAtLastError() 除外），因为它们不会指示此错误的存在或在状态 return 值中启动内核失败。
以上大部分内容都可以通过仔细使用 cuda 分析工具来发现，例如nvprof，在通过和失败的情况下，以及cuda-memcheck。在过去的案例中，cuda-memcheck 没有报告任何错误，分析器显示了对 cudaLaunch 的 8 次调用以及在 GPU 上实际执行的 8 个内核。在失败的案例中，cuda-memcheck 报告了 2 个上述类型的内核启动失败，并且分析器显示了 8 次 cudaLaunch 调用，但实际上只有 6 个内核在 GPU 上执行。当在 cc2.x GPU 上 运行 时，失败的内核配置了 65535 的网格 X 维度，而在 cc3.x 上 运行 时配置了更大的数字GPU.
因此，通过适当的 cuda 错误检查，上述序列虽然不一定合乎需要，但至少会因显式错误而失败。但是 OP 的代码无声地失败了——它 return 在失败的情况下是一个不正确的结果，但是 thrust 不会抛出任何类型的错误。
事实证明，在幕后，对从集合操作（至少是这个，特别是）启动的内核的推力错误检查有这个特殊的错误检查差距。
通过仔细研究分析器输出，我们可以发现哪些文件包含 thrust 在这种情况下用于启动关闭的代码（即内核启动实际来自哪里）。 （您也可以通过仔细跟踪模板序列来解决这个问题。）在特定的失败案例中，我相信内核启动是由 here. If we look at one of the kernel launches there, we see something like this:
引起的#ifndef __CUDA_ARCH__
kernel<<<(unsigned int) num_blocks, (unsigned int) block_size, (unsigned int) smem_size, stream(thrust::detail::derived_cast(exec))>>>(f);
#else
...
#endif // __CUDA_ARCH__
synchronize_if_enabled("launch_closure_by_value");
synchronize_if_enabled（在此特定代码路径中）将在内核启动后立即调用。可以找到该函数 here:
inline __host__ __device__
void synchronize_if_enabled(const char *message)
{
// XXX this could potentially be a runtime decision
// note we always have to synchronize in __device__ code
#if __THRUST_SYNCHRONOUS || defined(__CUDA_ARCH__)
synchronize(message);
#else
// WAR "unused parameter" warning
(void) message;
#endif
调用 synchronize():
inline __host__ __device__
void synchronize(const char *message)
{
throw_on_error(cudaDeviceSynchronize(), message);
} // end synchronize()
我们在 synchronize() 中看到 throw_on_error 调用 cudaDeviceSynchronize() 消除了先前的非粘性错误 11，表示错误配置的内核启动尝试，事实上 return s cudaSuccess（因为 cudaDeviceSynchronize() 操作本身实际上是成功的。）
所以总结是存在 2 个问题：
Thrust（在这种情况下）对内核启动配置做出运行时间决定，如果执行设备是 cc3.0 或更高版本并且代码是为 cc 编译的，这将是不正确的2.x（仅）。
对这个特定 set_intersection 调用的推力错误检查是有缺陷的，因为它没有适当的机制来捕获非粘性 CUDA 运行时间 API 与错误配置的内核启动相关的错误（错误 11）。
那么，如果您打算 运行 cc3.0 或更高版本的设备。 （当然，您可以指定两者 a cc 2.x和a cc3.x目标，并选择nvcc命令行开关。）引擎盖下的各种启动机制，并非所有（也许大多数）不受此特定缺陷的影响（#2），但（对我而言）似乎这个特定的 set_intersection 调用 受此不足，此时（推力v1.8）。
（对我而言）不清楚是否有系统地解决上述第一个问题 (#1) 的方法。我已将上面的第二个问题 (#2) 提请推力开发人员注意（通过 RFE 或错误报告。）
作为解决方法，推力开发人员可以在他们的推力应用程序中插入对 cudaGetLastError() 的调用（可能在最后），以防止此类错误成为 "silent".
有一件事让我感到困惑，并导致了奇怪的行为（我的 Cuda 程序的某些部分在调用 Thrust 函数后根本没有执行，也没有打印任何错误消息），是许多 Thrust 函数被可选的重载第一个参数 const thrust::detail::execution_policy_base< DerivedPolicy > & exec。在使用 Cuda 和 Thrust 处理 single-source 文件示例时，省略这个可选的第一个参数似乎没问题。但是，当使用多个源文件并在不同源文件的不同函数之间传递指向 GPU 内存中向量的指针时，我发现 Thrust 会无声地失败，除非我将 thrust::device 指定为这些 Thrust 函数的第一个参数。
TL;DR：只要您有这个选项，请将 thrust::device 指定为任何预期使用 GPU 内存的 Thrust 函数的第一个参数，并且 运行 在 GPU 上。