Error when compiling kernel (using the example python script)

Jebsta commented 1 year ago

I installed deepdrr and all dependencies. Ran the example python script with the path changed to a CT Scan on mine. Get following Error:

<deepdrr.vol.volume.Volume object at 0x7f18a5c78340>
Traceback (most recent call last):
  File "/home/user/deepdrr-test/app/proj_surgery-simulator/generate_xray.py", line 8, in <module>
    with Projector(ct, carm=carm) as projector:
  File "/home/user/deepdrr-test/.venv/lib/python3.10/site-packages/deepdrr/projector/projector.py", line 302, in __init__
    self.mod = _get_kernel_projector_module(
  File "/home/user/deepdrr-test/.venv/lib/python3.10/site-packages/deepdrr/projector/projector.py", line 134, in _get_kernel_projector_module
    return SourceModule(
  File "/home/user/deepdrr-test/.venv/lib/python3.10/site-packages/pycuda/compiler.py", line 355, in __init__
    cubin = compile(
  File "/home/user/deepdrr-test/.venv/lib/python3.10/site-packages/pycuda/compiler.py", line 304, in compile
    return compile_plain(source, options, keep, nvcc, cache_dir, target)
  File "/home/user/deepdrr-test/.venv/lib/python3.10/site-packages/pycuda/compiler.py", line 154, in compile_plain
    raise CompileError(
pycuda.driver.CompileError: nvcc compilation of /tmp/tmp9_1owmrm/kernel.cu failed
[command: nvcc --cubin -D NUM_VOLUMES=1 -D NUM_MATERIALS=3 -D ATTENUATE_OUTSIDE_VOLUME=0 -D AIR_INDEX=0 -arch sm_89 -I/home/user/deepdrr-test/.venv/lib/python3.10/site-packages/deepdrr/projector/cubic -I/home/user/deepdrr-test/.venv/lib/python3.10/site-packages/deepdrr/projector -I/home/user/deepdrr-test/.venv/lib/python3.10/site-packages/pycuda/cuda kernel.cu]
[stderr:
/home/user/deepdrr-test/.venv/lib/python3.10/site-packages/deepdrr/projector/cubic/cubicTex3D.cu(54): error: identifier "texture" is undefined
  __attribute__((device)) float linearTex3D(texture<T, 3, mode> tex, float3 coord)
                                            ^

/home/user/deepdrr-test/.venv/lib/python3.10/site-packages/deepdrr/projector/cubic/cubicTex3D.cu(54): error: type name is not allowed
  __attribute__((device)) float linearTex3D(texture<T, 3, mode> tex, float3 coord)
                                                    ^

/home/user/deepdrr-test/.venv/lib/python3.10/site-packages/deepdrr/projector/cubic/cubicTex3D.cu(54): error: identifier "tex" is undefined
  __attribute__((device)) float linearTex3D(texture<T, 3, mode> tex, float3 coord)
                                                                ^

/home/user/deepdrr-test/.venv/lib/python3.10/site-packages/deepdrr/projector/cubic/cubicTex3D.cu(54): error: type name is not allowed
  __attribute__((device)) float linearTex3D(texture<T, 3, mode> tex, float3 coord)
                                                                     ^

/home/user/deepdrr-test/.venv/lib/python3.10/site-packages/deepdrr/projector/cubic/cubicTex3D.cu(54): error: too many initializer values
  __attribute__((device)) float linearTex3D(texture<T, 3, mode> tex, float3 coord)
                                                       ^

/home/user/deepdrr-test/.venv/lib/python3.10/site-packages/deepdrr/projector/cubic/cubicTex3D.cu(54): error: expected a ")"
  __attribute__((device)) float linearTex3D(texture<T, 3, mode> tex, float3 coord)
                                                                            ^

/home/user/deepdrr-test/.venv/lib/python3.10/site-packages/deepdrr/projector/cubic/cubicTex3D.cu(55): error: expected a ";"
  {
  ^

kernel.cu(2358): error: identifier "cubicTex3D" is undefined
      do { if (do_trace[0]) { do { px[0] = sx_ijk[0] + alpha * rx_ijk[0] - 0.5; py[0] = sy_ijk[0] + alpha * ry_ijk[0] - 0.5; pz[0] = sz_ijk[0] + alpha * rz_ijk[0] - 0.5; } while (0); do { do { seg_at_alpha[0][0] = round( cubicTex3D(seg_0_0, px[0], py[0], pz[0])); } while (0); do { seg_at_alpha[0][1] = round( cubicTex3D(seg_0_1, px[0], py[0], pz[0])); } while (0); do { seg_at_alpha[0][2] = round( cubicTex3D(seg_0_2, px[0], py[0], pz[0])); } while (0); } while (0); } } while (0);
                                                                                                                                                                                                                             ^

kernel.cu(2358): error: identifier "cubicTex3D" is undefined
      do { if (do_trace[0]) { do { px[0] = sx_ijk[0] + alpha * rx_ijk[0] - 0.5; py[0] = sy_ijk[0] + alpha * ry_ijk[0] - 0.5; pz[0] = sz_ijk[0] + alpha * rz_ijk[0] - 0.5; } while (0); do { do { seg_at_alpha[0][0] = round( cubicTex3D(seg_0_0, px[0], py[0], pz[0])); } while (0); do { seg_at_alpha[0][1] = round( cubicTex3D(seg_0_1, px[0], py[0], pz[0])); } while (0); do { seg_at_alpha[0][2] = round( cubicTex3D(seg_0_2, px[0], py[0], pz[0])); } while (0); } while (0); } } while (0);
                                                                                                                                                                                                                                                                                                                      ^

kernel.cu(2358): error: identifier "cubicTex3D" is undefined
      do { if (do_trace[0]) { do { px[0] = sx_ijk[0] + alpha * rx_ijk[0] - 0.5; py[0] = sy_ijk[0] + alpha * ry_ijk[0] - 0.5; pz[0] = sz_ijk[0] + alpha * rz_ijk[0] - 0.5; } while (0); do { do { seg_at_alpha[0][0] = round( cubicTex3D(seg_0_0, px[0], py[0], pz[0])); } while (0); do { seg_at_alpha[0][1] = round( cubicTex3D(seg_0_1, px[0], py[0], pz[0])); } while (0); do { seg_at_alpha[0][2] = round( cubicTex3D(seg_0_2, px[0], py[0], pz[0])); } while (0); } while (0); } } while (0);
                                                                                                                                                                                                                                                                                                                                                                                                               ^

kernel.cu(2409): error: identifier "volume_0" is undefined
        do { if (do_trace[0] && (priority[0] == curr_priority)) { do { do { area_density[(0)] += (weight)*tex3D(volume_0, px[0], py[0], pz[0]) * seg_at_alpha[0][0]; } while (0); do { area_density[(1)] += (weight)*tex3D(volume_0, px[0], py[0], pz[0]) * seg_at_alpha[0][1]; } while (0); do { area_density[(2)] += (weight)*tex3D(volume_0, px[0], py[0], pz[0]) * seg_at_alpha[0][2]; } while (0); } while (0); } } while (0);
                                                                                                                ^

kernel.cu(2409): error: no instance of overloaded function "tex3D" matches the argument list
            argument types are: (<error-type>, float, float, float)
        do { if (do_trace[0] && (priority[0] == curr_priority)) { do { do { area_density[(0)] += (weight)*tex3D(volume_0, px[0], py[0], pz[0]) * seg_at_alpha[0][0]; } while (0); do { area_density[(1)] += (weight)*tex3D(volume_0, px[0], py[0], pz[0]) * seg_at_alpha[0][1]; } while (0); do { area_density[(2)] += (weight)*tex3D(volume_0, px[0], py[0], pz[0]) * seg_at_alpha[0][2]; } while (0); } while (0); } } while (0);
                                                                                                          ^

kernel.cu(2409): error: identifier "volume_0" is undefined
        do { if (do_trace[0] && (priority[0] == curr_priority)) { do { do { area_density[(0)] += (weight)*tex3D(volume_0, px[0], py[0], pz[0]) * seg_at_alpha[0][0]; } while (0); do { area_density[(1)] += (weight)*tex3D(volume_0, px[0], py[0], pz[0]) * seg_at_alpha[0][1]; } while (0); do { area_density[(2)] += (weight)*tex3D(volume_0, px[0], py[0], pz[0]) * seg_at_alpha[0][2]; } while (0); } while (0); } } while (0);
                                                                                                                                                                                                                           ^

kernel.cu(2409): error: no instance of overloaded function "tex3D" matches the argument list
            argument types are: (<error-type>, float, float, float)
        do { if (do_trace[0] && (priority[0] == curr_priority)) { do { do { area_density[(0)] += (weight)*tex3D(volume_0, px[0], py[0], pz[0]) * seg_at_alpha[0][0]; } while (0); do { area_density[(1)] += (weight)*tex3D(volume_0, px[0], py[0], pz[0]) * seg_at_alpha[0][1]; } while (0); do { area_density[(2)] += (weight)*tex3D(volume_0, px[0], py[0], pz[0]) * seg_at_alpha[0][2]; } while (0); } while (0); } } while (0);
                                                                                                                                                                                                                     ^

kernel.cu(2409): error: identifier "volume_0" is undefined
        do { if (do_trace[0] && (priority[0] == curr_priority)) { do { do { area_density[(0)] += (weight)*tex3D(volume_0, px[0], py[0], pz[0]) * seg_at_alpha[0][0]; } while (0); do { area_density[(1)] += (weight)*tex3D(volume_0, px[0], py[0], pz[0]) * seg_at_alpha[0][1]; } while (0); do { area_density[(2)] += (weight)*tex3D(volume_0, px[0], py[0], pz[0]) * seg_at_alpha[0][2]; } while (0); } while (0); } } while (0);
                                                                                                                                                                                                                                                                                                                                      ^

kernel.cu(2409): error: no instance of overloaded function "tex3D" matches the argument list
            argument types are: (<error-type>, float, float, float)
        do { if (do_trace[0] && (priority[0] == curr_priority)) { do { do { area_density[(0)] += (weight)*tex3D(volume_0, px[0], py[0], pz[0]) * seg_at_alpha[0][0]; } while (0); do { area_density[(1)] += (weight)*tex3D(volume_0, px[0], py[0], pz[0]) * seg_at_alpha[0][1]; } while (0); do { area_density[(2)] += (weight)*tex3D(volume_0, px[0], py[0], pz[0]) * seg_at_alpha[0][2]; } while (0); } while (0); } } while (0);
                                                                                                                                                                                                                                                                                                                                ^

kernel.cu(2786): error: identifier "volume_0" is undefined
              do { density_sample[0] = tex3D(volume_0, inp_x, inp_y, inp_z); mat_sample[0][0] = cubicTex3D(seg_0_0, inp_x, inp_y, inp_z); mat_sample[0][1] = cubicTex3D(seg_0_1, inp_x, inp_y, inp_z); mat_sample[0][2] = cubicTex3D(seg_0_2, inp_x, inp_y, inp_z); } while (0);
                                             ^

kernel.cu(2786): error: no instance of overloaded function "tex3D" matches the argument list
            argument types are: (<error-type>, float, float, float)
              do { density_sample[0] = tex3D(volume_0, inp_x, inp_y, inp_z); mat_sample[0][0] = cubicTex3D(seg_0_0, inp_x, inp_y, inp_z); mat_sample[0][1] = cubicTex3D(seg_0_1, inp_x, inp_y, inp_z); mat_sample[0][2] = cubicTex3D(seg_0_2, inp_x, inp_y, inp_z); } while (0);
                                       ^

kernel.cu(2786): error: identifier "cubicTex3D" is undefined
              do { density_sample[0] = tex3D(volume_0, inp_x, inp_y, inp_z); mat_sample[0][0] = cubicTex3D(seg_0_0, inp_x, inp_y, inp_z); mat_sample[0][1] = cubicTex3D(seg_0_1, inp_x, inp_y, inp_z); mat_sample[0][2] = cubicTex3D(seg_0_2, inp_x, inp_y, inp_z); } while (0);
                                                                                                ^

19 errors detected in the compilation of "kernel.cu".
]

Following Driver and Cuda Version

Thu Jul 13 11:36:28 2023       
+---------------------------------------------------------------------------------------+
| NVIDIA-SMI 530.30.02              Driver Version: 530.30.02    CUDA Version: 12.1     |
|-----------------------------------------+----------------------+----------------------+
| GPU  Name                  Persistence-M| Bus-Id        Disp.A | Volatile Uncorr. ECC |
| Fan  Temp  Perf            Pwr:Usage/Cap|         Memory-Usage | GPU-Util  Compute M. |
|                                         |                      |               MIG M. |
|=========================================+======================+======================|
|   0  NVIDIA GeForce RTX 4090         On | 00000000:10:00.0 Off |                  Off |
|  0%   42C    P8               29W / 450W|     27MiB / 24564MiB |      0%      Default |
|                                         |                      |                  N/A |
+-----------------------------------------+----------------------+----------------------+

henrykrumb commented 1 year ago

The error stems from the "texture" type being deprecated since CUDA 5.0 1. Support was entirely dropped in version 12.0.

deepdrr/projector/cubic/cubicTex3D.cu(54): error: identifier "texture" is undefined __attribute__((device)) float linearTex3D(texture<T, 3, mode> tex, float3 coord)

This issue is quite nasty to refactor, as the API significantly changed. Grepping for "texture<", I found about 360 occurrences in the whole DeepDRR project. Further, I learned that the project uses third party code from 2010 for performing cubic interpolation (http://www.dannyruijters.nl/cubicinterpolation/) which relies on texture objects.

I am not very familiar with CUDA programming, but I see the following options to resolve:

Refactor the existing code to use the new cudaTextureObject_t API, maybe under the help of a LMM.
Replace the CUDA kernels by an abstraction layer like cupy, which is basically equal to a rewrite.
Ask users to downgrade CUDA 11.x. However, this is not a sustainable solution and might not be possible in many setups (e.g. for students who don't have admin rights on the university machines).

mathiasunberath commented 1 year ago

This of course is unfortunate. We are reviewing this, and will likely be implementing a change. @benjamindkilleen and @liamjwang may have additional thoughts/updates on this now or soon.

benjamindkilleen commented 1 year ago

We have overhauled the Projector to use cupy, since Textures are officially deprecated with CUDA 12 and their replacement is not supported by PyCUDA. We will push this release after further testing, but if you would like to try it out, you can check out the dev branch and install from source.

arcadelab / deepdrr

Error when compiling kernel (using the example python script) #92