c++ - CUDA——简单的代码，但我的一些 warp 不运行

Question

编辑:当我自己阅读这个问题时，我想通了。

问题的根源很可能是我没有分配足够的内存。我会尝试考虑一下并正确地做，然后回答我的问题。傻我。 :-[ 它并没有解释 warp 没有出现在 stdout 中......

原始问题

我在 CUDA 中创建了一个模板化内核，在其中迭代全局内存中的灰度图像数据部分(共享内存优化是在我开始工作时进行的)以实现具有圆盘形结构元素的形态学操作。每个线程对应于图像的一个像素。当数据类型为 char 时，一切都按预期工作，我的所有线程都在做它们应该做的事情。当我将它更改为 unsigned short 时，它开始运行并且只计算图像的上半部分。当我输入一些 printfs(我的设备有 2.0 CC)时，我发现一些应该运行的扭曲甚至没有被计算。

这是相关代码。

在我的 main.cpp 中，我调用了 gcuda::ErodeGpuGray8(img, radius); 和 gcuda::ErodeGpuGray16(img, radius); 以下函数:

// gcuda.h
…
i3d::Image3d<i3d::GRAY8> ErodeGpuGray8(i3d::Image3d<i3d::GRAY8> img, const unsigned int radius);
i3d::Image3d<i3d::GRAY16> ErodeGpuGray16(i3d::Image3d<i3d::GRAY16> img, const unsigned int radius);
…

// gcuda.cu
…
// call this from outside
Image3d<GRAY8> ErodeGpuGray8(Image3d<GRAY8> img, const unsigned int radius) {
    return ErodeGpu<GRAY8>(img, radius);
}

// call this from outside
Image3d<GRAY16> ErodeGpuGray16(Image3d<GRAY16> img, const unsigned int radius) {
    return ErodeGpu<GRAY16>(img, radius);
}
…

我使用的库将 GRAY8 定义为 char 并将 GRAY16 定义为 unsigned short。

这是我调用内核的方式(blockSize 是一个在相关命名空间中设置为 128 的 const int):

// gcuda.cu
template<typename T> Image3d<T> ErodeGpu(Image3d<T> img, const unsigned int radius) {
    unsigned int width = img.GetWidth();
    unsigned int height = img.GetHeight();
    unsigned int w = nextHighestPower2(width);
    unsigned int h = nextHighestPower2(height);
    const size_t n = width * height;
    const size_t N = w * h;

    Image3d<T>* rslt = new Image3d<T>(img);
    T *vx = rslt->GetFirstVoxelAddr();

    // kernel parameters
    dim3 dimBlock( blockSize );
    dim3 dimGrid( ceil( N / (float)blockSize) );

    // source voxel array on device (orig)
    T *vx_d;

    // result voxel array on device (for result of erosion)
    T *vxr1_d;

    // allocate memory on device
    gpuErrchk( cudaMalloc( (void**)&vx_d, n ) );
    gpuErrchk( cudaMemcpy( vx_d, vx, n, cudaMemcpyHostToDevice ) );

    gpuErrchk( cudaMalloc( (void**)&vxr1_d, n ) );
    gpuErrchk( cudaMemcpy( vxr1_d, vx_d, n, cudaMemcpyDeviceToDevice ) );

    ErodeGpu<T><<<dimGrid, dimBlock>>>(vx_d, vxr1_d, n, width, radius);

    gpuErrchk( cudaMemcpy( vx, vxr1_d, n, cudaMemcpyDeviceToHost ) );

    // free device memory
    gpuErrchk( cudaFree( vx_d ) );
    gpuErrchk( cudaFree( vxr1_d ) );

    // for debug purposes
    rslt->SaveImage("../erodegpu.png");
    return rslt;
}

我的测试图像的尺寸是 82x82，所以 n = 82*82 = 6724 和 N = 128*128 = 16384。

这是我的内核:

// gcuda.cu
// CUDA Kernel -- used for image erosion with a circular structure element of radius "erosionR"
template<typename T> __global__ void ErodeGpu(const T *in, T *out, const unsigned int n, const int width, const int erosionR)
{
    ErodeOrDilateCore<T>(ERODE, in, out, n, width, erosionR);
}

// The core of erosion or dilation. Operation is determined by the first parameter
template<typename T> __device__ void ErodeOrDilateCore(operation_t operation, const T *in, T *out, const unsigned int n, const int width, const int radius) {
    // get thread number, this method is overkill for my purposes but generally should be bulletproof, right?
    int blockId = blockIdx.x + blockIdx.y * gridDim.x + gridDim.x * gridDim.y * blockIdx.z;
    int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z) + (threadIdx.z * (blockDim.x * blockDim.y)) + (threadIdx.y * blockDim.x) + threadIdx.x;
    int tx = threadId;

    if (tx >= n) {
        printf("[%d > %d]", tx, n);
        return;
    } else {
        printf("{%d}", tx);
    }

    … (erosion implementation, stdout is the same when this is commented out so it's probably not the root of the problem)
}

据我了解，此代码应将一组随机排序的 [X > N] 和 {X} 字符串写入标准输出，其中 X = 线程 ID 并且应该是 n 大括号数字(即索引为 <n 的线程的输出)和其余的 N - n，但是当我运行它并使用正则表达式计算大括号内的数字，我发现我只得到其中的 256 个。此外，它们似乎出现在 32 个成员的组中，这告诉我有些 warp 在运行，有些没有。

我真的很困惑。当我不注释掉侵 eclipse 实现部分时，GRAY8 侵 eclipse 起作用而 GRAY16 侵 eclipse 不起作用，这无济于事，即使两种情况下的 stdout 输出完全相同(可能取决于输入，我只用 2 张图片试过)。

我错过了什么？这可能是什么原因造成的？我是否存在一些内存管理错误，或者某些扭曲不运行并且侵 eclipse 的东西可能只是图像库中的错误，只发生在 GRAY16 类型上？

Answer 1

所以这只是一个愚蠢的 malloc 错误。

代替

const size_t n = width * height;
const size_t N = w * h;

我用过

const int n = width * height;
const int N = w * h;

而不是错误的

gpuErrchk( cudaMalloc( (void**)&vx_d, n ) );
gpuErrchk( cudaMemcpy( vx_d, vx, n, cudaMemcpyHostToDevice ) );

gpuErrchk( cudaMalloc( (void**)&vxr1_d, n ) );
gpuErrchk( cudaMemcpy( vxr1_d, vx_d, n, cudaMemcpyDeviceToDevice ) );

…

gpuErrchk( cudaMemcpy( vx, vxr1_d, n, cudaMemcpyDeviceToHost ) );

我用过

gpuErrchk( cudaMalloc( (void**)&vx_d, n * sizeof(T) ) );
gpuErrchk( cudaMemcpy( vx_d, vx, n * sizeof(T), cudaMemcpyHostToDevice ) );

gpuErrchk( cudaMalloc( (void**)&vxr1_d, n * sizeof(T) ) );
gpuErrchk( cudaMemcpy( vxr1_d, vx_d, n * sizeof(T), cudaMemcpyDeviceToDevice ) );

…

gpuErrchk( cudaMemcpy( vx, vxr1_d, n * sizeof(T), cudaMemcpyDeviceToHost ) );

现在侵 eclipse 工作正常，这是我试图解决的主要问题。不过，我仍然没有得到我期望的 stdout 输出，所以如果有人可以阐明这一点，请这样做。