注意：我使用的是GT 740，每个SM有2个SM和192个CUDA核心。
我有一个可运行的CUDA内核，它被执行了4次：
__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);
我希望两个活动线程块都能访问d_mem的不同部分，尽管特定的部分并不重要，因为块之间不共享数据。
例如，以下是几种理想的访问模式之一：
- 区块1:
d_mem[0]和区块2:d_mem[1] - 块3:
d_mem[0]和块4:d_mem[1]
虽然这是不可取的：
- 区块1:
d_mem[0]和区块2:d_mem[0] - 块3:
d_mem[1]和块4:d_mem[1]
本质上，我想要一种寻址d_mem的方法，以便任何活动块的组合都可以访问它的不同部分。
我认为用块的SM ID寻址d_mem可能会起作用，但似乎不能保证这个ID在块的整个生命周期中保持不变。
我还考虑过用线程的全局ID模2(threadIdx.x + blockIdx.x * blockDim.x) % 2来寻址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个SMs*192(个以上的计算处于活动状态，从而遇到内存问题，试着想想，你是否可以为同一个工作包使用多个线程(threadIdx.x+i的值(，线程是否可以在warp边界内轻松协作或使用共享内存。有时，只为内核的读取和写入全局内存的部分使用更多的线程是有益的。这里出现了最大的延迟。
所以你可以把你的内核称为
foo<<<2, dim3(4, 192)>>>
并且具有4个螺纹而不是1个。对于图形，这4个可以是rgba通道或xyz坐标或三角形角。它们还可以在整个内核中更改用途。
作为性能优化，这使得一些计算更加复杂。
您对带有一个块的当前实现的if语句可能应该是
if(threadIdx.x + i < 1500)