Skip to content

Commit 3aafde5

Browse files
fwyzardpsychocoderHPC
authored andcommitted
[do not merge] Add ALPAKA_DEFAULT_AMD_WAVEFRONT_SIZE
1 parent f71bcbf commit 3aafde5

File tree

1 file changed

+12
-2
lines changed

1 file changed

+12
-2
lines changed

include/alpaka/warp/WarpUniformCudaHipBuiltIn.hpp

Lines changed: 12 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -43,7 +43,7 @@ namespace alpaka::warp
4343
template<>
4444
struct GetSizeCompileTime<WarpUniformCudaHipBuiltIn>
4545
{
46-
static constexpr auto getSizeCompileTime() -> std::int32_t
46+
__device__ static constexpr auto getSizeCompileTime() -> std::int32_t
4747
{
4848
# if defined(__CUDA_ARCH__)
4949
// CUDA always has a warp size of 32
@@ -58,7 +58,12 @@ namespace alpaka::warp
5858
return 32;
5959
# else
6060
// Unknown AMD GPU architecture
61+
# ifdef ALPAKA_DEFAULT_AMD_WAVEFRONT_SIZE
62+
return ALPAKA_DEFAULT_AMD_WAVEFRONT_SIZE
63+
# else
64+
# error The current AMD GPU architucture is not supported by this version of alpaka. You can define a default wavefront size setting the preprocessor macro ALPAKA_DEFAULT_AMD_WAVEFRONT_SIZE
6165
return 0;
66+
# endif
6267
# endif
6368
# endif
6469
// Host compilation
@@ -69,7 +74,7 @@ namespace alpaka::warp
6974
template<>
7075
struct GetSizeUpperLimit<WarpUniformCudaHipBuiltIn>
7176
{
72-
static constexpr auto getSizeUpperLimit() -> std::int32_t
77+
__device__ static constexpr auto getSizeUpperLimit() -> std::int32_t
7378
{
7479
# if defined(__CUDA_ARCH__)
7580
// CUDA always has a warp size of 32
@@ -84,7 +89,12 @@ namespace alpaka::warp
8489
return 32;
8590
# else
8691
// Unknown AMD GPU architecture
92+
# ifdef ALPAKA_DEFAULT_AMD_WAVEFRONT_SIZE
93+
return ALPAKA_DEFAULT_AMD_WAVEFRONT_SIZE
94+
# else
95+
# error The current AMD GPU architucture is not supported by this version of alpaka. You can define a default wavefront size setting the preprocessor macro ALPAKA_DEFAULT_AMD_WAVEFRONT_SIZE
8796
return 64;
97+
# endif
8898
# endif
8999
# endif
90100
// Host compilation

0 commit comments

Comments
 (0)