Ale*_*sen 1 c++ algorithm parallel-processing reduce cuda
我是C++/CUDA 的新手。我尝试通过递归内核的输出（在内核包装器中）来实现并行算法“ reduce ”，它能够处理任何类型的输入大小和线程大小，而不会增加渐近并行运行时间。
例如，在 Cuda 中实现 Max Reduce 是这个问题的最佳答案，当线程大小足够小时，他/她的实现基本上是顺序的。
但是，当我编译和运行它时，我不断收到“分段错误”..？
>> nvcc -o mycode mycode.cu
>> ./mycode
Segmentail fault.
在带有 cuda 6.5 的 K40 上编译
这是内核，基本上与我将检查器链接为“越界”的SO帖子相同，但不同：
#include <stdio.h>
/* -------- KERNEL -------- */
__global__ void reduce_kernel(float * d_out, float * d_in, const int size)
{
// position and threadId
int pos = blockIdx.x * blockDim.x + threadIdx.x;
int tid = threadIdx.x;
// do reduction in global memory
for (unsigned int s = blockDim.x / 2; s>0; s>>=1)
{
if (tid < s)
{
if (pos+s < size) // Handling out of bounds
{
d_in[pos] = d_in[pos] + d_in[pos+s];
}
}
}
// only thread 0 writes result, as thread
if (tid==0)
{
d_out[blockIdx.x] = d_in[pos];
}
}
我提到的内核包装器在 1 个块不包含所有数据时处理。
/* -------- KERNEL WRAPPER -------- */
void reduce(float * d_out, float * d_in, const int size, int num_threads)
{
// setting up blocks and intermediate result holder
int num_blocks = ((size) / num_threads) + 1;
float * d_intermediate;
cudaMalloc(&d_intermediate, sizeof(float)*num_blocks);
// recursively solving, will run approximately log base num_threads times.
do
{
reduce_kernel<<<num_blocks, num_threads>>>(d_intermediate, d_in, size);
// updating input to intermediate
cudaMemcpy(d_in, d_intermediate, sizeof(float)*num_blocks, cudaMemcpyDeviceToDevice);
// Updating num_blocks to reflect how many blocks we now want to compute on
num_blocks = num_blocks / num_threads + 1;
// updating intermediate
cudaMalloc(&d_intermediate, sizeof(float)*num_blocks);
}
while(num_blocks > num_threads); // if it is too small, compute rest.
// computing rest
reduce_kernel<<<1, num_blocks>>>(d_out, d_in, size);
}
初始化输入/输出并创建用于测试的虚假数据的主程序。
/* -------- MAIN -------- */
int main(int argc, char **argv)
{
// Setting num_threads
int num_threads = 512;
// Making bogus data and setting it on the GPU
const int size = 1024;
const int size_out = 1;
float * d_in;
float * d_out;
cudaMalloc(&d_in, sizeof(float)*size);
cudaMalloc((void**)&d_out, sizeof(float)*size_out);
const int value = 5;
cudaMemset(d_in, value, sizeof(float)*size);
// Running kernel wrapper
reduce(d_out, d_in, size, num_threads);
printf("sum is element is: %.f", d_out[0]);
}
我会用你的代码指出一些事情。
作为一般规则/样板，我总是建议使用适当的 cuda 错误检查并使用 运行您的代码cuda-memcheck，无论何时您遇到 cuda 代码问题。然而，这些方法对 seg 错误没有多大帮助，尽管它们稍后可能会有所帮助（见下文）。
实际段错误发生在这条线上：
printf("sum is element is: %.f", d_out[0]);
您违反了 CUDA 编程的基本规则：不得在设备代码中取消引用主机指针，不得在主机代码中取消引用设备指针。后一个条件适用于此。 d_out是一个设备指针（通过 分配cudaMalloc）。如果您尝试在主机代码中取消引用这些指针，则这些指针没有任何意义，这样做会导致段错误。
解决办法是先把数据复制回主机再打印出来：
float result;
cudaMemcpy(&result, d_out, sizeof(float), cudaMemcpyDeviceToHost);
printf("sum is element is: %.f", result);
cudaMalloc在循环中使用同一个变量，不做任何cudaFree操作，不是好的做法，可能会导致长时间运行的循环出现内存不足错误，还可能导致程序内存泄漏，如果这样的构造用于更大的程序：
do
{
...
cudaMalloc(&d_intermediate, sizeof(float)*num_blocks);
}
while...
在这种情况下，我认为更好的方法和微不足道的解决方法是cudaFree d_intermediate在重新分配之前：
do
{
...
cudaFree(d_intermediate);
cudaMalloc(&d_intermediate, sizeof(float)*num_blocks);
}
while...
这可能不是您认为的那样：
const int value = 5;
cudaMemset(d_in, value, sizeof(float)*size);
可能您已经意识到这一点，但是cudaMemset，像 一样memset，对字节数量进行操作。因此，您正在d_in使用对应于的值填充数组0x05050505（并且我不知道该位模式在解释为float数量时对应的是什么）。由于您指的是虚假值，因此您可能已经意识到这一点。但这是一个常见的错误（例如，如果您实际上试图在每个float位置用 5 的值初始化数组），所以我想我会指出它。
您的代码也有其他问题（如果您进行上述修复，然后使用 运行您的代码，您会发现这些问题cuda-memcheck）。要了解如何进行良好的并行缩减，我建议您学习 CUDA 并行缩减示例代码和演示。出于性能原因，不建议并行减少全局内存。
为了完整起见，以下是我发现的一些其他问题：
您的内核代码需要一个适当的__syncthreads()语句，以确保在任何线程进入 for 循环的下一次迭代之前，块中所有线程的工作都已完成。
您对内核中全局内存的最终写入还需要以入界读取位置为条件。否则，您始终启动额外块的策略将允许从此行读取越界（cuda-memcheck将显示这一点）。
reduce函数中循环中的归约逻辑通常是混乱的，需要以多种方式重新处理。
我并不是说这段代码没有缺陷，但它似乎适用于给定的测试用例并产生正确的答案 (1024)：
#include <stdio.h>
/* -------- KERNEL -------- */
__global__ void reduce_kernel(float * d_out, float * d_in, const int size)
{
// position and threadId
int pos = blockIdx.x * blockDim.x + threadIdx.x;
int tid = threadIdx.x;
// do reduction in global memory
for (unsigned int s = blockDim.x / 2; s>0; s>>=1)
{
if (tid < s)
{
if (pos+s < size) // Handling out of bounds
{
d_in[pos] = d_in[pos] + d_in[pos+s];
}
}
__syncthreads();
}
// only thread 0 writes result, as thread
if ((tid==0) && (pos < size))
{
d_out[blockIdx.x] = d_in[pos];
}
}
/* -------- KERNEL WRAPPER -------- */
void reduce(float * d_out, float * d_in, int size, int num_threads)
{
// setting up blocks and intermediate result holder
int num_blocks = ((size) / num_threads) + 1;
float * d_intermediate;
cudaMalloc(&d_intermediate, sizeof(float)*num_blocks);
cudaMemset(d_intermediate, 0, sizeof(float)*num_blocks);
int prev_num_blocks;
// recursively solving, will run approximately log base num_threads times.
do
{
reduce_kernel<<<num_blocks, num_threads>>>(d_intermediate, d_in, size);
// updating input to intermediate
cudaMemcpy(d_in, d_intermediate, sizeof(float)*num_blocks, cudaMemcpyDeviceToDevice);
// Updating num_blocks to reflect how many blocks we now want to compute on
prev_num_blocks = num_blocks;
num_blocks = num_blocks / num_threads + 1;
// updating intermediate
cudaFree(d_intermediate);
cudaMalloc(&d_intermediate, sizeof(float)*num_blocks);
size = num_blocks*num_threads;
}
while(num_blocks > num_threads); // if it is too small, compute rest.
// computing rest
reduce_kernel<<<1, prev_num_blocks>>>(d_out, d_in, prev_num_blocks);
}
/* -------- MAIN -------- */
int main(int argc, char **argv)
{
// Setting num_threads
int num_threads = 512;
// Making non-bogus data and setting it on the GPU
const int size = 1024;
const int size_out = 1;
float * d_in;
float * d_out;
cudaMalloc(&d_in, sizeof(float)*size);
cudaMalloc((void**)&d_out, sizeof(float)*size_out);
//const int value = 5;
//cudaMemset(d_in, value, sizeof(float)*size);
float * h_in = (float *)malloc(size*sizeof(float));
for (int i = 0; i < size; i++) h_in[i] = 1.0f;
cudaMemcpy(d_in, h_in, sizeof(float)*size, cudaMemcpyHostToDevice);
// Running kernel wrapper
reduce(d_out, d_in, size, num_threads);
float result;
cudaMemcpy(&result, d_out, sizeof(float), cudaMemcpyDeviceToHost);
printf("sum is element is: %.f\n", result);
}