【发布时间】:2014-06-05 15:17:44
【问题描述】:
我正在用 CUDA 编写程序,问题如下:
两个矩阵 A (n * 128) 和 B (m * 128)
我取 A 的第一行,并逐一计算该向量与 B 的所有行之间的距离。
-
我将每个距离的结果写在矩阵C的一行上,所以C的元素C(i,j)包含A的i行和B的j行之间的距离。
-然后我继续下一行 A。
我是这样实现的: 我有一个由 ( n * m ) 个块组成的网格,每个块有 128 个线程。 ( 1 * 128 )。
程序正在编译,但问题是它没有给出好的距离。 我不知道哪里错了...
PS:我有 CUDA 6.0 和 NVIDIA GTX 650(计算能力 3.0)
__global__ void EuclidianDistances( float *A, float *B , float *C , int n , int m)
{
// SIZE is equal to 128
__shared__ float accumResult[SIZE];
__shared__ float sA[SIZE];
__shared__ float sB[SIZE];
// MAPPING
int bx = blockIdx.x; // n
int by = blockIdx.y; // m
int ty = threadIdx.y; // 128
int tx = threadIdx.x; // 1
sA[ty] = A [bx * SIZE + ty];
sB[ty] = B [by * SIZE + ty];
__syncthreads();
accumResult[ty] = (sA[ty] - sB[ty])*(sA[ty] - sB[ty]);
__syncthreads();
// Parallel tree-reduction
for (int stride = SIZE/2 ; stride < 0 ; stride >>= 1)
if (ty < stride)
{
accumResult[ty] += accumResult [stride + ty];
__syncthreads();
}
// Writing results to output matrix
if ((threadIdx.y == 0))
C [bx * m + by] = accumResult[ty];
__syncthreads();
}
【问题讨论】:
-
(ty < pas),这是什么pas?将__syncthreads();放在if语句中取决于threadIdx对我来说看起来很危险。 -
另外:条件似乎不对:
for (int stride = SIZE/2 ; stride < 0 ; stride >>= 1) -
@Levans:对不起,
pas是stride。我刚刚更正了。