/************************************************************************* * Copyright (c) 2017-2022, NVIDIA CORPORATION. All rights reserved. * Modifications Copyright (c) 2019-2023 Advanced Micro Devices, Inc. All rights reserved. * * See LICENSE.txt for license information ************************************************************************/ #ifndef NCCL_PARAM_H_ #define NCCL_PARAM_H_ #include const char* userHomeDir(); void setEnvFile(const char* fileName); void initEnv(); void ncclLoadParam(char const* env, int64_t deftVal, int64_t uninitialized, int64_t* cache); #define NCCL_PARAM(name, env, deftVal) \ int64_t ncclParam##name() { \ constexpr int64_t uninitialized = INT64_MIN; \ static_assert(deftVal != uninitialized, "default value cannot be the uninitialized value."); \ static int64_t cache = uninitialized; \ if (__builtin_expect(__atomic_load_n(&cache, __ATOMIC_RELAXED) == uninitialized, false)) { \ ncclLoadParam("NCCL_" env, deftVal, uninitialized, &cache); \ } \ return cache; \ } #define RCCL_PARAM_DECLARE(name) \ int64_t rcclParam##name() #define RCCL_PARAM(name, env, deftVal) \ pthread_mutex_t rcclParamMutex##name = PTHREAD_MUTEX_INITIALIZER; \ int64_t rcclParam##name() { \ constexpr int64_t uninitialized = INT64_MIN; \ static_assert(deftVal != uninitialized, "default value cannot be the uninitialized value."); \ static int64_t cache = uninitialized; \ if (__builtin_expect(__atomic_load_n(&cache, __ATOMIC_RELAXED) == uninitialized, false)) { \ ncclLoadParam("RCCL_" env, deftVal, uninitialized, &cache); \ } \ return cache; \ } #endif