CUBLAS gemm 中的 memset 总是在默认流中启动

Question

我注意到，当从主机每次调用 gemm 调用 cublasSgemm 函数时，有 3 个内核调用：memset、scal_kernel 和 gemm 内核本身（例如 sgemm_large）。即使我使用在设备内存中分配的常量 alpha/beta 也会发生这种情况。虽然 memset 和 scal_kernel 的开销相对较小，但问题是 memset 总是在默认流中启动，导致不必要的同步。

代码：

__constant__ __device__ float alpha = 1;
__constant__ __device__ float beta = 1;

int main()
{
    // ... memory allocation skipped ...
    float* px = thrust::raw_pointer_cast(x.data());
    float* py = thrust::raw_pointer_cast(y.data());
    float* pmat = thrust::raw_pointer_cast(mat.data());
    for (int iter = 0; iter < 3; ++iter)
    {
        cbstatus = cublasSgemm(cbh, CUBLAS_OP_N, CUBLAS_OP_N, crow, ccol, cshared, &alpha, px, crow, py, cshared, &beta, pmat, crow);
        assert(0 == cbstatus);
    }
}

这是我在分析器中看到的：

问题：有没有办法避免 memset 或让它在分配给 CUBLAS 句柄的流中运行？ 一种想法是使用 DP 并运行设备版本的 gemm 函数，但这仅适用于 CC 3.0 及更高版本。

Answer 1

试试下面的代码。除了不可避免的内存分配和副本之外，该代码被设想为只有一个cublasSgemm 调用。你会看到的

1. 您只启动了一个内核 (gemm_kernel1x1_core)；
2. 对cublasSgemm 的两次调用在两个不同的流中完美运行。

在图片中，显示了 Visual Profiler 时间线。

我的系统：GeForce 540M、Windows 7、CUDA 5.5。

#include <conio.h>
#include <stdio.h>
#include <assert.h>

#include <cublas_v2.h> 

/********************/
/* CUDA ERROR CHECK */
/********************/
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
    if (code != cudaSuccess) 
    {
        fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
        if (abort) { getchar(); exit(code); }
    }
}

/**********************/
/* cuBLAS ERROR CHECK */
/**********************/
#ifndef cublasSafeCall
#define cublasSafeCall(err)     __cublasSafeCall(err, __FILE__, __LINE__)
#endif

inline void __cublasSafeCall(cublasStatus_t err, const char *file, const int line)
{
    if( CUBLAS_STATUS_SUCCESS != err) {
        fprintf(stderr, "CUBLAS error in file '%s', line %d\n \nerror %d \nterminating!\n",__FILE__, __LINE__,err); 
        getch(); cudaDeviceReset(); assert(0); 
    }
}

/********/
/* MAIN */
/********/
int main()
{
    int N = 5;

    float *A1, *A2, *B1, *B2, *C1, *C2;
    float *d_A1, *d_A2, *d_B1, *d_B2, *d_C1, *d_C2;

    A1 = (float*)malloc(N*N*sizeof(float));
    B1 = (float*)malloc(N*N*sizeof(float));
    C1 = (float*)malloc(N*N*sizeof(float));

    A2 = (float*)malloc(N*N*sizeof(float));
    B2 = (float*)malloc(N*N*sizeof(float));
    C2 = (float*)malloc(N*N*sizeof(float));

    gpuErrchk(cudaMalloc((void**)&d_A1,N*N*sizeof(float)));
    gpuErrchk(cudaMalloc((void**)&d_B1,N*N*sizeof(float)));
    gpuErrchk(cudaMalloc((void**)&d_C1,N*N*sizeof(float)));
    gpuErrchk(cudaMalloc((void**)&d_A2,N*N*sizeof(float)));
    gpuErrchk(cudaMalloc((void**)&d_B2,N*N*sizeof(float)));
    gpuErrchk(cudaMalloc((void**)&d_C2,N*N*sizeof(float)));

    for (int i=0; i<N*N; i++) {
        A1[i] = ((float)rand()/(float)RAND_MAX);
        A2[i] = ((float)rand()/(float)RAND_MAX);
        B1[i] = ((float)rand()/(float)RAND_MAX);
        B2[i] = ((float)rand()/(float)RAND_MAX);
    }
    gpuErrchk(cudaMemcpy(d_A1, A1, N*N*sizeof(float), cudaMemcpyHostToDevice));
    gpuErrchk(cudaMemcpy(d_B1, B1, N*N*sizeof(float), cudaMemcpyHostToDevice));
    gpuErrchk(cudaMemcpy(d_A2, A2, N*N*sizeof(float), cudaMemcpyHostToDevice));
    gpuErrchk(cudaMemcpy(d_B2, B2, N*N*sizeof(float), cudaMemcpyHostToDevice));

    cublasHandle_t handle;
    cublasSafeCall(cublasCreate(&handle));

    cudaStream_t stream1, stream2;
    gpuErrchk(cudaStreamCreate(&stream1));
    gpuErrchk(cudaStreamCreate(&stream2));

    float alpha = 1.f;
    float beta = 1.f;

    cublasSafeCall(cublasSetStream(handle,stream1));
    cublasSafeCall(cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, N, N, N, &alpha, d_A1, N, d_B1, N, &beta, d_C1, N));
    cublasSafeCall(cublasSetStream(handle,stream2));
    cublasSafeCall(cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, N, N, N, &alpha, d_A2, N, d_B2, N, &beta, d_C2, N));

    gpuErrchk(cudaDeviceReset());

    return 0;

 }

Answer 2

CUBLAS5.5 中存在一个错误，其中在 k >> m,n 的专用路径中使用 cudaMemset 而不是 cudaMemsetAsync。

在 CUBLAS6.0 RC 中已修复。如果您是注册开发者，则可以访问它。

顺便说一句，我想知道您为什么使用 __constant__ __device__ 表示 alpha、beta。 你在用pointerMode = DEVICE吗？

如果没有，您可以简单地在主机上使用 alpha,beta。