0

CUDALink我目前使用Mathematica 提供的包装器在我的 GPU (GTX560Ti) 上运行一个简单的 3D 模板转换。块尺寸现在对我来说并不重要，因为我没有使用任何共享内存或寻找优化（现在）。

blockDim.x因此，我可以为和设置任何合理的数字blockDim.y。无论我设置什么维度，包装器都会启动适当数量的块，没问题。然而，在 z 维度中，只有一个块被启动。因此blockDim.z限制了我可以在该方向计算的总点数。

为什么z方向只有一个方块？我该如何解决这个问题？

作为参考，这是我正在使用的内核：

__global__ void conv(Real_t in[48][48][48], Real_t out[48][48][48], mint stencil[13][13][13], mint length, mint rad) {
    int x = threadIdx.x + blockIdx.x*blockDim.x;
    int y = threadIdx.y + blockIdx.y*blockDim.y;
    int z = threadIdx.z + blockIdx.z*blockDim.z;
    while (x<length||y<length||z<length) {
        out[x][y][z] = 0;
        for (int ix = -rad; ix <= rad; ix++) {
        for (int iy = -rad; iy <= rad; iy++) {
        for (int iz = -rad; iz <= rad; iz++) {
            if ( (fminf(x,fminf(y,z))-rad >= 0)
                && (fmaxf(x,fmaxf(y,z))+rad < length) )
                {out[x][y][z] += stencil[ix+rad][iy+rad][iz+rad]*in[ix+x][iy+y][iz+z];}
        }   }   }
        if (x<length) {
            x+= blockDim.x * gridDim.x;
        } else if (y<length) {
            y+= blockDim.y * gridDim.y;
        } else if (z<length) {
            z+= blockDim.z * gridDim.z;
        }
    }
}

请注意：变量length对应于数组的维度（例如 48）。rad与模板有关，小于length. stencil只是一个 0 和 1 的数组，用于从in我想要总结的东西中选择out.

我正在使用以下代码在 Mathematica 中运行内核：

Needs["CUDALink`"];
conv = CUDAFunctionLoad[code (*the kernel above, stored as a string*), "conv", {{_Real, _, "Input"}, {_Real, _, "Output"}, {_Integer , _, "Input"}, _Integer, _Integer}, {4, 4, 10}, "TargetPrecision" -> "Single", "XCompilerInstallation" -> "/usr/local/gcc44/bin/", "CleanIntermediate" -> False];
output = ConstantArray[1, {length, length, length}];
result =  conv[input, output, stencil, length, rad];

为了说明我的问题，这是我输出的一部分（显然我不能发布图像）：

0.  0.  0.  0.  0.  0.  0.  0.  0.  0.  1.  1.  1.
0.  0.  0.  0.  0.  0.  0.  0.  0.  0.  1.  1.  1.
0.  0.  0.  0.  0.  0.  0.  0.  0.  0.  1.  1.  1.
0.  0.  0.  0.  0.  0.  0.  0.  0.  0.  1.  1.  1.
0.  0.  0.  0.  0.  0.  0.  0.  0.  0.  1.  1.  1.
0.  0.  0.  0.  0.  0.  0.  0.  0.  0.  1.  1.  1.
0.  0.  0.  0.  0.  0.  0.  0.  0.  0.  1.  1.  1.
0.  0.  0.  0.  0.  0.  0.  0.  0.000578704 0.00173611  1.  1.  1.
0.  0.  0.  0.  0.  0.  0.  0.000289352 0.000868056 0.00173611  1.  1.  1.
0.  0.  0.  0.  0.  0.  0.  0.000578704 0.00144676  0.00260417  1.  1.  1.
0.  0.  0.  0.  0.  0.  0.  0.00115741  0.00202546  0.00347222  1.  1.  1.
0.  0.  0.  0.  0.  0.  0.  0.00115741  0.00202546  0.00347222  1.  1.  1.
0.  0.  0.  0.  0.  0.  0.  0.000578704 0.00144676  0.00289352  1.  1.  1.
0.  0.  0.  0.  0.  0.  0.  0.000578704 0.00144676  0.00289352  1.  1.  1.

这是用blockDim.z = 10. 零和分数是有用的值，但它们只是我初始化out数组的值。仅计算前 10 列，对应于 z 方向的单个块。blockDim.z（对于介于1和64（费米 GPU 的上限）之间的任何值，此行为都是可重现的。

4

1 回答 1

1

好的，我猜这种行为只是 CUDAResources 中的一个错误，而不是实际的编程问题。（不过，只有一个块。我现在拥有的是一种解决方法。）

我用 删除了 CUDAResources CUDAResourcesUninstall[]，重新启动了 Mathematica，使用重新安装CUDAResourcesInstall["/path/to/paclet/file",Update->True]并再次重新启动了 Mathematica。

然后我将内核更改为以下代码：

__global__ void conv(Real_t in[48][48][48], Real_t out[48][48][48], \
mint stencil[13][13][13], mint length, mint rad) {
    int x = threadIdx.x + blockIdx.x*blockDim.x;
    int y = threadIdx.y + blockIdx.y*blockDim.y;
    int z = threadIdx.z + blockIdx.z*blockDim.z;
    while (z<length) {
        out[x][y][z] = 0;
        for (int ix = -rad; ix <= rad; ix++) {
        for (int iy = -rad; iy <= rad; iy++) {
        for (int iz = -rad; iz <= rad; iz++) {
            if ( (fminf(x,fminf(y,z))-rad >= 0)
                && (fmaxf(x,fmaxf(y,z))+rad < length) )
                {out[x][y][z] += stencil[ix+rad][iy+rad][iz+rad]*in[ix+x][iy+y][iz+z];}
        }   }   }
        if (z<length) {
            z+= blockDim.z * gridDim.z;
        }
    }
}

现在它可以工作了。希望它保持这种状态。这当然意味着在 z 方向上进行的并行化较少，因为基本上有一个线程块在网格上按顺序运行，而不是多个块并行工作。但这很好，代码对于我的目的来说已经足够快了。

非常感谢所有帮助过的人。

于 2014-01-31T13:49:43.543 回答