Created
March 30, 2023 13:54
-
-
Save jszuppe/79645a85a442cc0c46556d1aadc5c782 to your computer and use it in GitHub Desktop.
libcucxx nvrtc nvc++: floating constant is out of range
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
#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(¤t_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