如何解决CUDA在执行过程中组合线程独立??变量
| 伙计们,如果标题令人困惑,我深表歉意。我虽然费时又辛苦,却无法想出正确的方式在一行中表达问题。因此,这里有更多详细信息。我正在做一个基本的图像减法,其中第二幅图像已被修改,我需要找到对该图像进行了多少更改的比率。为此,我使用了以下代码。两张图片均为128x1024。for(int i = 0; i < 128; i++)
{
for(int j = 0; j < 1024; j++)
{
den++;
diff[i * 1024 + j] = orig[i * 1024 + j] - modified[i * 1024 + j];
if(diff[i * 1024 + j] < error)
{
num++;
}
}
}
ratio = num/den;
上面的代码在CPU上工作正常,但我想尝试在CUDA上执行此操作。为此,我可以设置CUDA进行图像的基本减法(下面的代码),但是我不知道如何进行条件if语句来得出比率。
__global__ void calcRatio(float *orig,float *modified,int size,float *result)
{
int index = threadIdx.x + blockIdx.x * blockDim.x;
if(index < size)
result[index] = orig[index] - modified[index];
}
因此,到目前为止,它仍然有效,但是我无法弄清楚如何在每个线程执行结束时对每个线程中的num和den计数器进行并行计算来计算比率。在我看来,num和den counders独立于线程,因为每次我尝试使用它们时,似乎它们只会增加一次。
当我刚开始使用CUDA时,任何帮助将不胜感激,而我在网上看到的每个示例似乎都不适用于我需要做的事情。
编辑:修复了我的天真代码。忘记键入代码中的主要条件之一。这是漫长的一天。
for(int i = 0; i < 128; i++)
{
for(int j = 0; j < 1024; j++)
{
if(modified[i * 1024 + j] < 400.0) //400.0 threshold value to ignore noise
{
den++;
diff[i * 1024 + j] = orig[i * 1024 + j] - modified[i * 1024 + j];
if(diff[i * 1024 + j] < error)
{
num++;
}
}
}
}
ratio = num/den;
解决方法
在所有线程上执行全局求和所需的操作称为“并行归约”。尽管您可以使用原子操作来执行此操作,但我不建议这样做。有一个还原内核和一篇很好的文章,讨论了CUDA SDK中的技术,值得一读。
如果我正在编写代码以执行所需的操作,则可能看起来像这样:
template <int blocksize>
__global__ void calcRatio(float *orig,float *modified,int size,float *result,int *count,const float error)
{
__shared__ volatile float buff[blocksize];
int index = threadIdx.x + blockIdx.x * blockDim.x;
int stride = blockDim.x * gridDim.x;
int count = 0;
for(int i=index; i<n; i+=stride) {
val = orig[index] - modified[index];
count += (val < error);
result[index] = val;
}
buff[threadIdx.x] = count;
__syncthreads();
// Parallel reduction in shared memory using 1 warp
if (threadId.x < warpSize) {
for(int i=threadIdx.x + warpSize; i<blocksize; i+= warpSize) {
buff[threadIdx.x] += buff[i];
if (threadIdx.x < 16) buff[threadIdx.x] +=buff[threadIdx.x + 16];
if (threadIdx.x < 8) buff[threadIdx.x] +=buff[threadIdx.x + 8];
if (threadIdx.x < 4) buff[threadIdx.x] +=buff[threadIdx.x + 4];
if (threadIdx.x < 2) buff[threadIdx.x] +=buff[threadIdx.x + 2];
if (threadIdx.x == 0) count[blockIdx.x] = buff[0] + buff[1];
}
}
第一个节执行您的串行代码所执行的操作-计算差和小于错误的元素的线程局部总数。请注意,我已经编写了此版本,以便每个线程都被设计为处理输入数据的多个条目。这样做是为了帮助抵消随后并行缩减的计算成本,并且其想法是,与输入数据集条目相比,您将使用更少的块和线程。
第二个节是归约本身,它是在共享内存中完成的。它实际上是一种“树状”操作,其中将单个线程块中的线程局部小计集的大小首先求和为32个小计,然后将这些小计合并,直到有该块的最终小计为止,然后存储的是该块的总数。最后,您会看到一小部分总数的清单,每个启动的块一个,可以复制回主机,并在此处计算最终结果。
请注意,我在浏览器中对此代码进行了编码,但尚未编译,可能会出现错误,但是它应该使您了解您尝试执行的“高级”版本如何工作。
,分母非常简单,因为它只是大小。
分子比较麻烦,因为给定线程的值取决于所有先前的值。您将不得不顺序执行该操作。
您正在寻找的东西可能是atomicAdd。不过,它非常慢。
我认为您会发现此问题相关。您的num基本上是全局数据。
CUDA阵列到阵列总和
或者,您可以将错误检查的结果转储到数组中。然后可以对结果计数进行并行化。这有点棘手,但我认为类似的事情会扩大:http://tekpool.wordpress.com/2006/09/25/bit-count-parallel-counting-mit-hakmem/
版权声明:本文内容由互联网用户自发贡献,该文观点与技术仅代表作者本人。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如发现本站有涉嫌侵权/违法违规的内容, 请发送邮件至 dio@foxmail.com 举报,一经查实,本站将立刻删除。