Skip to content

Fix errors and warnings with CUDA 9.0 builds #2923

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
9prady9 merged 4 commits into from
Jun 5, 2020
Merged

Fix errors and warnings with CUDA 9.0 builds #2923

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
cc87488
Fix libs for the CUDA 9 Toolkit. Remove rdc and dlink flags
umar456 Jun 5, 2020
b23adf6
Fix several error messages when compiling against CUDA 9.0
umar456 Jun 5, 2020
4466bec
Fix CUSOLVER_CHECK error message
umar456 Jun 5, 2020
d2b3273
Fix several warnings with older compilers
umar456 Jun 5, 2020
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
16 changes: 15 additions & 1 deletion src/api/cpp/common.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,6 +15,10 @@
#include "half.hpp"
#pragma GCC diagnostic pop

#ifdef AF_CUDA
#include <cuda_fp16.h>
#endif

#include <cstring>

namespace af {
Expand All @@ -36,10 +40,20 @@ static inline dim_t getFNSD(const int dim, af::dim4 dims) {
namespace {
// casts from one type to another. Needed for af_half conversions specialization
template<typename To, typename T>
To cast(T in) {
inline To cast(T in) {
return static_cast<To>(in);
}

#if defined(AF_CUDA) && CUDA_VERSION < 10000
template<>
inline __half cast<__half, double>(double in) {
__half_raw out;
half_float::half h(in);
memcpy(&out, &h, sizeof(__half_raw));
return out;
}
#endif

template<>
[[gnu::unused]] af_half cast<af_half, double>(double in) {
half_float::half tmp = static_cast<half_float::half>(in);
Expand Down
23 changes: 11 additions & 12 deletions src/backend/cuda/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,7 +7,7 @@

dependency_check(CUDA_FOUND "CUDA not found.")
if(AF_WITH_CUDNN)
dependency_check(cuDNN_FOUND "CUDA not found.")
dependency_check(cuDNN_FOUND "CUDNN not found.")
endif()

include(AFcuda_helpers)
Expand All @@ -34,7 +34,7 @@ endif()

# Find if CUDA Toolkit is at least 10.0 to use static
# lapack library. Otherwise, we have to use regular shared library
if(UNIX AND CUDA_VERSION_MAJOR VERSION_GREATER 10 OR CUDA_VERSION_MAJOR VERSION_EQUAL 10)
if(UNIX AND (CUDA_VERSION_MAJOR VERSION_GREATER 10 OR CUDA_VERSION_MAJOR VERSION_EQUAL 10))
set(use_static_cuda_lapack ON)
else()
set(use_static_cuda_lapack OFF)
Expand All @@ -52,7 +52,6 @@ if(UNIX)
# FIXME When NVCC resolves this particular issue.
# NVCC doesn't like -l<full_path_static_lib>, hence we cannot
# use ${CMAKE_*_LIBRARY} variables in the following flags.
set(af_cuda_static_flags "-rdc=true;-dlink")
set(af_cuda_static_flags "${af_cuda_static_flags};-lculibos")
set(af_cuda_static_flags "${af_cuda_static_flags};-lcublas_static")
set(af_cuda_static_flags "${af_cuda_static_flags};-lcublasLt_static")
Expand All @@ -71,7 +70,7 @@ if(UNIX)

set(af_cuda_static_flags "${af_cuda_static_flags};-lcusolver_static")
else()
set(cusolver_lib "${CUDA_cusolver_LIBRARY}")
set(cusolver_lib "${CUDA_cusolver_LIBRARY}" OpenMP::OpenMP_CXX)
endif()
endif()

Expand All @@ -89,12 +88,6 @@ message(STATUS "CUDA_architecture_build_targets: ${CUDA_architecture_build_targe

set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};${cuda_architecture_flags})

if(${CUDA_SEPARABLE_COMPILATION})
# Enable relocatable device code generation for separable
# compilation which is in turn required for any device linking done.
set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-rdc=true)
endif()

mark_as_advanced(
CUDA_LIBRARIES_PATH
CUDA_architecture_build_targets)
Expand Down Expand Up @@ -301,13 +294,19 @@ if(UNIX)
-Wl,--start-group
${CUDA_culibos_LIBRARY} #also a static libary
${CUDA_cublas_static_LIBRARY}
${CUDA_cublasLt_static_LIBRARY}
${CUDA_cufft_static_LIBRARY}
${CUDA_lapack_static_LIBRARY}
${CUDA_cusparse_static_LIBRARY}
${cusolver_static_lib}
-Wl,--end-group
)

if(CUDA_VERSION VERSION_GREATER 9.5)

@9prady9 9prady9 Jun 5, 2020

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

nit: I think a better check to avoid any unexpected issues (with decimals and all against different cmake versions) with cmake is to do

if(CUDA_VERSION_MAJOR VERSION_GREATER 9)

On another note, I will take care of this with some other change.

target_link_libraries(af_cuda_static_cuda_library
PRIVATE
${CUDA_cublasLt_static_LIBRARY}
${CUDA_lapack_static_LIBRARY})

@9prady9 9prady9 Jun 5, 2020

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Makes sense, these are not there prior to toolkit 10 - I should have thought about this

endif()

set(CUDA_SEPARABLE_COMPILATION ${pior_val_CUDA_SEPARABLE_COMPILATION})
else()
target_link_libraries(af_cuda_static_cuda_library
Expand Down
4 changes: 2 additions & 2 deletions src/backend/cuda/ThrustArrayFirePolicy.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -34,8 +34,8 @@ get_temporary_buffer(ThrustArrayFirePolicy, std::ptrdiff_t n) {
}

template<typename Pointer>
void return_temporary_buffer(ThrustArrayFirePolicy, Pointer p) {
memFree(p.get());
inline void return_temporary_buffer(ThrustArrayFirePolicy, Pointer p) {
memFree(thrust::raw_pointer_cast(p));
}

} // namespace cuda
11 changes: 10 additions & 1 deletion src/backend/cuda/blas.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -216,7 +216,8 @@ cublasStatus_t gemmDispatch(BlasHandle handle, cublasOperation_t lOpts,
const Array<T> &rhs, dim_t rStride, const T *beta,
Array<T> &out, dim_t oleading) {
auto prop = getDeviceProp(getActiveDeviceId());
if (prop.major > 3) {
#if __CUDACC_VER_MAJOR__ >= 10
if (prop.major > 3 && __CUDACC_VER_MAJOR__ >= 10) {
return cublasGemmEx(
blasHandle(), lOpts, rOpts, M, N, K, alpha, lhs.get(), getType<T>(),
lStride, rhs.get(), getType<T>(), rStride, beta, out.get(),
Expand All @@ -233,11 +234,15 @@ cublasStatus_t gemmDispatch(BlasHandle handle, cublasOperation_t lOpts,
// type is CUDA_R_32F?
selectGEMMAlgorithm<T>());
} else {
#endif
using Nt = typename common::kernel_type<T>::native;
return gemm_func<Nt>()(blasHandle(), lOpts, rOpts, M, N, K, (Nt *)alpha,
(Nt *)lhs.get(), lStride, (Nt *)rhs.get(),
rStride, (Nt *)beta, (Nt *)out.get(), oleading);

#if __CUDACC_VER_MAJOR__ >= 10
}
#endif
}

template<typename T>
Expand All @@ -248,6 +253,7 @@ cublasStatus_t gemmBatchedDispatch(BlasHandle handle, cublasOperation_t lOpts,
const T *beta, T **optrs, int oStrides,
int batchSize) {
auto prop = getDeviceProp(getActiveDeviceId());
#if __CUDACC_VER_MAJOR__ >= 10
if (prop.major > 3) {
return cublasGemmBatchedEx(
blasHandle(), lOpts, rOpts, M, N, K, alpha, (const void **)lptrs,
Expand All @@ -264,12 +270,15 @@ cublasStatus_t gemmBatchedDispatch(BlasHandle handle, cublasOperation_t lOpts,
// type is CUDA_R_32F?
selectGEMMAlgorithm<T>());
} else {
#endif
using Nt = typename common::kernel_type<T>::native;
return gemmBatched_func<Nt>()(
blasHandle(), lOpts, rOpts, M, N, K, (const Nt *)alpha,
(const Nt **)lptrs, lStrides, (const Nt **)rptrs, rStrides,
(const Nt *)beta, (Nt **)optrs, oStrides, batchSize);
#if __CUDACC_VER_MAJOR__ >= 10
}
#endif
}

template<typename T>
Expand Down
20 changes: 10 additions & 10 deletions src/backend/cuda/cusolverDn.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,16 +14,16 @@ namespace cuda {

const char* errorString(cusolverStatus_t err);

#define CUSOLVER_CHECK(fn) \
do { \
cusolverStatus_t _error = fn; \
if (_error != CUSOLVER_STATUS_SUCCESS) { \
char _err_msg[1024]; \
snprintf(_err_msg, sizeof(_err_msg), "CUBLAS Error (%d): %s\n", \
(int)(_error), cuda::errorString(_error)); \
\
AF_ERROR(_err_msg, AF_ERR_INTERNAL); \
} \
#define CUSOLVER_CHECK(fn) \
do { \
cusolverStatus_t _error = fn; \
if (_error != CUSOLVER_STATUS_SUCCESS) { \
char _err_msg[1024]; \
snprintf(_err_msg, sizeof(_err_msg), "CUSOLVER Error (%d): %s\n", \
(int)(_error), cuda::errorString(_error)); \
\
AF_ERROR(_err_msg, AF_ERR_INTERNAL); \
} \
} while (0)

} // namespace cuda
9 changes: 5 additions & 4 deletions src/backend/cuda/jit/kernel_generators.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -71,8 +71,9 @@ void generateBufferRead(std::stringstream& kerStream, int id,
<< "];\n";
}

void generateShiftNodeOffsets(std::stringstream& kerStream, int id,
bool is_linear, const std::string& type_str) {
inline void generateShiftNodeOffsets(std::stringstream& kerStream, int id,

@9prady9 9prady9 Jun 5, 2020

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

nit: looks like this is just to make the compiler silent even though this won't be actually inlined.

bool is_linear,
const std::string& type_str) {
UNUSED(is_linear);
std::string idx_str = std::string("idx") + std::to_string(id);
std::string info_str = std::string("in") + std::to_string(id);
Expand All @@ -99,8 +100,8 @@ void generateShiftNodeOffsets(std::stringstream& kerStream, int id,
kerStream << type_str << " *in" << id << "_ptr = in" << id << ".ptr;\n";
}

void generateShiftNodeRead(std::stringstream& kerStream, int id,
const std::string& type_str) {
inline void generateShiftNodeRead(std::stringstream& kerStream, int id,
const std::string& type_str) {
kerStream << type_str << " val" << id << " = in" << id << "_ptr[idx" << id
<< "];\n";
}
Expand Down
2 changes: 1 addition & 1 deletion src/backend/cuda/kernel/mean.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -26,7 +26,7 @@

namespace cuda {

__host__ __device__ auto operator*(float lhs, __half rhs) -> __half {
__device__ auto operator*(float lhs, __half rhs) -> __half {
return __float2half(lhs * __half2float(rhs));
}

Expand Down
2 changes: 2 additions & 0 deletions src/backend/cuda/memory.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -131,6 +131,8 @@ INSTANTIATE(short)
INSTANTIATE(ushort)
INSTANTIATE(half)

template void memFree(void *ptr);

Allocator::Allocator() { logger = common::loggerFactory("mem"); }

void Allocator::shutdown() {
Expand Down
40 changes: 20 additions & 20 deletions src/backend/cuda/types.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -47,69 +47,69 @@ using data_t = typename common::kernel_type<T>::data;
#ifndef __CUDACC_RTC__
namespace {
template<typename T>
const char *shortname(bool caps = false) {
inline const char *shortname(bool caps = false) {
return caps ? "Q" : "q";
}
template<>
const char *shortname<float>(bool caps) {
inline const char *shortname<float>(bool caps) {
return caps ? "S" : "s";
}
template<>
const char *shortname<double>(bool caps) {
inline const char *shortname<double>(bool caps) {
return caps ? "D" : "d";
}
template<>
const char *shortname<cfloat>(bool caps) {
inline const char *shortname<cfloat>(bool caps) {
return caps ? "C" : "c";
}
template<>
const char *shortname<cdouble>(bool caps) {
inline const char *shortname<cdouble>(bool caps) {
return caps ? "Z" : "z";
}
template<>
const char *shortname<int>(bool caps) {
inline const char *shortname<int>(bool caps) {
return caps ? "I" : "i";
}
template<>
const char *shortname<uint>(bool caps) {
inline const char *shortname<uint>(bool caps) {
return caps ? "U" : "u";
}
template<>
const char *shortname<char>(bool caps) {
inline const char *shortname<char>(bool caps) {
return caps ? "J" : "j";
}
template<>
const char *shortname<uchar>(bool caps) {
inline const char *shortname<uchar>(bool caps) {
return caps ? "V" : "v";
}
template<>
const char *shortname<intl>(bool caps) {
inline const char *shortname<intl>(bool caps) {
return caps ? "X" : "x";
}
template<>
const char *shortname<uintl>(bool caps) {
inline const char *shortname<uintl>(bool caps) {
return caps ? "Y" : "y";
}
template<>
const char *shortname<short>(bool caps) {
inline const char *shortname<short>(bool caps) {
return caps ? "P" : "p";
}
template<>
const char *shortname<ushort>(bool caps) {
inline const char *shortname<ushort>(bool caps) {
return caps ? "Q" : "q";
}
template<>
const char *shortname<common::half>(bool caps) {
inline const char *shortname<common::half>(bool caps) {
return caps ? "H" : "h";
}

template<typename T>
const char *getFullName();
inline const char *getFullName();

#define SPECIALIZE(T) \
template<> \
const char *getFullName<T>() { \
return #T; \
#define SPECIALIZE(T) \
template<> \
inline const char *getFullName<T>() { \
return #T; \
}

SPECIALIZE(float)
Expand All @@ -126,7 +126,7 @@ SPECIALIZE(unsigned long long)
SPECIALIZE(long long)

template<>
const char *getFullName<common::half>() {
inline const char *getFullName<common::half>() {
return "half";
}
#undef SPECIALIZE
Expand Down