-1

我想在 CUDA 中实现一个简单的矩阵乘法。矩阵的维度是在运行时确定的，我还想使用共享内存来提高性能。我已经实现了这样的功能，但是每次运行它时，都会出现此错误：

mulKernel launch failed: an illegal memory access was encountered

我也不确定是否可以使用 malloc 分配共享内存。但是，如果我想使用这样的东西

__shared__ float matrM_sm[tile_width][tile_width];

编译器抱怨必须在运行时知道 tile_width ......

我已经尝试了我能想到的一切，也尝试了各种建议，但没有一个奏效。这是函数（完整的工作文件可以在这里找到）：

__global__ void mulKernelSM(float *matrR, const float *matrM, const float *matrN,
   const int m_x, const int m_y, const int n_x, const int n_y, const int tile_width)
{
    int i, j;
    extern __shared__ float shared[];
    float *matrM_sm = shared;
    float *matrN_sm = &shared[tile_width * tile_width];

    int bx = blockIdx.x;
    int by = blockIdx.y;
    int tx = threadIdx.x;
    int ty = threadIdx.y;

    int row = by * tile_width + ty;
    int col = bx * tile_width + tx;

    float tmp;
    int limit = ceil(m_y / (float) tile_width);
    for (i = 0; i < limit; i++)
    {
        tmp = 0.0;

        if (i * tile_width + tx < m_y && row < m_x)
            matrM_sm[ty * tile_width + tx] = matrM[row * m_y + (i * tile_width + tx)];
        else
            matrM_sm[ty * tile_width + tx] = 0.0;

        if (i * tile_width + ty < n_x && col < n_y)
            matrN_sm[ty * tile_width + tx] = matrN[col + (i * tile_width + ty) * n_y];
        else
            matrN_sm[ty * tile_width + tx] = 0.0;

        __syncthreads();

        for (j = 0; j < tile_width; j++)
            tmp += matrM_sm[ty * tile_width + j] * matrN_sm[j * tile_width + tx];

        __syncthreads();
    }

    if (row < m_x && col < n_y)
        matrR[row * n_y + col] = tmp;
}

基本布局应该可以工作，因为我还实现了一个没有共享内存的版本，它工作得很好。不带共享内存的函数如下：

__global__ void mulKernel(float *matrR, const float *matrM, const float *matrN,
    const int m_x, const int m_y, const int n_x, const int n_y)
{
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int i;

    if ((row < m_x) && (col < n_y))
    {
        float tmp = 0.0;
        for (i = 0; i < m_y; i++)
        {
            tmp += matrM[row * m_y + i] * matrN[col + n_y * i];
        }

        matrR[row * n_y + col] = tmp;
    }
}

如果有任何信息丢失，我会立即提供。

4

1 回答 1

-1

你换了row, col。此外，我相信要获得全局线程索引，您应该这样做int x_global = threadIdx.x + blockDim.x * threadIdx.y

于 2016-11-15T19:08:40.487 回答