Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

ENH: Add support for 16 and 32-bit radiograph images #76

Merged
merged 1 commit into from
Jan 16, 2024

Conversation

NicerNewerCar
Copy link
Contributor

@NicerNewerCar NicerNewerCar commented Jan 13, 2023

Copy link
Collaborator

@amymmorton amymmorton left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Successful build with 12.0, but when opening trial - autoscoper crashed:

C:\Users\amorton1\Documents\GitHub\Autoscoper\libautoscoper\src\View.cpp(430) : cudaSafeCall() Runtime API error : OS call failed or operation not supported on this OS.

@NicerNewerCar
Copy link
Contributor Author

NicerNewerCar commented Jan 17, 2023

Successful build with 12.0, but when opening trial - autoscoper crashed:

@amymmorton what trial are you opening? Did you make sure to build with a clean build tree? I have opened and tracked WN00106 and SOL001A successfully on this branch.

@amymmorton
Copy link
Collaborator

Yes. tried again - changed to main then back to PR 76:
Release configuration
same error on trial open:

Filename: P:/iTWA_Instrumented_Total_Wrist/Subjects/NORMAL/WN00106/Autoscoper/flx_ext.cfg
Volume Name: P:\iTWA_Instrumented_Total_Wrist\Subjects\NORMAL\WN00106\Models\mc2_mc3\mc2_mc3_dcm_cropped.tif
Volume Name: P:\iTWA_Instrumented_Total_Wrist\Subjects\NORMAL\WN00106\Models\rad\rad_dcm_cropped.tif
Volume Name: P:\iTWA_Instrumented_Total_Wrist\Subjects\NORMAL\WN00106\Models\mc3\mc3_dcm_cropped.tif
Volume Name: P:\iTWA_Instrumented_Total_Wrist\Subjects\NORMAL\WN00106\Models\uln\uln_dcm_cropped.tif
Volume Name: P:\iTWA_Instrumented_Total_Wrist\Subjects\NORMAL\WN00106\Models\mc3_mc4\mc3_mc4_dcm_cropped.tif
C:\Users\amorton1\Documents\GitHub\Autoscoper\libautoscoper\src\View.cpp(430) : cudaSafeCall() Runtime API error : OS call failed or operation not supported on this OS.

@amymmorton
Copy link
Collaborator

I tried uninstalling v12.0 (revert to 11.8) and the same error on trial open persists:
confirmv118

C:\Users\amorton1\Documents\GitHub\Autoscoper\build_PR76_11p8\Autoscoper-build\install\bin\Release>autoscoper.exe
Autoscoper v2.8

Setting to Autoscoper_RANDOM_SEED to 100
Graphics Card Vendor: Intel
Intel(R) UHD Graphics 630
4.6.0 - Build 30.0.101.1692
Initializing OpenGL...
Graphics Card Vendor: Intel
Intel(R) UHD Graphics 630
4.6.0 - Build 30.0.101.1692
Initializing OpenGL...

Setting up a new trial...
Filename: P:/iTWA_Instrumented_Total_Wrist/Subjects/NORMAL/WN00106/Autoscoper/flx_ext.cfg
Volume Name: P:\iTWA_Instrumented_Total_Wrist\Subjects\NORMAL\WN00106\Models\mc2_mc3\mc2_mc3_dcm_cropped.tif
Volume Name: P:\iTWA_Instrumented_Total_Wrist\Subjects\NORMAL\WN00106\Models\rad\rad_dcm_cropped.tif
Volume Name: P:\iTWA_Instrumented_Total_Wrist\Subjects\NORMAL\WN00106\Models\mc3\mc3_dcm_cropped.tif
Volume Name: P:\iTWA_Instrumented_Total_Wrist\Subjects\NORMAL\WN00106\Models\uln\uln_dcm_cropped.tif
Volume Name: P:\iTWA_Instrumented_Total_Wrist\Subjects\NORMAL\WN00106\Models\mc3_mc4\mc3_mc4_dcm_cropped.tif
C:\Users\amorton1\Documents\GitHub\Autoscoper\libautoscoper\src\View.cpp(430) : cudaSafeCall() Runtime API error : OS call failed or operation not supported on this OS.

@amymmorton
Copy link
Collaborator

pulled from main and had the same error- I believe the instll of v12 created issues in my enviroment- attempting a reinstall of 11.8..

@amymmorton
Copy link
Collaborator

My reinstall of 11.8 allowed successful build of this branch. I will have our engineer Maddie and @kaitohl test v12.0 too

@amymmorton amymmorton closed this Jan 20, 2023
@amymmorton amymmorton reopened this Jan 20, 2023
@NicerNewerCar
Copy link
Contributor Author

@jcfr I got these changes to build and run successfully using the docker container on Ubuntu 22.04. Have you been able to test them yet? If so, this can be merged in.

Copy link
Collaborator

@amymmorton amymmorton left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Awesome- confirmed- mayo shoulder data loads and tracks:
shoulderCudaLoadSuccess

@jcfr
Copy link
Contributor

jcfr commented Dec 20, 2023

Suggest commit message:

ENH: Add support for 16 and 32-bit radiograph images

Prior to these changes, the supported CUDA texture types (and by extension
radiograph images) was hard-coded to 8-bit (unsigned char) and managed as
a static global variable.
  
Since texture references needed to be known at compile time and the current
type was hard-coded to 8-bit, the use of them prevented from being able to
easily support 8, 16, and 32 bit radiograph images known at runtime. 

To address this, this commit removes the use of `cudaBindTextureToArray`
(deprecated since CUDA 5.0 and removed in CUDA 12.x) and switch to using
texture object instead by introducing a convenience helper function
`createTextureObjectFromArray`.

Approach implemented here is based of recommendation the post titled
"Kepler Texture Objects Improve Performance and Flexibility" published
in 2013.
See https://developer.nvidia.com/blog/cuda-pro-tip-kepler-texture-objects-improve-performance-and-flexibility/

For reference, before this changes, the error message reported at build time
while attempting to use autoscoper built against CUDA 12.x was:

```
3 Errors under file path `...\libautoscoper\src\gpu\cuda\RadRenderer_kernels.cu:
- no instance of overloaded function "tex2D" matches the argument list 
- identifier "cudaBindTextureToArray" is undefined
- texture is not a template
```

Can you revisit to:

Comment on lines +92 to +93
dim3 gridDim(((unsigned int)width+blockDim.x-1)/blockDim.x,
((unsigned int)height+blockDim.y-1)/blockDim.y);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Should the outcome of the division be truncated to unsigned int ? The use of parenthesis seems off.

Also, could you confirm truncation is what we want here ?

See https://stackoverflow.com/questions/65225888/difference-among-static-cast-int-and-float2int-rn-in-cuda

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is not performing truncation on the division, but converting the size_t width and height variables to prevent a compile time warning (see #42)

C:\SAM\Autoscoper\libautoscoper\src\gpu\cuda\Compositor_kernels.cu(70): warning C4267: 'argument': conversion from 'size_t' to 'unsigned int', possible loss of data [C:\SAM\Autoscoper\build-test\Autoscoper-build\libautoscoper\libautoscoper.vcxproj] [C:\SAM\Autoscoper\build-test\Autoscoper.vcxproj]

Prior to these changes, the supported CUDA texture types (and by extension
radiograph images) was hard-coded to 8-bit (unsigned char) and managed as
a static global variable.

Since texture references needed to be known at compile time and the current
type was hard-coded to 8-bit, the use of them prevented from being able to
easily support 8, 16, and 32 bit radiograph images known at runtime.

To address this, this commit removes the use of `cudaBindTextureToArray`
(Deprecated since CUDA 11.3 and removed in CUDA 12.x)
and switch to using texture object instead by introducing a convenience helper
function `createTextureObjectFromArray`.

Approach implemented here is based of recommendation the post titled
"Kepler Texture Objects Improve Performance and Flexibility" published
in 2013.
See https://developer.nvidia.com/blog/cuda-pro-tip-kepler-texture-objects-improve-performance-and-flexibility/

For reference, before this changes, the error message reported at build time
while attempting to use autoscoper built against CUDA 12.x was:

```
3 Errors under file path `...\libautoscoper\src\gpu\cuda\RadRenderer_kernels.cu:
- no instance of overloaded function "tex2D" matches the argument list
- identifier "cudaBindTextureToArray" is undefined
- texture is not a template
```
@jcfr jcfr merged commit 866a5d2 into BrownBiomechanics:main Jan 16, 2024
3 checks passed
@NicerNewerCar NicerNewerCar deleted the cuda-texture-objects branch January 16, 2024 17:50
@jcfr jcfr changed the title Update texture references to texture objects ENH: Add support for 16 and 32-bit radiograph images Feb 5, 2024
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
3 participants