一个二维矩阵的列求和问题

[richardingding]

RT：

1 2 3 4 5 6 7 8 9
1 2 3 4 5 6 7 8 9
1 2 3 4 5 6 7 8 9
1 2 3 4 5 6 7 8 9
1 2 3 4 5 6 7 8 9 ===》 9 18 27 36 45 54 63 72 81
1 2 3 4 5 6 7 8 9
1 2 3 4 5 6 7 8 9
1 2 3 4 5 6 7 8 9
1 2 3 4 5 6 7 8 9

例如这个矩阵，我想把每一列进行求和，请问这个kernel改怎么写呢？

[richardingding]

[Quote=引用 5 楼 sfft 的回复:]

重新dim3，把grid,block都dim成一维的，用开勇写的kernel: testAdd
[/Quote]
请问在kernel里面可以重新定义dim3么？能够说稍微具体一点么？在哪里个地方重新定义dim3？

[richardingding]

[Quote=引用 6 楼 sfft 的回复:]

l7331014 说的很对，因为你按列操作，最好用smem，否则效率比较低。
[/Quote]
开勇的方法用到的global memory，是coalesce的访问，share memory容易出现bank conflict。如果单纯计算的话，开勇的kernel效率高一些。

[sfft]

kernel外面啦

   dim3 dimBlock(BLOCK_SIZE);    
    dim3 dimGrid(size+BLOCK_SIZE-1 / dimBlock.x); 
    testAdd<<<dimGrid,dimBlock>>>( d_data, width, depth,  d_ret);

[无心人_过过小日子]

每列一个thread.
程序构架如下:
读M行到smem.
线程同步.
累计列.

[OpenHero]

__global__ void testAdd(float* d_data, int width, int depth, float* d_ret)
{
	float ret = 0; 
	int index = blockDim.x*blockDim.x + threadIdx.x;

	if (index < width)
	{
		d_data += index;
		d_ret += index;
		for (int i =0; i < depth; i++)
		{
			ret += d_data[0];
			d_data += width;
		}

		d_ret[0] = ret;
	}
}

这样写会比用shared memory高效的:)
没测试过，自己测试

[richardingding]

[Quote=引用 2 楼 openhero 的回复:]
C/C++ code

__global__ void testAdd(float* d_data, int width, int depth, float* d_ret)
{
float ret = 0;
int index = blockDim.x*blockDim.x + threadIdx.x;

if (index < width)
{
……
[/Quote]
谢谢开勇，我能够理解这样的kernel。对不起，是我没有说明白。
求数组中每一列的和是我kernel函数的最后一步需要做的。在做这个求和操作之前，我使用的数据结构是
int x = threadIdx.x + blockIdx.x*BLOCK_SIZE_X;
int y = threadIdx.y + blockIdx.y*BLOCK_SIZE_Y;
用来定位index用来计算。可是使用了这样的数据结构后，在对于每一列求和时候遇到了问题，我不知道如何设计这个kernel的后面的求和计算。希望多多指教。

[richardingding]

[Quote=引用 1 楼 l7331014 的回复:]

每列一个thread.
程序构架如下:
读M行到smem.
线程同步.
累计列.
[/Quote]
谢谢你，我的数据结构见上面的说明，谢谢你的回帖，也希望多交流:)

[sfft]

重新dim3，把grid,block都dim成一维的，用开勇写的kernel: testAdd

[sfft]

l7331014 说的很对，因为你按列操作，最好用smem，否则效率比较低。

[fantacy1127]

[Quote=引用 2 楼 openhero 的回复:]
C/C++ code

__global__ void testAdd(float* d_data, int width, int depth, float* d_ret)
{
float ret = 0;
int index = blockDim.x*blockDim.x + threadIdx.x;

if (index < width)
{
……
[/Quote]
我是初学者，有些地方不是很清楚。我想请问下int index = blockDim.x*blockDim.x + threadIdx.x;
这句话是做什么用？我觉得int index = blockDim.x*blockIdx.x + threadIdx.x;好像也是对的。。。希望大家能指点下·