blender
/
blender
128
396
Fork 569
Code Issues 6.2k Pull Requests 755 Projects 19 Wiki Activity

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:
Sayak Biswas 2022-06-28 16:55:27 +02:00 committed by Brecht Van Lommel
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
6 changed files with 9 additions and 6 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
CMakeLists.txt
View File

@ -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()

2
build_files/config/pipeline_config.yaml
View File

@ -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:

4
intern/cycles/blender/addon/properties.py
View File

@ -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')

2
intern/cycles/device/hip/util.h
View File

@ -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

2
intern/cycles/kernel/device/hip/compat.h
View File

@ -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)

3
intern/cycles/util/math.h
View File

@ -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