diff --git a/src/backend/common/half.hpp b/src/backend/common/half.hpp index 885664798e..ce06eedf02 100644 --- a/src/backend/common/half.hpp +++ b/src/backend/common/half.hpp @@ -832,9 +832,7 @@ class alignas(2) half { #endif public: -#if CUDA_VERSION >= 10000 AF_CONSTEXPR -#endif half() = default; /// Constructor. diff --git a/src/backend/cuda/compile_module.cpp b/src/backend/cuda/compile_module.cpp index 06a96e1f29..c4c3315d0a 100644 --- a/src/backend/cuda/compile_module.cpp +++ b/src/backend/cuda/compile_module.cpp @@ -143,20 +143,29 @@ Module compileModule(const string &moduleKey, const vector &sources, const vector &kInstances, const bool sourceIsJIT) { nvrtcProgram prog; if (sourceIsJIT) { - array headers = { + constexpr const char *header_names[] = { + "utility", + "cuda_fp16.hpp", + "cuda_fp16.h", + }; + constexpr size_t numHeaders = extent::value; + array headers = { + "", cuda_fp16_hpp, cuda_fp16_h, }; - array header_names = {"cuda_fp16.hpp", "cuda_fp16.h"}; + static_assert(headers.size() == numHeaders, + "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", @@ -183,12 +192,13 @@ Module compileModule(const string &moduleKey, const vector &sources, "minmax_op.hpp", }; - constexpr size_t NumHeaders = extent::value; - static const array sourceStrings = {{ + constexpr size_t numHeaders = extent::value; + static const array 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), @@ -230,11 +240,11 @@ Module compileModule(const string &moduleKey, const vector &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::value == NumHeaders, + sourceStrings[28].c_str()}; + static_assert(extent::value == numHeaders, "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)); } @@ -246,6 +256,7 @@ Module compileModule(const string &moduleKey, const vector &sources, vector compiler_options = { arch.data(), "--std=c++14", + "--device-as-default-execution-space", #if !(defined(NDEBUG) || defined(__aarch64__) || defined(__LP64__)) "--device-debug", "--generate-line-info" @@ -256,7 +267,6 @@ Module compileModule(const string &moduleKey, const vector &sources, back_insert_iterator>(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())); } @@ -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}; @@ -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()); } diff --git a/src/backend/cuda/device_manager.cpp b/src/backend/cuda/device_manager.cpp index 1493e5e432..9ec832fe59 100644 --- a/src/backend/cuda/device_manager.cpp +++ b/src/backend/cuda/device_manager.cpp @@ -18,6 +18,7 @@ #include #include #include +#include #include // needed for af/cuda.h #include #include @@ -44,10 +45,12 @@ #include #include #include +#include #include using std::begin; using std::end; +using std::find; using std::find_if; using std::make_pair; using std::pair; @@ -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 compute) { + int version = compute.first * 1000 + compute.second * 10; + return end(jetsonComputeCapabilities) != + find(begin(jetsonComputeCapabilities), + end(jetsonComputeCapabilities), version); +} + bool checkDeviceWithRuntime(int runtime, pair compute) { auto rt = find_if( begin(Toolkit2MaxCompute), end(Toolkit2MaxCompute), @@ -88,7 +109,7 @@ bool checkDeviceWithRuntime(int runtime, pair 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; } @@ -105,50 +126,66 @@ bool checkDeviceWithRuntime(int runtime, pair compute) { } /// Check for compatible compute version based on runtime cuda toolkit version -void checkAndSetDevMaxCompute(pair &prop) { - auto originalCompute = prop; - UNUSED(originalCompute); - int rtCudaVer = 0; +void checkAndSetDevMaxCompute(pair &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) { + 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); } } @@ -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) { @@ -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 @@ -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); } } diff --git a/src/backend/cuda/kernel/random_engine.hpp b/src/backend/cuda/kernel/random_engine.hpp index 0ef218ad93..e52e78d354 100644 --- a/src/backend/cuda/kernel/random_engine.hpp +++ b/src/backend/cuda/kernel/random_engine.hpp @@ -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::max() + - static_cast(1.0l))); + static_cast(1.0))); constexpr double half_factor((0.5) * factor); -#pragma diag_default 3245 return fma(static_cast(num), factor, half_factor); } diff --git a/src/backend/opencl/compile_module.cpp b/src/backend/opencl/compile_module.cpp index 2f6d374db1..35f992fe02 100644 --- a/src/backend/opencl/compile_module.cpp +++ b/src/backend/opencl/compile_module.cpp @@ -194,12 +194,12 @@ Module compileModule(const string &moduleKey, const vector &sources, // before the current thread. if (!renameFile(tempFile, cacheFile)) { removeFile(tempFile); } } catch (const cl::Error &e) { - AF_TRACE("{{{:<30} : Failed to fetch opencl binary for {}, {}}}", + AF_TRACE("{{{:<20} : Failed to fetch opencl binary for {}, {}}}", moduleKey, opencl::getDevice(device).getInfo(), e.what()); } catch (const std::ios_base::failure &e) { - AF_TRACE("{{{:<30} : Failed writing binary to {} for {}, {}}}", + AF_TRACE("{{{:<20} : Failed writing binary to {} for {}, {}}}", moduleKey, cacheFile, opencl::getDevice(device).getInfo(), e.what()); @@ -207,7 +207,7 @@ Module compileModule(const string &moduleKey, const vector &sources, } #endif - AF_TRACE("{{{:<30} : {{ compile:{:>5} ms, {{ {} }}, {} }}}}", moduleKey, + AF_TRACE("{{{:<20} : {{ compile:{:>5} ms, {{ {} }}, {} }}}}", moduleKey, duration_cast(compileEnd - compileBegin).count(), fmt::join(options, " "), getDevice(getActiveDeviceId()).getInfo()); @@ -250,26 +250,26 @@ Module loadModuleFromDisk(const int device, const string &moduleKey, program = Program(opencl::getContext(), {dev}, {clbin}); program.build(); - AF_TRACE("{{{:<30} : loaded from {} for {} }}", moduleKey, cacheFile, + AF_TRACE("{{{:<20} : loaded from {} for {} }}", moduleKey, cacheFile, dev.getInfo()); retVal.set(program); } 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, dev.getInfo()); } else { - AF_TRACE("{{{:<30} : Unable to open {} for {}}}", moduleKey, + AF_TRACE("{{{:<20} : Unable to open {} for {}}}", moduleKey, cacheFile, dev.getInfo()); } removeFile(cacheFile); } catch (const std::ios_base::failure &e) { - AF_TRACE("{{{:<30} : IO failure while loading {} for {}; {}}}", + AF_TRACE("{{{:<20} : IO failure while loading {} for {}; {}}}", moduleKey, cacheFile, dev.getInfo(), e.what()); removeFile(cacheFile); } catch (const cl::Error &e) { AF_TRACE( - "{{{:<30} : Loading OpenCL binary({}) failed for {}; {}, Build " + "{{{:<20} : Loading OpenCL binary({}) failed for {}; {}, Build " "Log: {}}}", moduleKey, cacheFile, dev.getInfo(), e.what(), getProgramBuildLog(program));