NVRTC with ptxas-options not working

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.

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 - cuda-samples/Samples/vectorAdd_nvrtc at master · NVIDIA/cuda-samples · GitHub

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.