Use CUML_KERNEL for remaining kernels - #8465
Conversation
|
Could a maintainer please add the improvement and non-breaking labels? Label Checker requires one category and one compatibility label, and GitHub does not allow outside contributors to apply them. Thanks! |
divyegala
left a comment
There was a problem hiding this comment.
Thank you for the contribution!
|
/ok to test c74847c |
|
/ok to test 64858f6 |
|
The lint fix now passes, but the remaining jobs fail before building during AWS OIDC authentication. It looks infrastructure-related, but if there’s anything I can correct on the PR side, I’m happy to implement it. Thank you! @divyegala |
|
Thanks @ahmedtaha100 , we just changed organizations for this repository so we may see some transient errors. I'll follow up and run CI when it is ready, and let you know if there are any real failures related to this PR. |
|
Great thank you for your help! |
@divyegala @csadorf
Closes #6532
Summary
This fixes NVIDIA/cuml issue #6532, "Create preprocessor macros to standardize annotations."
CUML_KERNELalready exists incpp/include/cuml/common/utils.hpp, so this patch uses it instead of defining another macro.CUML_KERNELfor all 16 remaining C++ CUDA kernels.<cuml/common/utils.hpp>directly in each affected file.Ten of these kernels were already declared
static __global__. The other six gain internal linkage throughCUML_KERNEL, so they no longer produce externally linked kernel symbols.The Python NVRTC
RawKernelstrings remain unchanged because Python loads those kernels by name at runtime. Theirextern "C"entry points must remain visible.Kernel bodies and public APIs are unchanged.
Validation
cuml,SG_ISOLATION_FOREST_TEST, andSG_RF_TEST.detect-weak-linkingreported no externally linked CUDA kernel symbols in the builtlibcuml.so. Upstreamconda-cpp-checksuses the same detector.git diff --checkpassed.__global__ voidkernel definitions.This patch adds no unit test because runtime tests do not directly inspect linkage. The targeted builds compile every changed translation unit, and
detect-weak-linkingchecks the resultinglibcuml.so.