memory_sm75.h

Go to the documentation of this file.

/***************************************************************************************************
 * Copyright (c) 2017-2019, NVIDIA CORPORATION.  All rights reserved.
 *
 * Redistribution and use in source and binary forms, with or without modification, are permitted
 * provided that the following conditions are met:
 *     * Redistributions of source code must retain the above copyright notice, this list of
 *       conditions and the following disclaimer.
 *     * Redistributions in binary form must reproduce the above copyright notice, this list of
 *       conditions and the following disclaimer in the documentation and/or other materials
 *       provided with the distribution.
 *     * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
 *       to endorse or promote products derived from this software without specific prior written
 *       permission.
 *
 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
 * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
 * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
 * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
 * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
 * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
 * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
 * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 *
 **************************************************************************************************/
#pragma once

#include "cutlass/array.h"
#include "cutlass/layout/matrix.h"

namespace cutlass {
namespace arch {


template <
  typename Layout,
  int MatrixCount
>
inline __device__ void ldsm(Array<unsigned, MatrixCount> & D, void const* ptr);

//
// Specializations
//

#if (__CUDACC_VER_MAJOR__ == 10) && (__CUDACC_VER_MINOR__ == 2)
  #define CUDA_NVVM_GET_SHARED_POINTER_SUPPORTED 1
#else
  #define CUDA_NVVM_GET_SHARED_POINTER_SUPPORTED 0
#endif

#if ! defined(CUDA_NVVM_GET_SHARED_POINTER_ENABLED)
  #define CUDA_NVVM_GET_SHARED_POINTER_ENABLED (CUDA_NVVM_GET_SHARED_POINTER_SUPPORTED)
#endif

#if ! defined(CUDA_LDMATRIX_SUPPORTED)
  #define CUDA_LDMATRIX_SUPPORTED ((__CUDACC_VER_MAJOR__ == 10) && (__CUDACC_VER_MINOR__ >= 2))
#endif

#if ! defined(CUDA_LDMATRIX_ENABLED)
  #define CUDA_LDMATRIX_ENABLED (CUDA_LDMATRIX_SUPPORTED)
#endif

#if (CUDA_LDMATRIX_ENABLED && defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 750))
  #define CUDA_LDMATRIX_ACTIVATED 1
#else
  #define CUDA_LDMATRIX_ACTIVATED 0
#endif

#if defined(CUTLASS_GET_SMEM_POINTER)
  // Use the existing implementation
#elif CUDA_NVVM_GET_SHARED_POINTER_ENABLED
  #if ! defined(NVVM_GET_SMEM_POINTER)
    #define NVVM_GET_SMEM_POINTER
  extern "C" {
    //
    // This NVVM intrinsic is subject to change in future versions of CUDA.
    // Clients should not call it directly. Rather, they should use the 
    // cutlass::arch::ldsm<>() template.
    //
    __device__ uint32_t __nvvm_get_smem_pointer(void*);
  }
  #endif
  #define CUTLASS_GET_SMEM_POINTER(ptr) __nvvm_get_smem_pointer((void*)ptr)
#endif


template <>
inline __device__ void ldsm<layout::RowMajor, 1>(
    Array<unsigned, 1> & D,
    void const* ptr) {

  #if CUDA_LDMATRIX_ACTIVATED

    unsigned addr = CUTLASS_GET_SMEM_POINTER(ptr);

    int x;
    asm volatile ("ldmatrix.sync.aligned.x1.m8n8.shared.b16 {%0}, [%1];" : "=r"(x) : "r"(addr));
    reinterpret_cast<int &>(D) = x;

  #else

    assert(0);

  #endif
}


template <>
inline __device__ void ldsm<layout::RowMajor, 2>(
    Array<unsigned, 2> & D,
    void const* ptr) {

  #if CUDA_LDMATRIX_ACTIVATED

    unsigned addr = CUTLASS_GET_SMEM_POINTER(ptr);

    int x, y;
    asm volatile ("ldmatrix.sync.aligned.x2.m8n8.shared.b16 {%0, %1}, [%2];" : "=r"(x), "=r"(y) : "r"(addr));
    reinterpret_cast<int2 &>(D) = make_int2(x, y);

  #else

    assert(0);

  #endif
}


template <>
inline __device__ void ldsm<layout::RowMajor, 4>(
    Array<unsigned, 4> & D,
    void const* ptr) {

  #if CUDA_LDMATRIX_ACTIVATED

    unsigned addr = CUTLASS_GET_SMEM_POINTER(ptr);

    int x, y, z, w;
    asm volatile ("ldmatrix.sync.aligned.x4.m8n8.shared.b16 {%0, %1, %2, %3}, [%4];" : "=r"(x), "=r"(y), "=r"(z), "=r"(w) : "r"(addr));
    reinterpret_cast<int4 &>(D) = make_int4(x, y, z, w);

  #else

    assert(0);

  #endif
}

//
// Transpose on 16b granularity
//

template <>
inline __device__ void ldsm<layout::ColumnMajor, 1>(
    Array<unsigned, 1> & D,
    void const* ptr) {
  #if CUDA_LDMATRIX_ACTIVATED

    unsigned addr = CUTLASS_GET_SMEM_POINTER(ptr);

    int x;
    asm volatile ("ldmatrix.sync.aligned.x1.trans.m8n8.shared.b16 {%0}, [%1];" : "=r"(x) : "r"(addr));
    reinterpret_cast<int &>(D) = x;

  #else

    assert(0);

  #endif
}


template <>
inline __device__ void ldsm<layout::ColumnMajor, 2>(
    Array<unsigned, 2> & D,
    void const* ptr) {

  #if CUDA_LDMATRIX_ACTIVATED

    unsigned addr = CUTLASS_GET_SMEM_POINTER(ptr);

    int x, y;
    asm volatile ("ldmatrix.sync.aligned.x2.trans.m8n8.shared.b16 {%0, %1}, [%2];" : "=r"(x), "=r"(y) : "r"(addr));
    reinterpret_cast<int2 &>(D) = make_int2(x, y);

  #else

    assert(0);

  #endif
}


template <>
inline __device__ void ldsm<layout::ColumnMajor, 4>(
    Array<unsigned, 4> & D,
    void const* ptr) {

  #if CUDA_LDMATRIX_ACTIVATED

    unsigned addr = CUTLASS_GET_SMEM_POINTER(ptr);

    int x, y, z, w;
    asm volatile ("ldmatrix.sync.aligned.x4.trans.m8n8.shared.b16 {%0, %1, %2, %3}, [%4];" : "=r"(x), "=r"(y), "=r"(z), "=r"(w) : "r"(addr));
    reinterpret_cast<int4 &>(D) = make_int4(x, y, z, w);

  #else

    assert(0);

  #endif
}

} // namespace arch
} // namespace cutlass