Skip to content

Instantly share code, notes, and snippets.

@jszuppe
Created March 30, 2023 13:54
Show Gist options
  • Star 0 You must be signed in to star a gist
  • Fork 0 You must be signed in to fork a gist
  • Save jszuppe/79645a85a442cc0c46556d1aadc5c782 to your computer and use it in GitHub Desktop.
Save jszuppe/79645a85a442cc0c46556d1aadc5c782 to your computer and use it in GitHub Desktop.
libcucxx nvrtc nvc++: floating constant is out of range
#include <vector>
#include <iostream>
#include <string>
#include <algorithm>
#include <memory>
#include <random>
#include <nvrtc.h>
#include <cuda.h>
#include <cuda_runtime_api.h>
#include <cuComplex.h>
#include <cstdlib>
#define CUDA_CHECK_AND_EXIT(x) x
#define NVRTC_SAFE_CALL(x) \
do { \
nvrtcResult result = x; \
if (result != NVRTC_SUCCESS) { \
std::cerr << "\nerror: " #x " failed with error " << nvrtcGetErrorString(result) << '\n'; \
exit(1); \
} \
} while (0)
#ifndef CU_CHECK_AND_EXIT
# define CU_CHECK_AND_EXIT(error) \
{ \
auto status = static_cast<CUresult>(error); \
if (status != CUDA_SUCCESS) { \
const char* pstr; \
cuGetErrorString(status, &pstr); \
std::cout << pstr << " " << __FILE__ << ":" << __LINE__ << std::endl; \
std::exit(status); \
} \
}
#endif // CU_CHECK_AND_EXIT
template<class Type>
inline Type get_global_from_module(CUmodule module, const char* name) {
CUdeviceptr value_ptr;
size_t value_size;
CU_CHECK_AND_EXIT(cuModuleGetGlobal(&value_ptr, &value_size, module, name));
Type value_host;
CU_CHECK_AND_EXIT(cuMemcpyDtoH(&value_host, value_ptr, value_size));
return value_host;
}
inline unsigned get_device_architecture(int device) {
int major = 0;
int minor = 0;
CUDA_CHECK_AND_EXIT(cudaDeviceGetAttribute(&major, cudaDevAttrComputeCapabilityMajor, device));
CUDA_CHECK_AND_EXIT(cudaDeviceGetAttribute(&minor, cudaDevAttrComputeCapabilityMinor, device));
return major * 10 + minor;
}
inline std::string get_device_architecture_option(int device) {
std::string gpu_architecture_option =
"--gpu-architecture=compute_" + std::to_string(get_device_architecture(device));
return gpu_architecture_option;
}
inline void print_program_log(const nvrtcProgram prog) {
size_t log_size;
NVRTC_SAFE_CALL(nvrtcGetProgramLogSize(prog, &log_size));
char* log = new char[log_size];
NVRTC_SAFE_CALL(nvrtcGetProgramLog(prog, log));
std::cout << log << '\n';
delete[] log;
}
const char* test_kernel = R"kernel(
#include <cuda/std/climits>
constexpr double a = __DBL_DENORM_MIN__;
using T = decltype(a);
)kernel";
// Problem:
// nvc++ -I/usr/local/cuda/include -DCUDA_INCLUDE_DIR=\"/usr/local/cuda/include\" -L/usr/local/cuda/lib64/ bug.cpp -lcuda -lcudart -lnvrtc
// OK:
// g++ -I/usr/local/cuda/include -DCUDA_INCLUDE_DIR=\"/usr/local/cuda/include\" -L/usr/local/cuda/lib64/ bug.cpp -lcuda -lcudart -lnvrtc
int main(int, char**) {
// Complex type according to precision
// using value_type = cuDoubleComplex; // or double2, or cuda::std::complex<double>
using value_type = double;
nvrtcProgram program;
NVRTC_SAFE_CALL(nvrtcCreateProgram(&program, // program
test_kernel, // buffer
"test_kernel.cu", // name
0, // numHeaders
NULL, // headers
NULL)); // includeNames
// Get current device
int current_device;
CUDA_CHECK_AND_EXIT(cudaGetDevice(&current_device));
// Prepare compilation options
std::vector<const char*> opts = {
"--std=c++17",
"--device-as-default-execution-space",
"--include-path=" CUDA_INCLUDE_DIR // Add path to CUDA include directory
};
// Add gpu-architecture to opts
std::string gpu_architecture_option = get_device_architecture_option(current_device);
opts.push_back(gpu_architecture_option.c_str());
nvrtcResult compileResult = nvrtcCompileProgram(program, // program
static_cast<int>(opts.size()), // numOptions
opts.data()); // options
// Obtain compilation log from the program
if (compileResult != NVRTC_SUCCESS) {
for (auto o : opts) {
std::cout << o << std::endl;
}
print_program_log(program);
std::exit(1);
}
// Destroy the program.
NVRTC_SAFE_CALL(nvrtcDestroyProgram(&program));
std::cout << "Success" << std::endl;
return 0;
}
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment