CUDA在内核代码中多次乘法运算

Posted

tags:

篇首语:本文由小常识网(cha138.com)小编为大家整理,主要介绍了CUDA在内核代码中多次乘法运算相关的知识,希望对你有一定的参考价值。

矩阵乘法的函数:

__global__ void gpu_matrix_mult(float *a, float *b, float *c, int m, int n, int k)
{
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    float sum = 0;
    if (col < k && row < m)
    {
        for (int i = 0; i < n; i++)
        {
            sum += a[row * n + i] * b[i * k + col];
        }
        c[row * k + col] = sum;
    }
}

然后在以下循环中调用该函数:

int currentActivityCount = -1;

while (activityCount != currentActivityCount)
{
    if (currentActivityCount > -1)
    {
        cudaMemcpy(d_b, h_b_new, sizeof(int)*m*k, cudaMemcpyHostToDevice);
    }

    gpu_matrix_mult << <dimGrid, dimBlock >> >(d_a, d_b, d_c, m, n, k);

    cudaMemcpy(h_c, d_c, sizeof(int)*m*k, cudaMemcpyDeviceToHost);

    currentActivityCount = activityCount;
    activityCount = 0;

    for (int i = 0; i < m; ++i)
    {
        for (int j = 0; j < k; ++j)
        {
            if (h_c[i*k + j] >= 0.5)
            {
                activityCount++;

                h_b_new[i * k + j] = 1;
            }
            else
            {
                h_b_new[i * k + j] = 0;
            }
        }
    }

    during++;
    printf("Count of activity: %d During: %d
", activityCount, during);
}

我的目标是将此循环移动到“gpu_matrix_mult”函数中,以便GPU之间的数据传输仅发生在调用函数之前和之后的两倍,而不是在循环的每次迭代中。我一直在尝试某些方法,但都没有效果。这种解决方案是否可行?

答案

您可以在内核中执行以下操作:

__device__ int activityCount;
__global__ void gpu_matrix_mult(float *a, float *b0, float *b1, float *c, int m, int n, int k)
{
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    float sum = 0;
    if (col < k && row < m)
    {
        for (int i = 0; i < n; i++)
        {
            sum += a[row * n + i] * b0[i * k + col];
        }
        c[row * k + col] = sum;
        if (sum >= 0.5)
        {
            atomicAdd(&activityCount, 1);
            b1[i * k + j] = 1;
        }
        else
        {
            b1[i * k + j] = 0;
        }
    }
}

// .............


int currentActivityCount = -1;
int activityCount_h = 0;
while (activityCount_h != currentActivityCount)
{
    if (currentActivityCount > -1)
    {
        float *tmp = d_b0;
        d_b0 = d_b1;
        d_b1 = tmp;
    }
    currentActivityCount = activityCount_h;
    activityCount_h = 0;
    cudaMemcpyToSymbol(activityCount, &activityCount_h, sizeof(int));
    gpu_matrix_mult<<<dimGrid, dimBlock>>>(d_a, d_b0, d_b1, d_c, m, n, k);
    cudaMemcpyfromSymbol(&activityCount_h, activity, sizeof(int));

    during++;
    printf("Count of activity: %d During: %d
", activityCount, during);
}

[显然从未编译或运行,使用风险自负]

即,在矩阵乘法之后,用于计算activityCount的内环可以在设备的内核中运行。这需要GPU上的两个b矩阵在内存中,但主机上只需要指针交换来更新它们,这基本上是零成本。每次外循环迭代,内存传输减少到一个整数两次,这将相当快。

以上是关于CUDA在内核代码中多次乘法运算的主要内容,如果未能解决你的问题,请参考以下文章

cuda编程-矩阵乘法

在 CUDA 的平铺矩阵乘法中访问矩阵作为其转置

使用 CUDA 进行矩阵乘法:2D 块与 1D 块

如何在 python 中并行化以下代码片段?

如何获取 CUDA 内核的汇编代码?

如何从主机代码中中断或取消 CUDA 内核