CUBLAS gemm의 memset은 항상 기본 스트림으로 시작됩니다.

Question

호스트에서 gemm을 호출 할 때마다 cublasSgemm 함수를 호출 할 때 memset, scal_kernel 및 gemm 커널 자체 (예 : sgemm_large)의 3 가지 커널 호출이 있음을 발견했습니다. 이것은 디바이스 메모리에 할당 된 상수 alpha/beta를 사용하더라도 발생합니다. memset과 scal_kernel의 오버 헤드가 비교적 적지 만 문제는 memset이 불필요한 동기화를 유발하는 기본 스트림에서 항상 시작된다는 것입니다.

코드는 : 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); 
    } 
} 

이것은 제가 프로파일 러에서 볼 수있다 CUBLAS가 처리합니까? 하나의 아이디어는 DP를 사용하고 gemm 기능의 장치 버전을 실행하는 것이지만 CC 3.0 이상에서만 작동합니다.

Answer 1

cudaMemset 어디 K >> m, n은 특별한 경로 대신 cudaMemsetAsync 사용한 CUBLAS5.5에서 문제가 있었다.

CUBLAS6.0 RC에서 수정되었습니다. 등록 된 개발자는 액세스 할 수 있습니다.

Btw, 왜 __constant__ __device__을 알파, 베타로 사용하는지 궁금합니다. pointerMode = DEVICE을 사용하고 있습니까?

그렇지 않은 경우 단순히 호스트에서 알파, 베타를 사용할 수 있습니다.

Answer 2

아래 코드를 사용해보십시오. 이 코드는 피할 수없는 메모리 할당 및 사본을 제외하고는 cublasSgemm 전화 만 가지고 있다고 생각합니다. 그러면 알 수 있습니다.

1. 커널이 하나만 시작되었습니다 (gemm_kernel1x1_core).
2. cublasSgemm에 대한 두 번의 호출은 서로 다른 두 스트림에서 완벽하게 실행됩니다.

그림에서 비주얼 프로파일 러 타임 라인이 표시됩니다.

내 시스템 : 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; 

}