Hendrik/xgboost
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

disable Optin Shared Mem

Browse Source
This commit is contained in:
amdsc21 2023-03-15 02:10:16 +01:00
parent 8207015e48
commit 4484c7f073
2 changed files with 3 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
src/common/device_helpers.hip.h
View File

@ -173,9 +173,11 @@ inline size_t MaxSharedMemory(int device_idx) {
inline size_t MaxSharedMemoryOptin(int device_idx) { inline size_t MaxSharedMemoryOptin(int device_idx) {
int max_shared_memory = 0; int max_shared_memory = 0;
#if 0 /* CUDA Only */
dh::safe_cuda(hipDeviceGetAttribute dh::safe_cuda(hipDeviceGetAttribute
(&max_shared_memory, hipDeviceAttributeSharedMemPerBlockOptin, (&max_shared_memory, hipDeviceAttributeSharedMemPerBlockOptin,
device_idx)); device_idx));
#endif
return static_cast<std::size_t>(max_shared_memory); return static_cast<std::size_t>(max_shared_memory);
} }

2
src/tree/gpu_hist/histogram.cu
View File

@ -294,7 +294,7 @@ void BuildGradientHistogram(CUDAContext const* ctx, EllpackDeviceAccessor const&
#if defined(XGBOOST_USE_CUDA) #if defined(XGBOOST_USE_CUDA)
dh::safe_cuda(cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, dh::safe_cuda(cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize,
max_shared_memory)); max_shared_memory));
#elif defined(XGBOOST_USE_HIP) #elif defined(XGBOOST_USE_HIP) && 0 /* CUDA Only */
dh::safe_cuda(hipFuncSetAttribute((const void *)kernel, hipFuncAttributeMaxDynamicSharedMemorySize, dh::safe_cuda(hipFuncSetAttribute((const void *)kernel, hipFuncAttributeMaxDynamicSharedMemorySize,
max_shared_memory)); max_shared_memory));
#endif #endif