请问这里报错如何修改。

Question

__global__ void gpuRecursiveReduce (int *g_idata, int *g_odata,
                                    unsigned int isize)
{
    // set thread ID
    unsigned int tid = threadIdx.x;

    // convert global data pointer to the local pointer of this block
    int *idata = g_idata + blockIdx.x * blockDim.x;
    int *odata = &g_odata[blockIdx.x];

    // stop condition
    if (isize == 2 && tid == 0)
    {
        g_odata[blockIdx.x] = idata[0] + idata[1];
        return;
    }

    // nested invocation
    int istride = isize >> 1;

    if(istride > 1 && tid < istride)
    {
        // in place reduction
        idata[tid] += idata[tid + istride];
    }

    // sync at block level
    __syncthreads();

    // nested invocation to generate child grids
    if(tid == 0)
    {
        gpuRecursiveReduce<<<1, istride>>>(idata, odata, istride);

        // sync all child grids launched in this block
        cudaDeviceSynchronize();
    }

    // sync at block level again
    __syncthreads();
}

// main from here
int main(int argc, char **argv)
{
    。。。。。
    gpuRecursiveReduce<<<grid, block>>>(d_idata, d_odata, block.x);
    。。。。。。
}

make
nvcc -O2 -arch=sm_86 -o nestedHelloWorld nestedHelloWorld.cu -lcudadevrt --relocatable-device-code true
nvcc -O2 -arch=sm_86 -o nestedReduce nestedReduce.cu -lcudadevrt --relocatable-device-code true
nestedReduce.cu(94): error: calling a host function("cudaDeviceSynchronize") from a global function("gpuRecursiveReduce") is not allowed

nestedReduce.cu(94): error: identifier "cudaDeviceSynchronize" is undefined in device code

2 errors detected in the compilation of "nestedReduce.cu".
make: *** [Makefile:8：nestedReduce] 错误 1

一本书里的示例代码，请问这里为啥有错？应该如何修改才行。gpuRecursiveReduce 既被host端调用，同时进入device里后，在device里递归调用，所以又被device调用，只能定义成__global__，还能怎么改呢？

Answer 1

参考通义千问和郭老师的小迷弟雅思莫了-编写提供，主要是为了帮助您快速解决问题，如下回答，您可以参考下，如有问题，评论区留言。

在 CUDA 中，cudaDeviceSynchronize() 是一个主机端函数，用于同步主机与设备间的执行。而在您的示例代码中，您尝试在一个设备端的 __global__ 函数（gpuRecursiveReduce）内部调用它，这是不允许的。

为了解决这个问题并实现设备端的同步，您可以移除对 cudaDeviceSynchronize() 的调用，并依赖于 __syncthreads() 来确保在同一块内的线程间完成同步。对于跨块的同步，CUDA 提供了动态协作组（Dynamic Parallelism），但请注意，这需要支持动态并行性的硬件架构（ Kepler 架构或更新的 GPU）。

针对您的代码示例，可以简单地删除 cudaDeviceSynchronize() 这一行：

// 删除这一行
// cudaDeviceSynchronize();

然而，由于这是一个递归并行化示例，要实现正确的跨块同步可能需要重新设计算法以适应没有 cudaDeviceSynchronize() 的情况。您可以通过使用事件（events）或者在主机端进行适当的 kernel 调用来间接实现同步。考虑到您目前的代码结构和设备限制，可能需要进一步调整整个递归逻辑才能正确同步不同层级的递归调用。