Errors building fatbin HIP kernels for AMD GPU with Rocm 5.1.3

[FogLizard]

Making package: blender-git 3.3.r115268.gae39abe7f3c-1

[224/4619] Generating kernel_gfx1010.fatbin FAILED: intern/cycles/kernel/kernel_gfx1010.fatbin /home/pink/.cache/pikaur/build/blender-git/src/build/intern/cycles/kernel/kernel_gfx1010.fatbin cd /home/pink/.cache/pikaur/build/blender-git/src/build/intern/cycles/kernel && /usr/bin/hipcc --amdgpu-target=gfx1010 --genco /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/device/hip/kernel.cpp -D CCL_NAMESPACE_BEGIN= -D CCL_NAMESPACE_END= -D HIPCC -I /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/.. -I /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/device/hip -Wno-parentheses-equality -Wno-unused-value --hipcc-func-supp -ffast-math -o /home/pink/.cache/pikaur/build/blender-git/src/build/intern/cycles/kernel/kernel_gfx1010.fatbin In file included from /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/device/hip/kernel.cpp:12: /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:55:1: error: unknown type name '__noinline__' ccl_device_noinline T kernel_tex_image_interp_bicubic(ccl_global const TextureInfo &info, ^ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/hip/compat.h:33:40: note: expanded from macro 'ccl_device_noinline' #define ccl_device_noinline __device__ __noinline__ ^ In file included from /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/device/hip/kernel.cpp:12: /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:55:21: warning: variable templates are a C++14 extension [-Wc++14-extensions] ccl_device_noinline T kernel_tex_image_interp_bicubic(ccl_global const TextureInfo &info, ^ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:55:22: error: expected ';' at end of declaration ccl_device_noinline T kernel_tex_image_interp_bicubic(ccl_global const TextureInfo &info, ^ ; /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:55:23: error: C++ requires a type specifier for all declarations ccl_device_noinline T kernel_tex_image_interp_bicubic(ccl_global const TextureInfo &info, ^ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:69:15: error: no matching function for call to 'cubic_g0' float g0x = cubic_g0(fx); ^~~~~~~~ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:34:18: note: candidate function not viable: call to __device__ function from __host__ function ccl_device float cubic_g0(float a) ^ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:70:15: error: no matching function for call to 'cubic_g1' float g1x = cubic_g1(fx); ^~~~~~~~ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:38:18: note: candidate function not viable: call to __device__ function from __host__ function ccl_device float cubic_g1(float a) ^ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:72:20: error: no matching function for call to 'cubic_h0' float x0 = (px + cubic_h0(fx) + 0.5f) / info.width; ^~~~~~~~ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:44:18: note: candidate function not viable: call to __device__ function from __host__ function ccl_device float cubic_h0(float a) ^ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:73:20: error: no matching function for call to 'cubic_h1' float x1 = (px + cubic_h1(fx) + 0.5f) / info.width; ^~~~~~~~ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:48:18: note: candidate function not viable: call to __device__ function from __host__ function ccl_device float cubic_h1(float a) ^ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:74:20: error: no matching function for call to 'cubic_h0' float y0 = (py + cubic_h0(fy) + 0.5f) / info.height; ^~~~~~~~ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:44:18: note: candidate function not viable: call to __device__ function from __host__ function ccl_device float cubic_h0(float a) ^ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:75:20: error: no matching function for call to 'cubic_h1' float y1 = (py + cubic_h1(fy) + 0.5f) / info.height; ^~~~~~~~ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:48:18: note: candidate function not viable: call to __device__ function from __host__ function ccl_device float cubic_h1(float a) ^ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:77:10: error: no matching function for call to 'cubic_g0' return cubic_g0(fy) * (g0x * ccl_gpu_tex_object_read_2D<T>(tex, x0, y0) + ^~~~~~~~ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:34:18: note: candidate function not viable: call to __device__ function from __host__ function ccl_device float cubic_g0(float a) ^ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:79:10: error: no matching function for call to 'cubic_g1' cubic_g1(fy) * (g0x * ccl_gpu_tex_object_read_2D<T>(tex, x0, y1) + ^~~~~~~~ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:38:18: note: candidate function not viable: call to __device__ function from __host__ function ccl_device float cubic_g1(float a) ^ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:85:1: error: unknown type name '__noinline__' ccl_device_noinline T ^ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/hip/compat.h:33:40: note: expanded from macro 'ccl_device_noinline' #define ccl_device_noinline __device__ __noinline__ ^ In file included from /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/device/hip/kernel.cpp:12: /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:85:21: warning: variable templates are a C++14 extension [-Wc++14-extensions] ccl_device_noinline T ^ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:85:22: error: expected ';' at end of declaration ccl_device_noinline T ^ ; /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:86:1: error: C++ requires a type specifier for all declarations kernel_tex_image_interp_tricubic(ccl_global const TextureInfo &info, float x, float y, float z) ^ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:101:15: error: no matching function for call to 'cubic_g0' float g0x = cubic_g0(fx); ^~~~~~~~ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:34:18: note: candidate function not viable: call to __device__ function from __host__ function ccl_device float cubic_g0(float a) ^ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:102:15: error: no matching function for call to 'cubic_g1' float g1x = cubic_g1(fx); ^~~~~~~~ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:38:18: note: candidate function not viable: call to __device__ function from __host__ function ccl_device float cubic_g1(float a) ^ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:103:15: error: no matching function for call to 'cubic_g0' float g0y = cubic_g0(fy); ^~~~~~~~ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:34:18: note: candidate function not viable: call to __device__ function from __host__ function ccl_device float cubic_g0(float a) ^ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:104:15: error: no matching function for call to 'cubic_g1' float g1y = cubic_g1(fy); ^~~~~~~~ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:38:18: note: candidate function not viable: call to __device__ function from __host__ function ccl_device float cubic_g1(float a) ^ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:105:15: error: no matching function for call to 'cubic_g0' float g0z = cubic_g0(fz); ^~~~~~~~ /home/pink/.cache/pikaur/build/blender-git/src/blender/intern/cycles/kernel/../kernel/device/gpu/image.h:34:18: note: candidate function not viable: call to __device__ function from __host__ function ccl_device float cubic_g0(float a) ^ fatal error: too many errors emitted, stopping now [-ferror-limit=] 2 warnings and 20 errors generated when compiling for gfx1010.

[FogLizard]

deleted all Rocm packages, fully updated system, rebuilt Rocm 5.1.3 and still blender-git fails to build with the same error.

[FogLizard]

any info on this? it just fails to build every day for me

removing Rocm 5.1.3 (rocm-hip-sdk rocm-opencl-sdk from AUR with all their dependencies) lets me build Blender successfully without AMD HIP GPU rendering supported.

[EvanWoodard]

I'm having this exact same problem. Maybe there's some flag that I'm supposed to set, but I can't find documentation that says so.

Any updates?

[EvanWoodard]

@FogLizard - Not that it helps with this issue. However, 3.2 has been released and downloading and running the Linux release from the Blender page https://www.blender.org/download/releases/3-2/ does work on Arch. I say this because I was more excited for the AMD GPU support of 3.2 than building blender-git itself, and that is working.

[FogLizard]

@EvanWoodard In my case the point is to get Vega to work, cause AMD took a genius decision to drop support for anything older than RDNA, and only now they are working on Vega support after 9 months since 3.0 alpha. So I need the alpha to build.