Mathematica 中的 CUDAFunctionLoad - 索引问题

Question

我正在尝试调试我在 CUDA 机器上遇到的索引问题

Cuda Machine Info:

{1->{Name->Tesla C2050,Clock Rate->1147000,Compute Capabilities->2.,GPU Overlap->1,Maximum Block Dimensions->{1024,1024,64},Maximum Grid Dimensions- >{65535,65535,65535},Maximum Threads Per Block->1024,Maximum Shared Memory Per Block->49152,Total Constant Memory->65536,Warp Size->32,Maximum Pitch->2147483647,Maximum Registers Per Block- >32768,Texture Alignment->512,Multiprocessor Count->14,Core Count->448,Execution Timeout->0,Integrated->False,Can Map Host Memory->True,Compute Mode->Default,Texture1D Width-> 65536,Texture2D 宽度->65536,Texture2D 高度->65535,Texture3D 宽度->2048,Texture3D 高度->2048,Texture3D 深度->2048,Texture2D 阵列宽度->16384,Texture2D 阵列高度->16384,Texture2D 阵列切片- >2048,Surface Alignment->512,Concurrent Kernels->True,ECC Enabled->True,Total Memory->2817982462},

此代码所做的只是将 3D 数组的值设置为等于 CUDA 正在使用的索引：

__global __ void cudaMatExp(
float *matrix1, float *matrixStore, int lengthx, int lengthy, int lengthz){

long UniqueBlockIndex = blockIdx.y * gridDim.x + blockIdx.x;

long index = UniqueBlockIndex * blockDim.z * blockDim.y * blockDim.x +
    threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x +
    threadIdx.x;

if (index < lengthx*lengthy*lengthz) {

matrixStore[index] =  index;

}
}

由于某种原因，一旦我的 3D 数组的维度变得太大，索引就会停止。

我尝试了不同的块尺寸（blockDim.x by blockDim.y by blockDim.z）：

8x8x8 只提供正确的索引到数组维度 12x12x12

9x9x9 只提供正确的索引到数组维度 14x14x14

10x10x10 仅提供正确的索引，最大数组维度为 15x15x15

对于大于这些的维度，所有不同的块大小最终都会再次开始增加，但它们永远不会达到 dim^3-1 的值（这是 cuda 线程应该达到的最大索引）

这里有一些图表可以说明这种行为：

例如：这是在 x 轴上绘制 3D 数组的维度（即 xxx），在 y 轴上绘制 cuda 执行期间处理的最大索引号.此特定图适用于 10x10x10 的块尺寸。

这是生成该图的 (Mathematica) 代码，但是当我运行这个时，我使用了 1024x1x1 的块尺寸：

CUDAExp = CUDAFunctionLoad[codeexp, "cudaMatExp",
  {{"Float", _,"Input"}, {"Float", _,"Output"},
    _Integer, _Integer, _Integer},
  {1024, 1, 1}]; (*These last three numbers are the block dimensions*)

max = 100; (* the maximum dimension of the 3D array *)
hold = Table[1, {i, 1, max}];
compare = Table[i^3, {i, 1, max}];
Do[
   dim = ii;
   AA  = CUDAMemoryLoad[ConstantArray[1.0, {dim, dim, dim}], Real, 
                                     "TargetPrecision" -> "Single"];
   BB  = CUDAMemoryLoad[ConstantArray[1.0, {dim, dim, dim}], Real, 
                                     "TargetPrecision" -> "Single"];

   hold[[ii]] = Max[Flatten[
                  CUDAMemoryGet[CUDAExp[AA, BB, dim, dim, dim][[1]]]]];

 , {ii, 1, max}]

ListLinePlot[{compare, Flatten[hold]}, PlotRange -> All]

这是同一个图，但现在绘制 x^3 以比较它应该在的位置。请注意，它在数组的维度>32

后发散

我测试了 3D 数组的维度，并查看索引的范围，并将其与 dim^3-1 进行比较。例如。对于 dim=32，cuda 最大索引为 32767（即 32^3 -1），但对于 dim=33，cuda 输出为 33791，而应为 35936（33^3 -1）。注意 33791-32767 = 1024 = blockDim.x

问题：

有没有办法正确索引一个维度大于 Mathematica 中块维度的数组？

现在，我知道有些人在他们的索引方程中使用 __mul24(threadIdx.y,blockDim.x) 来防止位乘法错误，但在我的情况下似乎没有帮助。

另外，我看到有人提到您应该使用 -arch=sm_11 编译代码，因为默认情况下它是为计算能力 1.0 编译的。不过，我不知道 Mathematica 是否是这种情况。我假设 CUDAFunctionLoad[] 知道使用 2.0 功能进行编译。有谁知道吗？

任何建议都会非常有帮助！

Answer 1

因此，Mathematica 有一种处理网格尺寸的隐藏方式，要将网格尺寸固定为可行的东西，您必须在调用的函数末尾添加另一个数字。

参数表示要启动的线程数（或网格维度乘以块维度）。

例如，在我上面的代码中：

CUDAExp = 
  CUDAFunctionLoad[codeexp, 
   "cudaMatExp", {
           {"Float", _, "Input"}, {"Float", _,"Output"}, 
                        _Integer, _Integer, _Integer}, 
     {8, 8, 8}, "ShellOutputFunction" -> Print];

(8,8,8) 表示块的维度。

当您在mathematica 中调用CUDAExp[] 时，您可以添加一个参数来表示要启动的线程数：

在这个例子中，我终于让它与以下内容一起工作：

// AA and BB are 3D arrays of 0 with dimensions dim^3
dim = 64;
CUDAExp[AA, BB, dim, dim, dim, 4089];

请注意，当您使用 CUDAFunctionLoad[] 进行编译时，它只需要 5 个输入，第一个是您传递给它的数组（维度为 dim x dim x dim），第二个是存储它的内存的位置。第三、四、五是维度。

当您传递第 6 个时，mathematica 将其转换为 gridDim.x * blockDim.x，因此，由于我知道我需要 gridDim.x = 512 才能处理数组中的每个元素，所以我将此数字设置为等于 512 * 8 = 4089。

我希望这对将来遇到此问题的人来说清楚且有用。