如何为每个活动线程块处理重用的全局内存的唯一部分？

如何解决如何为每个活动线程块处理重用的全局内存的唯一部分？

注意：我使用的是GT 740，每个SM具有2个SM和192个CUDA内核。

我有一个工作CUDA内核，该内核执行了4次：

from gekko import GEKKO

m = GEKKO()

# Variables
x1 = m.Var()
x2 = m.Var()

# Equation
FindX1 = 6 == x1*x2
m.Equation(FindX1)

# Fix x1 to 2
m.fix(x1,val=2)

# Solve X2
m.solve(disp=False)
print("X1: %s and X2: %s" % (x1.VALUE,x2.VALUE))

# Fix x2 to 2
m.fix(x2,val=2)

# Free x1
m.free(x1)

# Solve X1
m.solve(disp=False)
print("X1: %s and X2: %s" % (x1.VALUE,x2.VALUE))

由于内存限制，每个内核调用都会重用分配给__global__ void foo(float *d_a,int i) { if (i < 1500) { ... ... ... } } int main() { float *d_mem; cudaMalloc(&d_mem,lots_of_bytes); for (int i = 0; i < 1500; i += 384) foo<<<1,384>>>(d_mem,i); return 0; } 的内存。

我想将其修改为从单个语句执行，如下所示：

d_mem

我希望两个活动线程块都访问foo<<<8,192>>>(d_mem); 的不同部分，尽管特定的部分并不重要，因为在块之间不共享数据。

例如，以下是几种理想的访问模式之一：

1. 第1块：d_mem和第2块：d_mem[0]
2. 第3块：d_mem[1]和第4块：d_mem[0]
3. ...

这是不可取的：

1. 第1块：d_mem[1]和第2块：d_mem[0]
2. 第3块：d_mem[0]和第4块：d_mem[1]
3. ...

本质上，我想要一种寻址d_mem[1]的方法，以便活动块的任何组合都可以访问其中的不同部分。

我认为用块的SM ID寻址d_mem可能会起作用，but it appears that this ID is not guaranteed to remain the same throughout a block's life.

我还考虑过使用线程的全局ID模2 d_mem来寻址d_mem，但这依赖于以特定顺序处理的块。

这主要与每个SM使用1个块有关，但是我也对如何解决每个SM任意数目的块（如果有的话）感兴趣。

解决方法

最简单的方法是

foo<<<2,384>>> or foo<<<2,192>>>

并在内核中的计算周围放置一个for循环。然后，您可以使用blockIdx.x选择一半的内存。即使将两个块安排在同一SM上，也可以使用。此方法还适用于每个SM多个块，例如（每个块四分之一内存）

foo<<<4,96>>>

每个SM仅具有192个线程效率低下（至少384个，更好的是512、576、768、960或1024）。 SM需要隐藏等待时间并切换活动线程。如果通过同时进行超过384（= 2 SM * 192）个计算而陷入内存问题，请尝试考虑是否可以为同一工作包使用多个线程（threadIdx.x + i的值） ），线程可以轻松地在扭曲边界内协作或使用共享内存。有时，仅在内核部分（正在读写全局内存）使用更多线程是有益的。这里出现最大的延迟。

因此您可以将内核称为

foo<<<2,dim3(4,192)>>>

，有4个线程而不是一个。对于图形，这4个可能是rgba通道或xyz坐标或三角形角。他们还可以更改整个内核的用法。

作为性能优化，这会使一些计算更加复杂。

您当前的if语句可能只有一个块

if(threadIdx.x + i < 1500)