Skip to content

Compute capability fix #2996

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
umar456 merged 4 commits into from
Aug 15, 2020
Merged

Compute capability fix #2996

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
bb1066e
Fix checkAndSetDevMaxCompute when the device cc is greater than max
umar456 Aug 15, 2020
e8a69c2
Add utility header included from cuda_fp16.hpp for CUDA 9
umar456 Aug 15, 2020
fa29b90
Formatting and rewording of warning messages
umar456 Aug 15, 2020
5b350de
Fix errors and warnings in RNG for CUDA 9.0
umar456 Aug 14, 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
2 changes: 0 additions & 2 deletions src/backend/common/half.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -832,9 +832,7 @@ class alignas(2) half {
#endif

public:
#if CUDA_VERSION >= 10000
AF_CONSTEXPR
Comment thread
umar456 marked this conversation as resolved.
#endif
half() = default;

/// Constructor.
Expand Down
42 changes: 26 additions & 16 deletions src/backend/cuda/compile_module.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -143,20 +143,29 @@ Module compileModule(const string &moduleKey, const vector<string> &sources,
const vector<string> &kInstances, const bool sourceIsJIT) {
nvrtcProgram prog;
if (sourceIsJIT) {
array<const char *, 2> headers = {
constexpr const char *header_names[] = {
"utility",
"cuda_fp16.hpp",
"cuda_fp16.h",
};
constexpr size_t numHeaders = extent<decltype(header_names)>::value;
array<const char *, numHeaders> headers = {
"",
cuda_fp16_hpp,
cuda_fp16_h,
};
array<const char *, 2> header_names = {"cuda_fp16.hpp", "cuda_fp16.h"};
static_assert(headers.size() == numHeaders,
Comment thread
umar456 marked this conversation as resolved.
"headers array contains fewer sources than header_names");
NVRTC_CHECK(nvrtcCreateProgram(&prog, sources[0].c_str(),
moduleKey.c_str(), 2, headers.data(),
header_names.data()));
moduleKey.c_str(), numHeaders,
headers.data(), header_names));
} else {
constexpr static const char *includeNames[] = {
"math.h", // DUMMY ENTRY TO SATISFY cuComplex_h inclusion
"stdbool.h", // DUMMY ENTRY TO SATISFY af/defines.h inclusion
"stdlib.h", // DUMMY ENTRY TO SATISFY af/defines.h inclusion
"vector_types.h", // DUMMY ENTRY TO SATISFY cuComplex_h inclusion
"utility", // DUMMY ENTRY TO SATISFY cuda_fp16.hpp inclusion
"backend.hpp",
"cuComplex.h",
"jit.cuh",
Expand All @@ -183,12 +192,13 @@ Module compileModule(const string &moduleKey, const vector<string> &sources,
"minmax_op.hpp",
};

constexpr size_t NumHeaders = extent<decltype(includeNames)>::value;
static const array<string, NumHeaders> sourceStrings = {{
constexpr size_t numHeaders = extent<decltype(includeNames)>::value;
static const array<string, numHeaders> sourceStrings = {{
string(""), // DUMMY ENTRY TO SATISFY cuComplex_h inclusion
string(""), // DUMMY ENTRY TO SATISFY af/defines.h inclusion
string(""), // DUMMY ENTRY TO SATISFY af/defines.h inclusion
string(""), // DUMMY ENTRY TO SATISFY cuComplex_h inclusion
string(""), // DUMMY ENTRY TO SATISFY utility inclusion
string(backend_hpp, backend_hpp_len),
string(cuComplex_h, cuComplex_h_len),
string(jit_cuh, jit_cuh_len),
Expand Down Expand Up @@ -230,11 +240,11 @@ Module compileModule(const string &moduleKey, const vector<string> &sources,
sourceStrings[22].c_str(), sourceStrings[23].c_str(),
sourceStrings[24].c_str(), sourceStrings[25].c_str(),
sourceStrings[26].c_str(), sourceStrings[27].c_str(),
};
static_assert(extent<decltype(headers)>::value == NumHeaders,
sourceStrings[28].c_str()};
static_assert(extent<decltype(headers)>::value == numHeaders,
Comment thread
umar456 marked this conversation as resolved.
"headers array contains fewer sources than includeNames");
NVRTC_CHECK(nvrtcCreateProgram(&prog, sources[0].c_str(),
moduleKey.c_str(), NumHeaders, headers,
moduleKey.c_str(), numHeaders, headers,
includeNames));
}

Expand All @@ -246,6 +256,7 @@ Module compileModule(const string &moduleKey, const vector<string> &sources,
vector<const char *> compiler_options = {
arch.data(),
"--std=c++14",
"--device-as-default-execution-space",
Comment thread
umar456 marked this conversation as resolved.
#if !(defined(NDEBUG) || defined(__aarch64__) || defined(__LP64__))
"--device-debug",
"--generate-line-info"
Expand All @@ -256,7 +267,6 @@ Module compileModule(const string &moduleKey, const vector<string> &sources,
back_insert_iterator<vector<const char *>>(compiler_options),
[](const string &s) { return s.data(); });

compiler_options.push_back("--device-as-default-execution-space");
for (auto &instantiation : kInstances) {
NVRTC_CHECK(nvrtcAddNameExpression(prog, instantiation.c_str()));
}
Expand Down Expand Up @@ -391,8 +401,8 @@ Module loadModuleFromDisk(const int device, const string &moduleKey,
Module retVal{nullptr};
try {
std::ifstream in(cacheFile, std::ios::binary);
if (!in.is_open()) {
AF_TRACE("{{{:<30} : Unable to open {} for {}}}", moduleKey,
if (!in) {
AF_TRACE("{{{:<20} : Unable to open {} for {}}}", moduleKey,
cacheFile, getDeviceProp(device).name);
removeFile(cacheFile); // Remove if exists
return Module{nullptr};
Expand Down Expand Up @@ -438,23 +448,23 @@ Module loadModuleFromDisk(const int device, const string &moduleKey,

CU_CHECK(cuModuleLoadData(&modOut, cubin.data()));

AF_TRACE("{{{:<30} : loaded from {} for {} }}", moduleKey, cacheFile,
AF_TRACE("{{{:<20} : loaded from {} for {} }}", moduleKey, cacheFile,
getDeviceProp(device).name);

retVal.set(modOut);
} catch (const std::ios_base::failure &e) {
AF_TRACE("{{{:<30} : Unable to read {} for {}}}", moduleKey, cacheFile,
AF_TRACE("{{{:<20} : Unable to read {} for {}}}", moduleKey, cacheFile,
getDeviceProp(device).name);
removeFile(cacheFile);
} catch (const AfError &e) {
if (e.getError() == AF_ERR_LOAD_SYM) {
AF_TRACE(
"{{{:<30} : Corrupt binary({}) found on disk for {}, removed}}",
"{{{:<20} : Corrupt binary({}) found on disk for {}, removed}}",
moduleKey, cacheFile, getDeviceProp(device).name);
} else {
if (modOut != nullptr) { CU_CHECK(cuModuleUnload(modOut)); }
AF_TRACE(
"{{{:<30} : cuModuleLoadData failed with content from {} for "
"{{{:<20} : cuModuleLoadData failed with content from {} for "
"{}, {}}}",
moduleKey, cacheFile, getDeviceProp(device).name, e.what());
}
Expand Down
167 changes: 101 additions & 66 deletions src/backend/cuda/device_manager.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,6 +18,7 @@
#include <common/defines.hpp>
#include <common/graphics_common.hpp>
#include <common/host_memory.hpp>
#include <common/util.hpp>
#include <cublas_v2.h> // needed for af/cuda.h
#include <device_manager.hpp>
#include <driver.h>
Expand All @@ -44,10 +45,12 @@
#include <stdexcept>
#include <string>
#include <thread>
#include <utility>
#include <vector>

using std::begin;
using std::end;
using std::find;
using std::find_if;
using std::make_pair;
using std::pair;
Expand All @@ -63,21 +66,39 @@ struct cuNVRTCcompute {
int major;
/// Maximum minor compute flag supported by cudaVersion
int minor;
/// Maximum minor compute flag supported on the embedded(Jetson) platforms
int embedded_minor;
};

// clang-format off
static const int jetsonComputeCapabilities[] = {
7020,
6020,
5030,
3020,
};
// clang-format on

// clang-format off
static const cuNVRTCcompute Toolkit2MaxCompute[] = {
{10020, 7, 5},
{10010, 7, 5},
{10000, 7, 2},
{9020, 7, 2},
{9010, 7, 2},
{9000, 7, 2},
{8000, 5, 3},
{7050, 5, 3},
{7000, 5, 3}};
{10020, 7, 5, 2},
{10010, 7, 5, 2},
{10000, 7, 0, 2},
{ 9020, 7, 0, 2},
{ 9010, 7, 0, 2},
{ 9000, 7, 0, 2},
{ 8000, 5, 2, 3},
{ 7050, 5, 2, 3},
{ 7000, 5, 2, 3}};
// clang-format on

bool isEmbedded(pair<int, int> compute) {
int version = compute.first * 1000 + compute.second * 10;
return end(jetsonComputeCapabilities) !=
find(begin(jetsonComputeCapabilities),
end(jetsonComputeCapabilities), version);
}

bool checkDeviceWithRuntime(int runtime, pair<int, int> compute) {
auto rt = find_if(
begin(Toolkit2MaxCompute), end(Toolkit2MaxCompute),
Expand All @@ -88,7 +109,7 @@ bool checkDeviceWithRuntime(int runtime, pair<int, int> compute) {
"CUDA runtime version({}) not recognized. Please "
"create an issue or a pull request on the ArrayFire repository "
"to update the Toolkit2MaxCompute array with this version of "
"the CUDA Runtime. Continuing assuming everything is okay.",
"the CUDA Runtime. Continuing.",
int_version_to_string(runtime));
return true;
}
Expand All @@ -105,50 +126,66 @@ bool checkDeviceWithRuntime(int runtime, pair<int, int> compute) {
}

/// Check for compatible compute version based on runtime cuda toolkit version
void checkAndSetDevMaxCompute(pair<int, int> &prop) {
auto originalCompute = prop;
UNUSED(originalCompute);
int rtCudaVer = 0;
void checkAndSetDevMaxCompute(pair<int, int> &computeCapability) {
auto originalCompute = computeCapability;
int rtCudaVer = 0;
CUDA_CHECK(cudaRuntimeGetVersion(&rtCudaVer));
auto tkitMaxCompute = find_if(
begin(Toolkit2MaxCompute), end(Toolkit2MaxCompute),
[rtCudaVer](cuNVRTCcompute v) { return rtCudaVer == v.cudaVersion; });

bool embeddedDevice = isEmbedded(computeCapability);

// If runtime cuda version is found in toolkit array
// check for max possible compute for that cuda version
if (tkitMaxCompute != end(Toolkit2MaxCompute) &&
prop.first > tkitMaxCompute->major) {
prop = make_pair(tkitMaxCompute->major, tkitMaxCompute->minor);
#ifndef NDEBUG
char errMsg[] =
"Current device compute version (%d.%d) exceeds supported maximum "
"cuda runtime compute version (%d.%d). Using %d.%d.";
fprintf(stderr, errMsg, originalCompute.first, originalCompute.second,
prop.first, prop.second, prop.first, prop.second);
#endif
} else if (prop.first > Toolkit2MaxCompute[0].major) {
computeCapability.first >= tkitMaxCompute->major) {
int minorVersion = embeddedDevice ? tkitMaxCompute->embedded_minor
: tkitMaxCompute->minor;

if (computeCapability.second > minorVersion) {
Comment thread
umar456 marked this conversation as resolved.
computeCapability = make_pair(tkitMaxCompute->major, minorVersion);
spdlog::get("platform")
->warn(
"The compute capability for the current device({}.{}) "
"exceeds maximum supported by ArrayFire's CUDA "
"runtime({}.{}). Download or rebuild the latest version of "
"ArrayFire to avoid this warning. Using {}.{} for JIT "
"compilation kernels.",
originalCompute.first, originalCompute.second,
computeCapability.first, computeCapability.second,
computeCapability.first, computeCapability.second);
}
} else if (computeCapability.first >= Toolkit2MaxCompute[0].major) {
// If runtime cuda version is NOT found in toolkit array
// use the top most toolkit max compute
prop =
make_pair(Toolkit2MaxCompute[0].major, Toolkit2MaxCompute[0].minor);
#ifndef NDEBUG
char errMsg[] =
"Runtime cuda version not found in toolkit info array."
"Current device compute version (%d.%d) exceeds supported maximum "
"runtime cuda compute version (%d.%d) of latest known cuda toolkit."
"Using %d.%d.";
fprintf(stderr, errMsg, originalCompute.first, originalCompute.second,
prop.first, prop.second, prop.first, prop.second);
#endif
} else if (prop.first < 3) {
int minorVersion = embeddedDevice ? tkitMaxCompute->embedded_minor
: tkitMaxCompute->minor;
if (computeCapability.second > minorVersion) {
computeCapability =
make_pair(Toolkit2MaxCompute[0].major, minorVersion);
spdlog::get("platform")
->warn(
"CUDA runtime version({}) not recognized. Targeting "
"compute {}.{} for this device which is the latest compute "
"capability supported by ArrayFire's CUDA runtime({}.{}). "
"Please create an issue or a pull request on the ArrayFire "
"repository to update the Toolkit2MaxCompute array with "
"this version of the CUDA Runtime.",
int_version_to_string(rtCudaVer), originalCompute.first,
originalCompute.second, computeCapability.first,
computeCapability.second, computeCapability.first,
computeCapability.second);
}
} else if (computeCapability.first < 3) {
// all compute versions prior to Kepler, we don't support
// don't change the prop.
#ifndef NDEBUG
char errMsg[] =
"Current device compute version (%d.%d) lower than the"
"minimum compute version ArrayFire supports.";
fprintf(stderr, errMsg, originalCompute.first, originalCompute.second);
#endif
// don't change the computeCapability.
spdlog::get("platform")
->warn(
"The compute capability of the current device({}.{}) "
"lower than the minimum compute version ArrayFire "
"supports.",
originalCompute.first, originalCompute.second);
}
}

Expand Down Expand Up @@ -370,7 +407,6 @@ static const ToolkitDriverVersions
/// \note: only works in debug builds
void debugRuntimeCheck(spdlog::logger *logger, int runtime_version,
int driver_version) {
#ifndef NDEBUG
auto runtime_it =
find_if(begin(CudaToDriverVersion), end(CudaToDriverVersion),
[runtime_version](ToolkitDriverVersions ver) {
Expand All @@ -388,31 +424,28 @@ void debugRuntimeCheck(spdlog::logger *logger, int runtime_version,
// display a message in the trace. Do not throw an error unless this is
// a debug build
if (runtime_it == end(CudaToDriverVersion)) {
char buf[1024];
char buf[256];
char err_msg[] =
"CUDA runtime version(%s) not recognized. Please "
"create an issue or a pull request on the ArrayFire repository to "
"update the CudaToDriverVersion variable with this version of "
"the CUDA Toolkit.\n";
snprintf(buf, 1024, err_msg,
"CUDA runtime version(%s) not recognized. Please create an issue "
"or a pull request on the ArrayFire repository to update the "
"CudaToDriverVersion variable with this version of the CUDA "
"runtime.\n";
snprintf(buf, 256, err_msg,
int_version_to_string(runtime_version).c_str());
AF_TRACE("{}", buf);
#ifndef NDEBUG
AF_ERROR(buf, AF_ERR_RUNTIME);
#endif
}

if (driver_it == end(CudaToDriverVersion)) {
char buf[1024];
char err_msg[] =
"CUDA driver version(%s) not part of the "
"CudaToDriverVersion array. Please create an issue or a pull "
"request on the ArrayFire repository to update the "
"CudaToDriverVersion variable with this version of the CUDA "
"Toolkit.\n";
snprintf(buf, 1024, err_msg,
int_version_to_string(driver_version).c_str());
AF_TRACE("{}", buf);
AF_TRACE(
"CUDA driver version({}) not part of the CudaToDriverVersion "
"array. Please create an issue or a pull request on the ArrayFire "
"repository to update the CudaToDriverVersion variable with this "
"version of the CUDA runtime.\n",
int_version_to_string(driver_version).c_str());
}
#endif
}

// Check if the device driver version is recent enough to run the cuda libs
Expand Down Expand Up @@ -515,11 +548,13 @@ DeviceManager::DeviceManager()
compute2cores(dev.prop.major, dev.prop.minor) *
dev.prop.clockRate;
dev.nativeId = i;
AF_TRACE("Found device: {} ({:0.3} GB | ~{} GFLOPs | {} SMs)",
dev.prop.name,
dev.prop.totalGlobalMem / 1024. / 1024. / 1024.,
dev.flops / 1024. / 1024. * 2,
dev.prop.multiProcessorCount);
AF_TRACE(
"Found device: {} (sm_{}{}) ({:0.3} GB | ~{} GFLOPs | {} "
"SMs)",
dev.prop.name, dev.prop.major, dev.prop.minor,
dev.prop.totalGlobalMem / 1024. / 1024. / 1024.,
dev.flops / 1024. / 1024. * 2,
dev.prop.multiProcessorCount);
cuDevices.push_back(dev);
}
}
Expand Down
4 changes: 1 addition & 3 deletions src/backend/cuda/kernel/random_engine.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -134,12 +134,10 @@ __device__ static double getDouble01(uint num1, uint num2) {
uint64_t n2 = num2;
n1 <<= 32;
uint64_t num = n1 | n2;
#pragma diag_suppress 3245
constexpr double factor =
((1.0) / (std::numeric_limits<unsigned long long>::max() +
static_cast<long double>(1.0l)));
static_cast<double>(1.0)));
constexpr double half_factor((0.5) * factor);
#pragma diag_default 3245

return fma(static_cast<double>(num), factor, half_factor);
}
Expand Down
Loading