FazBrowse GitHub Viewer | Trending |
URL:
| Home
Tools: [Download Repo ZIP]   [Original HTTPS Page]

Fix errors and warnings with CUDA 9.0 builds by umar456 · Pull Request #2923 · arrayfire/arrayfire · GitHub

  1. arrayfire
  2. arrayfire

Repository navigation

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
Conversation 4 Commits 4 Checks 0 Files changed
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 .cpp  (1) .cu  (1) .hpp  (6) .txt  (1) All 4 file types selected
Viewed files
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Unified
Split
Hide whitespace
Diff view
Unified
Split
Hide whitespace
16 changes: 15 additions & 1 deletion src/api/cpp/common.hpp
Show comments
View file
Open in desktop
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters. Learn more about bidirectional Unicode characters
Show hidden characters
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
Show comments
View file
Open in desktop
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters. Learn more about bidirectional Unicode characters
Show hidden characters
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 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.

Choose a reason Spam Abuse Off Topic Outdated Duplicate Resolved Low Quality

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 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.

Choose a reason Spam Abuse Off Topic Outdated Duplicate Resolved Low Quality

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
Show comments
View file
Open in desktop
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters. Learn more about bidirectional Unicode characters
Show hidden characters
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
Show comments
View file
Open in desktop
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters. Learn more about bidirectional Unicode characters
Show hidden characters
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
Show comments
View file
Open in desktop
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters. Learn more about bidirectional Unicode characters
Show hidden characters
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
Show comments
View file
Open in desktop
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters. Learn more about bidirectional Unicode characters
Show hidden characters
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 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.

Choose a reason Spam Abuse Off Topic Outdated Duplicate Resolved Low Quality

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
Show comments
View file
Open in desktop
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters. Learn more about bidirectional Unicode characters
Show hidden characters
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
Show comments
View file
Open in desktop
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters. Learn more about bidirectional Unicode characters
Show hidden characters
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
Show comments
View file
Open in desktop
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters. Learn more about bidirectional Unicode characters
Show hidden characters
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
Back | FazBrowse Home | New Git URL