Error when evaluating vprintf

[CarloRamponi]

Hello,
I've noticed that cuasm does not work if I use the on-device printf function.

Here is how to reproduce the issue:

a.cu:

__device__ int foo(int a, int b) {
	int c = a + b;
        printf("%d\n", c);
	return c;
}

__global__ void kernel(int *a) {
	int index = threadIdx.x;
	a[index] = foo(a[index], a[index]);
}

Commands:

nvcc -arch=sm_80 a.cu -o a.cubin --cubin
cuasm a.cubin -o a.cuasm
cuasm a.cuasm -o a_assembled.cubin

Output (of the last command):

2024-09-02 16:36:07,646 -          - Running CuAsmParser.parse...
2024-09-02 16:36:07,646 -    ENTRY -     Parsing file a.cuasm
2024-09-02 16:36:07,646 -          -     Running CuAsmParser.__preScan...
2024-09-02 16:36:07,974 -          -     Running CuAsmParser.__gatherTextSectionSizeLabel...
2024-09-02 16:36:07,974 -          -     Running CuAsmParser.__buildInternalTables...
2024-09-02 16:36:07,975 -          -     Running CuAsmParser.__evalFixups...
2024-09-02 16:36:07,975 -    ERROR - Assertion failed in:
    File a.cuasm:777 :

    Error when evaluating fixup @line614: expr=vprintf, msg=Unknown expression vprintf
Traceback (most recent call last):
  File "/home/carlo/Projects/GPU/CuAssembler/CuAsm/CuAsmParser.py", line 1118, in __evalFixups
    val, vs = self.__evalExpr(expr)
              ^^^^^^^^^^^^^^^^^^^^^
  File "/home/carlo/Projects/GPU/CuAssembler/CuAsm/CuAsmParser.py", line 1580, in __evalExpr
    aval, a_issym = self.__evalVar(a)
                    ^^^^^^^^^^^^^^^^^
  File "/home/carlo/Projects/GPU/CuAssembler/CuAsm/CuAsmParser.py", line 1542, in __evalVar
    raise Exception('Unknown expression %s'%var)
Exception: Unknown expression vprintf

During handling of the above exception, another exception occurred:

Traceback (most recent call last):
  File "/home/carlo/Projects/GPU/CuAssembler/bin/cuasm", line 153, in <module>
    doProcess(infile, outfile, direction)
  File "/home/carlo/Projects/GPU/CuAssembler/bin/cuasm", line 94, in doProcess
    cuasm2cubin(src, dst)
  File "/home/carlo/Projects/GPU/CuAssembler/bin/cuasm", line 72, in cuasm2cubin
    cap.parse(asmname)
  File "/home/carlo/Projects/GPU/CuAssembler/CuAsm/CuAsmLogger.py", line 210, in wrapper
    ret = func(*args, **kwargs)
          ^^^^^^^^^^^^^^^^^^^^^
  File "/home/carlo/Projects/GPU/CuAssembler/CuAsm/CuAsmParser.py", line 766, in parse
    self.__evalFixups() #
    ^^^^^^^^^^^^^^^^^^^
  File "/home/carlo/Projects/GPU/CuAssembler/CuAsm/CuAsmLogger.py", line 241, in wrapper
    ret = func(*args, **kwargs)
          ^^^^^^^^^^^^^^^^^^^^^
  File "/home/carlo/Projects/GPU/CuAssembler/CuAsm/CuAsmParser.py", line 1137, in __evalFixups
    self.__assert(False, 'Error when evaluating fixup @line%d: expr=%s, msg=%s'
  File "/home/carlo/Projects/GPU/CuAssembler/CuAsm/CuAsmParser.py", line 1490, in __assert
    raise Exception(full_msg)
Exception: Assertion failed in:
    File a.cuasm:777 :

    Error when evaluating fixup @line614: expr=vprintf, msg=Unknown expression vprintf

If I remove the printf call everything works fine.

[cloudcores]

Sorry, supports for high level features(such as printf/debug/device runtime/dynamic parallelism, etc.) are not guaranteed. Usually their interfaces may subject to change from version to version, thus not that easy to follow.