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.

Sign up for GitHub

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

prefer the use of cdf over alias on gfx10xx and gfx11xx #541

Merged
merged 1 commit into from
Aug 26, 2024
Merged

prefer the use of cdf over alias on gfx10xx and gfx11xx #541

Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
36 changes: 36 additions & 0 deletions library/include/rocrand/rocrand_discrete.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -40,6 +40,13 @@

#include "rocrand/rocrand_discrete_types.h"

// On certain architectures such as NAVI2 and NAVI3, double arithmetic is significantly slower.
// In such cases we want to prefer the CDF method over the alias method.
// This macro is undefined at the end of the file.
#if defined(__HIP_DEVICE_COMPILE__) && (defined(__GFX10__) || defined(__GFX11__))
#define ROCRAND_PREFER_CDF_OVER_ALIAS
#endif

// Alias method
//
// Walker, A. J.
Expand Down Expand Up @@ -281,7 +288,11 @@ __forceinline__ __device__ unsigned int
rocrand_discrete(rocrand_state_mtgp32* state,
const rocrand_discrete_distribution discrete_distribution)
{
#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
#else
return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
#endif
}

/**
Expand Down Expand Up @@ -376,7 +387,11 @@ __forceinline__ __device__ __host__ unsigned int
rocrand_discrete(rocrand_state_lfsr113* state,
const rocrand_discrete_distribution discrete_distribution)
{
#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
#else
return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
#endif
}

/**
Expand All @@ -395,7 +410,11 @@ __forceinline__ __device__ __host__ unsigned int
rocrand_discrete(rocrand_state_threefry2x32_20* state,
const rocrand_discrete_distribution discrete_distribution)
{
#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
#else
return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
#endif
}

/**
Expand All @@ -414,7 +433,11 @@ __forceinline__ __device__ __host__ unsigned int
rocrand_discrete(rocrand_state_threefry2x64_20* state,
const rocrand_discrete_distribution discrete_distribution)
{
#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
#else
return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
#endif
}

/**
Expand All @@ -433,7 +456,11 @@ __forceinline__ __device__ __host__ unsigned int
rocrand_discrete(rocrand_state_threefry4x32_20* state,
const rocrand_discrete_distribution discrete_distribution)
{
#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
#else
return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
#endif
}

/**
Expand All @@ -452,9 +479,18 @@ __forceinline__ __device__ __host__ unsigned int
rocrand_discrete(rocrand_state_threefry4x64_20* state,
const rocrand_discrete_distribution discrete_distribution)
{
#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
#else
return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
#endif
}

/** @} */ // end of group rocranddevice

// Undefine the macro that may be defined at the top of the file!
#if defined(ROCRAND_PREFER_CDF_OVER_ALIAS)
#undef ROCRAND_PREFER_CDF_OVER_ALIAS
#endif

#endif // ROCRAND_DISCRETE_H_
Loading