Hello Guys,
I could get required information for offline compiled kernels by passing --ptxas-options=-v to nvcc. However, I am unable to achieve the same for runtime compilation using nvrtc. I can see that NVRTC documentation does mention the ptxas-options flag but it doesn’t seem to work. I was abel to reproduce this behavior in standalone cuda-sample vectorAdd_nvrtc.
The error simply says option unrecognized.
cp vectorAdd_nvrtc ../../bin/x86_64/linux/release
> Using CUDA Device [0]: NVIDIA GeForce GTX 1060 3GB
compilation log ---
nvrtc: error: unrecognized option --ptxas-options=-v found
end log ---
error: nvrtcCompileProgram failed with error NVRTC_ERROR_INVALID_OPTION
I have tried the short option also, I am clueless why only this option is not working even though it is in documentation.
Any help/hints greatly appreciated.
Thank you,
Pradeep.
1 Like
The format of the option is:
--ptxas-options <options>
not:
--ptxas-options=<options>
Otherwise, for a question like this, I suggest you provide a complete code example of what you are doing.
My bad, I forgot to mention what I have already tried.
I have tried all the below combinations of the passing the option, none of them work.
--ptxas-options --verbose--ptxas-options -v--ptxas-options=--verbose--ptxas-options=-v
Given below is the text in NVRTC online documentation,
NVRTC supports the compile options below. Option names with two preceding dashs (--) are long option names and option names with one preceding dash (-) are short option names. Short option names can be used instead of long option names. When a compile option takes an argument, an assignment operator (=) is used to separate the compile option argument from the compile option name, e.g., "--gpu-architecture=compute_60". Alternatively, the compile option name and the argument can be specified in separate strings without an assignment operator, .e.g, "--gpu-architecture""compute_60". Single-character short option names, such as -D, -U, and -I, do not require an assignment operator, and the compile option name and the argument can be present in the same string with or without spaces between them. For instance, "-D=<def>", "-D<def>", and "-D <def>" are all supported.
That gave me the impression that = should also work. Also, maxrregcount works fine with or without =. So, it doesn’t seem like that is the issue here.
Actually, I already did share(in the original description) the sample I used to check this - https://github.com/NVIDIA/cuda-samples/tree/master/Samples/vectorAdd_nvrtc
You can apply the following diff to the above CUDA example to try reproduce the problem on your end.
nvrtc.diff (1.4 KB)
Please disregard the previous suggestion to file a bug (which I have deleted.)
The arguments must be space-separated, = separation will not work.
And the --ptxas-options and -v arguments must be delivered via two separate compile options. Example:
$ cat t2.cpp
#include <sstream>
#include <fstream>
#include <string>
#include <cstring>
#include <iostream>
#include <nvrtc.h>
#include <cuda_runtime_api.h>
#include <cuda.h>
#define NVRTC_SAFE_CALL(Name, x) \
do { \
nvrtcResult result = x; \
if (result != NVRTC_SUCCESS) { \
std::cerr << "\nerror: " << Name << " failed with error " \
<< nvrtcGetErrorString(result) << std::endl; \
exit(1); \
} \
} while (0)
std::string tmpcode = "__global__ void k(int *d){ \n\
*d = 1; } \n\
";
int main(){
// Create nvrtcProgram
nvrtcProgram prog;
std::string my_code = tmpcode;
NVRTC_SAFE_CALL("create", nvrtcCreateProgram(&prog, my_code.c_str(), "k.cpp", 0, NULL, NULL));
const char *compileOptions[] = {"--gpu-architecture=sm_70", "--ptxas-options", " -v"};
size_t nOptions = NOPT;
auto res = nvrtcCompileProgram (prog, nOptions, compileOptions);
// dump log
size_t logSize;
NVRTC_SAFE_CALL("nvrtcGetProgramLogSize",
nvrtcGetProgramLogSize(prog, &logSize));
char *log = reinterpret_cast<char *>(malloc(sizeof(char) * logSize + 1));
NVRTC_SAFE_CALL("nvrtcGetProgramLog", nvrtcGetProgramLog(prog, log));
log[logSize] = '\x0';
if (strlen(log) >= 2) {
std::cerr << "\n compilation log ---\n";
std::cerr << log;
std::cerr << "\n end log ---\n";
}
free(log);
NVRTC_SAFE_CALL("compile", res);
}
$ nvcc -o t2 t2.cpp -lnvrtc -DNOPT=3
$ ./t2
compilation log ---
ptxas info : 0 bytes gmem
ptxas info : Compiling entry function '_Z1kPi' for 'sm_70'
ptxas info : Function properties for _Z1kPi
ptxas . 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 8 registers, 360 bytes cmem[0]
end log ---
$
I have been able to get usage info from your example, and from my edited cuda sample based on your suggestion. However, I didn’t need the space before -vand in fact it didn’t work if I added that space.
Nevertheless, isn’t this a bug given what is written NVRTC compile options documentation. Other compile options work fine with/without = but --ptxas-options requires special syntax. It seems like a bug to me, either in code or at the least documentation should reflect that --ptxas-options doesn’t work with = syntax.
Note: I have already raised a bug before your reply. Given that it doesn’t work as per documentation, I think CUDA dev team should still fix either the docs or the options behavior. I will add a comment in the bug that if the options are provided separately, they work fine.
sure. You might wish to comment in the bug that you suggest the documentation be fixed.