如何为八个维度设置我的 Cuda 线程块和网格？答案

【问题标题】：How do I setup my Cuda thread blocks and grids for eight dimensions?如何为八个维度设置我的 Cuda 线程块和网格？
【发布时间】：2013-06-17 09:15:22
【问题描述】：

我定期进行三维 cuda 工作，但现在我遇到了八个维度的问题。

如何为八个维度设置线程块和网格？

例如在 3d 中我使用：

    grid.x=int(321);  
    grid.y=int(321);

    foo<<<grid,int(321)>>>

所以每个线程都有自己的 x,y,z 地址，并尝试仅使用该 x,y,z 来解决问题。如果结果是积极的结果得到报告。我知道如何传递内存指针来报告结果。

现在我有一个新问题，我想在八个不同的轴上尝试 0 到 11 的值，而不是在三个轴上尝试 0 到 321。我假设每个线程都应该通过线程块网格获取它的一组 (a,b,c,d,e,f,g,h) 参数。

如何通过网格和块大小将八个维度映射到我的内核中？

【问题讨论】：

认为我已经回答了自己的问题。我将块设置为 32,32,1，我的网格设置为 32,32,32，并通过代码在内部执行其他三个参数。这听起来对吗？
你是认真的吗？您的“8 个不同轴 [sic] 上的 0 到 11 的值”将需要 12**8（>4.29 亿）个线程。

标签： cuda dimension

【解决方案1】：

网格的 2D 特性和块的 3D 特性只是 NVIDIA 的一种便利；他们还不如在那里只接受单个整数，并且硬件将以相同的方式工作。这就是为什么，如果您的问题本质上不是 2D 或 3D，我会建议使用单维索引并在需要时“拆分”索引。像这样的：

int grid = 65536;
int block = 256;
foo<<<grid,block>>>();

然后在您的设备代码中：

__device__ int globalIndex() { blockIdx.x * blockDim.x + threadIdx.x; }

__device__ int index8D(int dim) { return (globalIndex() >> (dim*3))%8 }

在上面的例子中，通过调用index8D(i)，你得到了当前线程对应的i-th坐标。但是，有效坐标仅在 0 到 7 的范围内。您将需要更多线程来增加该范围...

请注意！八维立方体，每个维度中只有很少的单元格是巨大的！如果您的 8D 空间仅在 [0..7] 上进行迭代，我们正在查看 8^8 单元格总数 (16777216)。您可能需要考虑使用一个线程来实际迭代多个单元格。

为了解释index8D 的内部结构：我本质上是将全局线程索引的二进制表示分成每维 3 位的组：

101 110 001 000 110 001 101 110

每个组现在代表一个维度中的一个索引。左移和取模用于提取相应的 3 位组。（计算模常数将由编译器优化为按位运算；为了便于阅读，我保留它）

【讨论】：

谢谢 CygnusX1，我现在看到了我需要知道的东西。 talonmies，已经在我的 I7 上进行了八轴搜索，范围为 0 到 13。花了一周的时间，但它做到了。当然，我的 cuda 卡可以击败我的 I7。它是纯数学，没有线程之间的通信。
@user2475221：我不是质疑进行 4.21 亿组计算的必要性。我质疑是否需要使用 4.21 亿个线程来进行这些计算。您也可能不需要在单个内核启动中进行计算，因此可能也不需要为所有八个维度制定完整的索引方案
我有一个与 OP 类似的问题，我不明白为什么将单个线程集成到多个单元格上而不是仅仅拥有许多线程会更好？

【解决方案2】：

这是我对我的问题的回答，这是一个简短的 Cuda 程序，可以在八轴上找到无单位方程。第二行的搜索大小应该是一个大于 5 的奇数。程序以行的形式输出 c0 h0 A0 R0 W0 G0 N0 F0 的幂。

我可能错了，但我认为该结构是老派的短路逻辑结构线程发散并报告正确答案；否则线程只会收敛并与错误的线程全速运行！

这是我正在求解的常数幂的无单位方程。我认为程序很粗糙，因为方程很粗糙。每个括号组必须归零才能使等式成立。可以像我的方法一样使用正方形或 abs 或纯逻辑。

(x4-x5)^2+(x3+x8)^2+(x4+x7+x8)^2+(x2+x4-x6)^2+(-2*(x4+x6)+x8 -x1-x2)^2+(2*(x2+x4)+x1+x5+3*x6)^2 = 0

如果有人可以帮助我让它运行得更快，我会很感兴趣...

#include "stdio.h"
#define searchsize 27
//nvcc helloworld.cu -o helloworld -arch=sm_21 -maxrregcount=20 -ccbin=gcc-4.4
__global__ void helloworld()
{
int x1,x2,x3,x4,x5,x6,x7,x8,rlow,rhgh;

rlow=-((gridDim.x-1)/2);
rhgh=((gridDim.x-1)/2);
x1=blockIdx.x+rlow;
x2=blockIdx.y+rlow;
x3=threadIdx.x+rlow;
x4=rlow;
x5=rlow;
x6=rlow;
x7=rlow;
x8=rlow;

while (x8<=rhgh)
{
if (x4 == x5)
{
if (x3 == -x8)
{
if (x4 + x7 == -x8)
{
if (x2+x4 == x6)
{
if (-2*( x4 + x6) + x8 == x1 + x2)
{
if (2*(x2+x4) + x1 + x5 == -3*x6)
{
printf("%+4d,%+4d,%+4d,%+4d,%+4d,%+4d,%+4d,%+4d \n", x1,x2,x3,x4,x5,x6,x7,x8);
}
}
}
}
}
}
x4=x4+1;
if (x4>rhgh)
{
x5=x5+1;
x4=rlow;
}
if (x5>rhgh)
{
x6=x6+1;
x5=rlow;
}
if (x6>rhgh)
{
x7=x7+1;
x6=rlow;
}
if (x7>rhgh)
{
x8=x8+1;
x7=rlow;
}
}
}

int main()
{
int rangeofsearch(searchsize);
dim3 grid,block;
grid.x=rangeofsearch;
grid.y=rangeofsearch;
block.x=rangeofsearch;
size_t buf=1e7;

cudaDeviceSetLimit(cudaLimitPrintfFifoSize, buf);
helloworld<<<grid,block>>>();
cudaDeviceSynchronize();
return 0;
}

Sample Output, powers in row form.

c0, h0, A0, R0, W0, G0, N0, F0

-14, -14, +0, +14, +14, +0, -14, +0
-13, -13, +0, +13, +13, +0, -13, +0
-12, -12, +0, +12, +12, +0, -12, +0
-11, -11, +0, +11, +11, +0, -11, +0
-7, -13, -2, +12, +12, -1, -14, +2
-6, -12, -2, +11, +11, -1, -13, +2
-5, -11, -2, +10, +10, -1, -12, +2
+0, -12, -4, +10, +10, -2, -14, +4
+1, -11, -4, +9, +9, -2, -13, +4
+7, -11, -6, +8, +8, -3, -14, +6
-14, -8, +2, +9, +9, +1, -7, -2
-13, -7, +2, +8, +8, +1, -6, -2
-12, -6, +2, +7, +7, +1, -5, -2
-11, -5, +2, +6, +6, +1, -4, -2 
...

【讨论】：