Open lawlietSxk opened 3 months ago
不止这里,得是所有的blockIdx.x和blockIdx.y都交换一下,并且还需要host代码的grid x和y也得相应的换一下。本质上这只是block x和y的一维简单映射,要一劳永逸的话直接引入block swizzle
不止这里,得是所有的blockIdx.x和blockIdx.y都交换一下,并且还需要host代码的grid x和y也得相应的换一下。本质上这只是block x和y的一维简单映射,要一劳永逸的话直接引入block swizzle
322 dim3 grid((n + BLOCK_DIM - 1) / BLOCK_DIM, (m + BLOCK_DIM - 1) / BLOCK_DIM);
不好意思,搞错了,看了下整个kernel逻辑是按照X表示N,Y表示M写的,host的代码这里有错误,grid.x表示了M,grid.y表示了N,322行修改成上面这样就可以了。
https://github.com/njuhope/cuda_sgemm/blob/4d19c09de2b19adb21d47ed8c84d33a80a46c4d3/gemm.cu#L40-L48, 请教一下,blockIdx.x 是展开M,blockIdx.y是展开N,计算offset,应该是blockIdx.x TILE_Y, blockIdx.y TILE_X?
不过试着改成上面代码,调了下不过无法pass