使用GPU计算和存储最大值的最佳解决方案是什么?我现在的不满意

Question

我的设备上运行以下内核：

__global__ void kernel1(float *Vb, int *sS, int *sE, int *bI, float *eR, int S, int K, int B, int N)
{
const unsigned long long int blockId = blockIdx.x //1D
+ blockIdx.y * gridDim.x //2D
+ gridDim.x * gridDim.y * blockIdx.z; //3D
const unsigned long long int threadId = blockId * blockDim.x + threadIdx.x;
int s = threadId / (K * B), k = (threadId - (s * K * B)) / B, b = threadId - (s * K * B) - (k * B);
if (s < S && k < K && b < B)
{
float sum = 0;
for (int t = sS[k]; t <= sE[k]; t++)
sum += eR[s * N + bI[b * N + t]];
if (sum > Vb[b * K + k])
{
Vb[b * K + k] = sum;
}
}
}

我基本上是基于eR[SxN]和bI[BxN]矩阵(被映射为简单的1D阵列)以及sE[K]和sS[K]阵列来计算一些和，并试图将每个(s，(K，b))对获得的最大值存储在Vb[BxK]矩阵(也被映射为1D数组)中。

我的问题是，最终，Vb矩阵不包含为每对计算的最大值。据我所知，问题的出现是因为所有GPU线程都是并行运行的(当然，这是一件好事)，并且它们都同时达到"if(sum>Vb[b*K+K])"语句，因此所有线程都基于Vb[b*K+K+元素的原始值来评估它。因此，存储在Vb[b*K+K]中的最终值是在设置元素值的最后一个线程中计算的和的值(最后一个和大于原始元素值)，而不是总体最大值。

为了纠正这一点，我尝试将Vb转换为[SxKxB]立方体，以便计算所有(s，k，b)对的和，然后最大化CPU上每个s的元素。内核如下所示：

__global__ void kernel2(float *Vb, int *sS, int *sE, int *bI, float *eR, int S, int K, int B, int N)
{
const unsigned long long int blockId = blockIdx.x //1D
+ blockIdx.y * gridDim.x //2D
+ gridDim.x * gridDim.y * blockIdx.z; //3D
const unsigned long long int threadId = blockId * blockDim.x + threadIdx.x;
int s = threadId / (K * B), k = (threadId - (s * K * B)) / B, b = threadId - (s * K * B) - (k * B);
if (s < S && k < K && b < B)
{
float sum = 0;
for (int t = sS[k]; t <= sE[k]; t++)
sum += eR[s * N + bI[b * N + t]];
Vb[s * K * B + k * B + b] = sum;
}
}

这对于相对较小的S、K和B来说很好，但当它们较大时(比如S＝100000，K＝12，B＝1000)，Vb矩阵的内存需求(约4.5GB)远远超过设备空闲内存(约600-700MB)。

所以我的问题是：1.有没有办法让第一个内核按预期工作(最终获得最大和)？2.当处理大数据集时，你认为解决这个问题的最佳方法是什么？a.将数据拆分为多个块，并运行kernel2的多个实例？(我认为这大大增加了计算所需的时间)b.投资于内存容量更大的硬件？c.我读到有可能直接使用设备中的主机内存(零内存拷贝)，但我现在不熟悉它的工作方式。这可能是一个解决方案吗？(这样我就可以专注于学习和实施它)d.另一种方法(请建议)。。。越简单越好。

对第一个问题采取积极有效的解决办法将是非常可取的。

我的设备是GeForce GT 220，总内存为1GB，计算能力为1.2(最新驱动程序)。我在Windows 8.1 64位上的VS2012中使用CUDA5.5。

Answer 1

您可以实现并使用atomicMax()的浮点版本，但性能可能不太好，尤其是在CC 1.2设备上。也许值得一试。

借款来源https://stackoverflow.com/a/17401122/442006:

__device__ static float atomicMax(float* address, float val)
{
int* address_as_i = (int*) address;
int old = *address_as_i, assumed;
do {
assumed = old;
old = ::atomicCAS(address_as_i, assumed,
__float_as_int(::fmaxf(val, __int_as_float(assumed))));
} while (assumed != old);
return __int_as_float(old);
}

然后：

__global__ void kernel1(float *Vb, int *sS, int *sE, int *bI, float *eR, int S, int K, int B, int N)
{
const unsigned long long int blockId = blockIdx.x //1D
+ blockIdx.y * gridDim.x //2D
+ gridDim.x * gridDim.y * blockIdx.z; //3D
const unsigned long long int threadId = blockId * blockDim.x + threadIdx.x;
int s = threadId / (K * B), k = (threadId - (s * K * B)) / B, b = threadId - (s * K * B) - (k * B);
if (s < S && k < K && b < B)
{
float sum = 0;
for (int t = sS[k]; t <= sE[k]; t++)
sum += eR[s * N + bI[b * N + t]];
atomicMax(Vb + b * K + k, sum);
}
}