From: Cordell Bloor Date: Thu, 13 Apr 2023 22:11:52 -0600 Subject: hide kernel symbols If not marked as static, the rocblas kernels would be weak public symbols. They are not intended to be visible, but are not affected by -fvisiblity=hidden and cannot be entirely hidden except by being marked as static. Applied-Uptream: https://github.com/ROCmSoftwarePlatform/rocBLAS/commit/c311f3ce684368091acae744c924bcddea4add33 --- library/src/blas2/rocblas_trsv_kernels.cpp | 2 +- library/src/include/macros.hpp | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/library/src/blas2/rocblas_trsv_kernels.cpp b/library/src/blas2/rocblas_trsv_kernels.cpp index 5b7a34b..5225dc7 100644 --- a/library/src/blas2/rocblas_trsv_kernels.cpp +++ b/library/src/blas2/rocblas_trsv_kernels.cpp @@ -458,7 +458,7 @@ void ROCBLAS_KERNEL_ILF rocblas_trsv_block_solve_upper(const T* __restrict__ A, } } -static ROCBLAS_KERNEL(1) rocblas_trsv_init(rocblas_int* w_completed_sec) +ROCBLAS_KERNEL(1) rocblas_trsv_init(rocblas_int* w_completed_sec) { // The last block section which has been completed (for each batch) w_completed_sec[blockIdx.x] = -1; diff --git a/library/src/include/macros.hpp b/library/src/include/macros.hpp index 93f54b4..36d75b0 100644 --- a/library/src/include/macros.hpp +++ b/library/src/include/macros.hpp @@ -31,8 +31,8 @@ //#else //#endif -#define ROCBLAS_KERNEL(lb_) __global__ __launch_bounds__((lb_)) void -#define ROCBLAS_KERNEL_NO_BOUNDS __global__ void +#define ROCBLAS_KERNEL(lb_) __global__ __launch_bounds__((lb_)) static void +#define ROCBLAS_KERNEL_NO_BOUNDS __global__ static void // A storage-class-specifier other than thread_local shall not be specified in an explicit specialization. // ROCBLAS_KERNEL_INSTANTIATE should be used where kernels are instantiated to avoid use of static.