Cuda Gemm вне вопроса в библиотеке магмы - PullRequest
Новая
       61

Cuda Gemm вне вопроса в библиотеке магмы

0 голосов
/ 13 апреля 2020

Я нахожу реализацию gemm cuda в magma, код такой: 

static __device__ 
void gemm_template_device_nn(
    int M, int N, int K,
    const T* __restrict__ A, int LDA,
    const T* __restrict__ B, int LDB,
    T*       __restrict__ C, int LDC,
    T alpha, T beta )
{
#if (__CUDA_ARCH__ >= 200)
    int idx = threadIdx.x;  // thread's m dimension
    int idy = threadIdx.y;  // thread's n dimension

    int idt = DIM_X * idy + idx;    // thread's global number

    int idxA = idt % DIM_XA;    // idx within A
    int idyA = idt / DIM_XA;    // idy within A

    int idxB = idt % DIM_XB;    // idx within B
    int idyB = idt / DIM_XB;    // idy within B

    int blx = blockIdx.x;   // block's m dimension
    int bly = blockIdx.y;   // block's n dimension

    __shared__ T sA[BLK_K][BLK_M+1];      // +1 only required if A is transposed
    __shared__ T sB[BLK_N][BLK_K+1];      // +1 always required

    // Registers for the innermost loop
    T rC[THR_N][THR_M];
    T rA[THR_M];
    T rB[THR_N];

    T ra[BLK_K/DIM_YA][BLK_M/DIM_XA];
    T rb[BLK_N/DIM_YB][BLK_K/DIM_XB];

    const T *offs_dA = A + blx*BLK_M     + idyA*LDA + idxA;
    ptrdiff_t boundA = (LDA*(K-1) + M) - ( blx*BLK_M  + idyA*LDA + idxA ) -1;

    const T *offs_dB = B + bly*BLK_N*LDB + idyB*LDB + idxB;
    ptrdiff_t boundB = (LDB*(N-1) + K) - ( bly*BLK_N*LDB + idyB*LDB + idxB ) -1;

    int m, n, k, kk;


    // Zero C
    #pragma unroll
    for (n = 0; n < THR_N; n++)
        #pragma unroll
        for (m = 0; m < THR_M; m++)
            rC[n][m] = make_FloatingPoint(0.0, 0.0);

    #pragma unroll
    for (n = 0; n < BLK_K; n += DIM_YA)
        #pragma unroll
        for (m = 0; m < BLK_M; m += DIM_XA)
            sA[n+idyA][m+idxA] = fetch(A, m, n, boundA);

    #pragma unroll
    for (n = 0; n < BLK_N; n += DIM_YB)
        #pragma unroll
        for (m = 0; m < BLK_K; m += DIM_XB)
            sB[n+idyB][m+idxB] = fetch(B, m, n, boundB);

    __syncthreads();

Мой вопрос заключается в извлечении из части глобальной памяти в общую память, магма использует макрос выборки для получения этой функции; определение выборки выглядит следующим образом: 

#ifdef TEXTURE_1D
    #define fetch(A, m, n, bound) tex_fetch(Mjoin1(tex_ref_##A##magma_,precision), coord_##A + n*LD##A+m)
    #define Mjoin1(Mname,Mp) Mjoin(Mname,Mp)
    #define Mjoin(Mname,Mp) Mname##Mp
#else
    #define fetch(A, m, n, bound) offs_d##A[min(n*LD##A+m, bound)]
#endif

первая использованная текстура для выборки, вторая - загрузка буфера использования; у меня вопрос к этим двум методам, как магма может убедиться, что выборка за пределами границы матрицы вернет 0?

...