Cycles: enable Vega GPU/APU support
Enables Vega and Vega II GPUs as well as Vega APU, using changes in HIP code to support 64-bit waves and a new HIP SDK version. Tested with Radeon WX9100, Radeon VII GPUs and Ryzen 7 PRO 5850U with Radeon Graphics APU. Ref T96740, T91571 Differential Revision: https://developer.blender.org/D15242
This commit is contained in:
parent
270ed1c716
commit
abfa09752f
Notes:
blender-bot
2023-02-14 02:27:51 +01:00
Referenced by issue #96740, Enable HIP on older devices Referenced by issue #91571, Cycles HIP device
|
@ -444,7 +444,7 @@ endif()
|
|||
if(NOT APPLE)
|
||||
option(WITH_CYCLES_DEVICE_HIP "Enable Cycles AMD HIP support" ON)
|
||||
option(WITH_CYCLES_HIP_BINARIES "Build Cycles AMD HIP binaries" OFF)
|
||||
set(CYCLES_HIP_BINARIES_ARCH gfx1010 gfx1011 gfx1012 gfx1030 gfx1031 gfx1032 gfx1034 gfx1035 CACHE STRING "AMD HIP architectures to build binaries for")
|
||||
set(CYCLES_HIP_BINARIES_ARCH gfx900 gfx906 gfx90c gfx902 gfx1010 gfx1011 gfx1012 gfx1030 gfx1031 gfx1032 gfx1034 gfx1035 CACHE STRING "AMD HIP architectures to build binaries for")
|
||||
mark_as_advanced(WITH_CYCLES_DEVICE_HIP)
|
||||
mark_as_advanced(CYCLES_HIP_BINARIES_ARCH)
|
||||
endif()
|
||||
|
|
|
@ -55,7 +55,7 @@ buildbot:
|
|||
cuda11:
|
||||
version: '11.4.1'
|
||||
hip:
|
||||
version: '5.0.20451'
|
||||
version: '5.2.21440'
|
||||
optix:
|
||||
version: '7.3.0'
|
||||
cmake:
|
||||
|
|
|
@ -1545,10 +1545,10 @@ class CyclesPreferences(bpy.types.AddonPreferences):
|
|||
elif device_type == 'HIP':
|
||||
import sys
|
||||
if sys.platform[:3] == "win":
|
||||
col.label(text="Requires discrete AMD GPU with RDNA architecture", icon='BLANK1')
|
||||
col.label(text="Requires AMD GPU with Vega or RDNA architecture", icon='BLANK1')
|
||||
col.label(text="and AMD Radeon Pro 21.Q4 driver or newer", icon='BLANK1')
|
||||
elif sys.platform.startswith("linux"):
|
||||
col.label(text="Requires discrete AMD GPU with RDNA architecture", icon='BLANK1')
|
||||
col.label(text="Requires AMD GPU with Vega or RDNA architecture", icon='BLANK1')
|
||||
col.label(text="and AMD driver version 22.10 or newer", icon='BLANK1')
|
||||
elif device_type == 'METAL':
|
||||
col.label(text="Requires Apple Silicon with macOS 12.2 or newer", icon='BLANK1')
|
||||
|
|
|
@ -51,7 +51,7 @@ static inline bool hipSupportsDevice(const int hipDevId)
|
|||
hipDeviceGetAttribute(&major, hipDeviceAttributeComputeCapabilityMajor, hipDevId);
|
||||
hipDeviceGetAttribute(&minor, hipDeviceAttributeComputeCapabilityMinor, hipDevId);
|
||||
|
||||
return (major > 10) || (major == 10 && minor >= 1);
|
||||
return (major >= 9);
|
||||
}
|
||||
|
||||
CCL_NAMESPACE_END
|
||||
|
|
|
@ -62,7 +62,7 @@ typedef unsigned long long uint64_t;
|
|||
#define ccl_gpu_block_idx_x (blockIdx.x)
|
||||
#define ccl_gpu_grid_dim_x (gridDim.x)
|
||||
#define ccl_gpu_warp_size (warpSize)
|
||||
#define ccl_gpu_thread_mask(thread_warp) uint(0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp))
|
||||
#define ccl_gpu_thread_mask(thread_warp) uint64_t((1ull << thread_warp) - 1)
|
||||
|
||||
#define ccl_gpu_global_id_x() (ccl_gpu_block_idx_x * ccl_gpu_block_dim_x + ccl_gpu_thread_idx_x)
|
||||
#define ccl_gpu_global_size_x() (ccl_gpu_grid_dim_x * ccl_gpu_block_dim_x)
|
||||
|
|
|
@ -793,6 +793,9 @@ ccl_device_inline uint popcount(uint x)
|
|||
return i & 1;
|
||||
}
|
||||
# endif
|
||||
#elif defined(__KERNEL_HIP__)
|
||||
/* Use popcll to support 64-bit wave for pre-RDNA AMD GPUs */
|
||||
# define popcount(x) __popcll(x)
|
||||
#elif !defined(__KERNEL_METAL__)
|
||||
# define popcount(x) __popc(x)
|
||||
#endif
|
||||
|
|
Loading…
Reference in New Issue