Compilation error for navi10 (use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD')

Hello, I have some trouble to compile composable_kernel for my AMD GPU architecture (gfx1010)

cmake                                                                                             \
-D CMAKE_PREFIX_PATH=/opt/rocm                                                                    \
-D CMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc                                                         \
-D CMAKE_CXX_FLAGS="-O3"                                                                          \
-D CMAKE_BUILD_TYPE=Release                                                                       \
-D GPU_TARGETS="gfx1010"                                                                    \

-- The C compiler identification is GNU 9.4.0
-- The CXX compiler identification is Clang 15.0.0
-- Detecting C compiler ABI info
-- Detecting C compiler ABI info - done
-- Check for working C compiler: /usr/bin/cc - skipped
-- Detecting C compile features
-- Detecting C compile features - done
-- Detecting CXX compiler ABI info
-- Detecting CXX compiler ABI info - done
-- Check for working CXX compiler: /opt/rocm/bin/hipcc - skipped
-- Detecting CXX compile features
-- Detecting CXX compile features - done
-- Looking for pthread.h
-- Looking for pthread.h - found
-- Performing Test CMAKE_HAVE_LIBC_PTHREAD
-- Performing Test CMAKE_HAVE_LIBC_PTHREAD - Failed
-- Check if compiler accepts -pthread
-- Check if compiler accepts -pthread - yes
-- Found Threads: TRUE  
CMAKE_CXX_COMPILER_ID: Clang
OpenMP_CXX_LIB_NAMES: libomp;libgomp;libiomp5
OpenMP_gomp_LIBRARY: 
OpenMP_pthread_LIBRARY: 
OpenMP_CXX_FLAGS: -fopenmp=libomp -Wno-unused-command-line-argument
-- hip::amdhip64 is SHARED_LIBRARY
-- Performing Test HIP_CLANG_SUPPORTS_PARALLEL_JOBS
-- Performing Test HIP_CLANG_SUPPORTS_PARALLEL_JOBS - Success
-- Build with HIP 5.4.22505
-- Clang tidy not found
-- Clang tidy checks: *,-abseil-*,-android-cloexec-fopen,-cert-msc30-c,-bugprone-exception-escape,-bugprone-macro-parentheses,-cert-env33-c,-cert-msc32-c,-cert-msc50-cpp,-cert-msc51-cpp,-cert-dcl37-c,-cert-dcl51-cpp,-clang-analyzer-alpha.core.CastToStruct,-clang-analyzer-optin.performance.Padding,-clang-diagnostic-deprecated-declarations,-clang-diagnostic-extern-c-compat,-clang-diagnostic-unused-command-line-argument,-cppcoreguidelines-avoid-c-arrays,-cppcoreguidelines-avoid-magic-numbers,-cppcoreguidelines-explicit-virtual-functions,-cppcoreguidelines-init-variables,-cppcoreguidelines-macro-usage,-cppcoreguidelines-non-private-member-variables-in-classes,-cppcoreguidelines-pro-bounds-array-to-pointer-decay,-cppcoreguidelines-pro-bounds-constant-array-index,-cppcoreguidelines-pro-bounds-pointer-arithmetic,-cppcoreguidelines-pro-type-member-init,-cppcoreguidelines-pro-type-reinterpret-cast,-cppcoreguidelines-pro-type-union-access,-cppcoreguidelines-pro-type-vararg,-cppcoreguidelines-special-member-functions,-fuchsia-*,-google-explicit-constructor,-google-readability-braces-around-statements,-google-readability-todo,-google-runtime-int,-google-runtime-references,-hicpp-vararg,-hicpp-braces-around-statements,-hicpp-explicit-conversions,-hicpp-named-parameter,-hicpp-no-array-decay,-hicpp-avoid-c-arrays,-hicpp-signed-bitwise,-hicpp-special-member-functions,-hicpp-uppercase-literal-suffix,-hicpp-use-auto,-hicpp-use-equals-default,-hicpp-use-override,-llvm-header-guard,-llvm-include-order,-llvmlibc-restrict-system-libc-headers,-llvmlibc-callee-namespace,-llvmlibc-implementation-in-namespace,-llvm-else-after-return,-llvm-qualified-auto,-misc-misplaced-const,-misc-non-private-member-variables-in-classes,-misc-no-recursion,-modernize-avoid-bind,-modernize-avoid-c-arrays,-modernize-pass-by-value,-modernize-use-auto,-modernize-use-default-member-init,-modernize-use-equals-default,-modernize-use-trailing-return-type,-modernize-use-transparent-functors,-performance-unnecessary-value-param,-readability-braces-around-statements,-readability-else-after-return,-readability-function-cognitive-complexity,-readability-isolate-declaration,-readability-magic-numbers,-readability-named-parameter,-readability-uppercase-literal-suffix,-readability-convert-member-functions-to-static,-readability-qualified-auto,-readability-redundant-string-init,-bugprone-narrowing-conversions,-cppcoreguidelines-narrowing-conversions,-altera-struct-pack-align,-cppcoreguidelines-prefer-member-initializer
CMAKE_CXX_FLAGS: -O3
adding instance device_batched_gemm_instance
adding instance device_batched_gemm_add_relu_gemm_add_instance
adding instance device_batched_gemm_bias_permute_instance
adding instance device_batched_gemm_gemm_instance
adding instance device_batched_gemm_multi_d_instance
adding instance device_batched_gemm_reduce_instance
adding instance device_batched_gemm_softmax_gemm_instance
adding instance device_batched_gemm_softmax_gemm_permute_instance
adding instance device_batchnorm_instance
adding instance device_contraction_bilinear_instance
adding instance device_contraction_scale_instance
adding instance device_conv1d_bwd_data_instance
adding instance device_conv2d_bwd_data_instance
adding instance device_conv2d_fwd_instance
adding instance device_conv2d_fwd_bias_relu_instance
adding instance device_conv2d_fwd_bias_relu_add_instance
adding instance device_conv3d_bwd_data_instance
adding instance device_elementwise_instance
adding instance device_elementwise_normalization_instance
adding instance device_gemm_instance
adding instance device_gemm_add_add_fastgelu_instance
adding instance device_gemm_add_fastgelu_instance
adding instance device_gemm_add_multiply_instance
adding instance device_gemm_add_relu_add_layernorm_instance
adding instance device_gemm_bias_add_reduce_instance
adding instance device_gemm_bilinear_instance
adding instance device_gemm_fastgelu_instance
adding instance device_gemm_reduce_instance
adding instance device_gemm_splitk_instance
adding instance device_grouped_conv1d_bwd_weight_instance
adding instance device_grouped_conv1d_fwd_instance
adding instance device_grouped_conv2d_bwd_data_instance
adding instance device_grouped_conv2d_bwd_weight_instance
adding instance device_grouped_conv2d_fwd_instance
adding instance device_grouped_conv3d_bwd_weight_instance
adding instance device_grouped_conv3d_fwd_instance
adding instance device_grouped_gemm_instance
adding instance device_grouped_gemm_fastgelu_instance
adding instance device_normalization_instance
adding instance device_pool_fwd_instance
adding instance device_quantization_instance
adding instance device_reduce_instance
adding instance device_softmax_instance
adding example example_gemm_dl_fp32
adding example example_gemm_dl_fp16
adding example example_gemm_dl_int8
adding example example_gemm_xdl_fp16
adding example example_gemm_xdl_wavelet_fp16
adding example example_gemm_xdl_bf16
adding example example_gemm_xdl_int8
adding example example_gemm_xdl_skip_b_lds_fp16
adding example example_gemm_xdl_fp64
adding example example_convnd_fwd_dl_fp16
adding example example_convnd_fwd_dl_fp32
adding example example_convnd_fwd_dl_int8
adding example example_reduce_blockwise
adding example example_reduce_multiblock_atomic_add
adding example example_reduce_blockwise_two_call
adding example example_pool2d_fwd_fp16
adding example example_pool2d_fwd_fp32
adding example example_gemm_dl_quantization_int8
adding example example_grouped_gemm_xdl_fp32
adding example example_grouped_gemm_xdl_fp16
adding example example_grouped_gemm_xdl_bfp16
adding example example_grouped_gemm_xdl_int8
adding example example_grouped_gemm_multiple_d_dl_fp16
adding example example_grouped_gemm_xdl_splitk_fp16
adding example example_convnd_bwd_data_dl_fp16
adding example example_broadcast_add_2d_amn_bn
adding example example_broadcast_add_3d_am_bmnk
adding example example_elementwise_add_1d
adding example example_elementwise_add_4d
adding example example_grouped_conv_bwd_weight_dl_fp16
adding example example_cgemm_xdl_bf16
adding example example_cgemm_xdl_fp16
adding example example_cgemm_xdl_fp32
adding example example_cgemm_xdl_int8
adding example example_softmax_blockwise
adding example example_batched_gemm_xdl_fp32
adding example example_batched_gemm_xdl_fp16
adding example example_batched_gemm_xdl_bfp16
adding example example_batched_gemm_xdl_int8
adding example example_gemm_bias_e_permute_g1m3n2k1_xdl_fp16
adding example example_gemm_bias_e_permute_g1m2n3k1_xdl_fp16
adding example example_contraction_bilinear_xdl_fp32
adding example example_contraction_scale_xdl_fp32
adding example example_contraction_bilinear_xdl_fp64
adding example example_contraction_scale_xdl_fp64
adding example example_layernorm_fp16
adding example example_layernorm_splitk_fp16
adding example example_grouped_gemm_bias_e_permute_xdl_fp16
adding example example_batched_gemm_bias_e_permute_xdl_fp16
adding example example_batched_gemm_scale_softmax_gemm_xdl_fp16
adding example example_batched_gemm_scale_softmax_gemm_xdl_bf16
adding example example_batched_gemm_scale_softmax_gemm_permute_xdl_fp16
adding example example_batched_gemm_scale_softmax_gemm_permute_xdl_bf16
adding example example_grouped_gemm_scale_softmax_gemm_permute_xdl_fp16
adding example example_batched_gemm_lower_triangle_scale_softmax_gemm_permute_xdl_fp16
adding example example_grouped_gemm_lower_triangle_scale_softmax_gemm_permute_xdl_fp16
adding example example_dual_reduce_multiblock
adding example example_dual_reduce_threadwise
adding example example_batchnorm_forward_training
adding example example_batchnorm_forward_inferring
adding example example_batchnorm_backward
adding example example_sparse_embedding3_forward_layernorm
adding example example_batched_gemm_add_add_relu_gemm_add_xdl_fp16
adding example example_permute_1xHxW_fp16
adding example example_permute_NxHxW_fp16
adding example example_permute_HxWx4_fp16
adding example example_conv2d_fwd_dl_perlayer_quantization_int8
adding example example_conv2d_fwd_dl_perchannel_quantization_int8
adding example example_conv2d_fwd_dl_bias_relu_perlayer_quantization_int8
adding example example_conv2d_fwd_dl_bias_relu_perchannel_quantization_int8
adding example example_conv2d_fwd_dl_bias_tanh_perlayer_quantization_int8
adding example example_conv2d_fwd_dl_bias_tanh_perchannel_quantization_int8
adding example example_groupnorm_sigmoid_mul_fp16
adding example example_groupnorm_splitk_fp16
adding example example_groupnorm_swish_fp16
adding example example_splitk_gemm_bias_e_permute_xdl_fp16
adding example example_splitk_gemm_bias_e_permute_xdl_fp32
adding example example_elementwise_permute_4D_fp16
adding example example_elementwise_permute_4D_fp16_2d
adding example example_elementwise_layernorm_blockwise
adding example example_gemm_add_multiply_dl_fp16
adding example example_gemm_add_multiply_xdl_fp16
adding example example_pool3d_fwd_fp16
adding example example_maxpool2d_bwd_bf16
adding example example_maxpool2d_bwd_fp16
adding example example_maxpool2d_bwd_fp32
adding example example_put_element_fp16
-- Fetching GoogleTest
-- Suppressing googltest warnings with flags: -Wno-undef;-Wno-reserved-identifier;-Wno-global-constructors;-Wno-missing-noreturn;-Wno-disabled-macro-expansion;-Wno-used-but-marked-unused;-Wno-switch-enum;-Wno-zero-as-null-pointer-constant;-Wno-unused-member-function;-Wno-comma;-Wno-old-style-cast;-Wno-deprecated;-Wno-unsafe-buffer-usage
-- Found Python: /usr/local/bin/python3 (found version "3.8.10") found components: Interpreter 
adding test test_magic_number_division
adding test test_space_filling_curve
adding gtest test_conv_util
adding gtest test_reference_conv_fwd
adding test test_gemm_fp32
adding test test_gemm_fp16
adding test test_gemm_bf16
adding test test_gemm_int8
adding test test_gemm_standalone_xdl_fp16
adding test test_gemm_reduce_fp16
adding test test_reduce_no_index
adding test test_reduce_with_index
adding gtest test_grouped_convnd_fwd
adding gtest test_block_to_ctile_map
adding gtest test_softmax_rank3
adding gtest test_softmax_rank4
adding gtest test_softmax_interface
adding gtest test_layernorm2d_fp32
adding gtest test_layernorm2d_fp16
adding gtest test_groupnorm_fp16
adding gtest test_groupnorm_fp32
adding gtest test_fp8
adding gtest test_elementwise_layernorm_fp16
adding gtest test_batchnorm_fwd_rank_4
adding gtest test_batchnorm_bwd_rank_4
adding gtest test_batchnorm_infer_rank_4
adding gtest test_contraction
adding gtest test_avg_pool2d_fwd
adding gtest test_avg_pool3d_fwd
adding gtest test_max_pool2d_fwd
adding gtest test_max_pool3d_fwd
adding gtest test_batched_gemm_multi_d
RPM version 4.14.2.1
-- Configuring done
-- Generating done
-- Build files have been written to: /home/tyra/rocm/composable_kernel/build

Scanning dependencies of target device_softmax_instance
[  0%] Building CXX object library/src/tensor_operation_instance/gpu/softmax/CMakeFiles/device_softmax_instance.dir/device_softmax_f16_f16_instance_rank3_reduce1.cpp.o
In file included from /home/tyra/rocm/composable_kernel/library/src/tensor_operation_instance/gpu/softmax/device_softmax_f16_f16_instance_rank3_reduce1.cpp:9:
In file included from /home/tyra/rocm/composable_kernel/library/include/ck/library/tensor_operation_instance/gpu/softmax/device_softmax_f16_f16_instance_type.hpp:8:
In file included from /home/tyra/rocm/composable_kernel/include/ck/tensor_operation/gpu/device/impl/device_softmax_impl.hpp:12:
In file included from /home/tyra/rocm/composable_kernel/include/ck/tensor_operation/gpu/device/impl/device_reduce_common.hpp:9:
In file included from /home/tyra/rocm/composable_kernel/include/ck/utility/common_header.hpp:36:
/home/tyra/rocm/composable_kernel/include/ck/utility/amd_buffer_addressing.hpp:32:48: error: use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD'
    wave_buffer_resource.config(Number<3>{}) = CK_BUFFER_RESOURCE_3RD_DWORD;
                                               ^
/home/tyra/rocm/composable_kernel/include/ck/utility/amd_buffer_addressing.hpp:47:48: error: use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD'
    wave_buffer_resource.config(Number<3>{}) = CK_BUFFER_RESOURCE_3RD_DWORD;
                                               ^
2 errors generated when compiling for gfx1010.
make[2]: *** [library/src/tensor_operation_instance/gpu/softmax/CMakeFiles/device_softmax_instance.dir/build.make:82: library/src/tensor_operation_instance/gpu/softmax/CMakeFiles/device_softmax_instance.dir/device_softmax_f16_f16_instance_rank3_reduce1.cpp.o] Error 1
make[1]: *** [CMakeFiles/Makefile2:14807: library/src/tensor_operation_instance/gpu/softmax/CMakeFiles/device_softmax_instance.dir/all] Error 2
make: *** [Makefile:182: all] Error 2

Any ideas about a solution?

@Tyravex Thanks for reporting this issue. Please let us know if it works or not. We do not have Navi10 GPUs to test it.

In the
5.7 branch

There is missing support for gfx941, gfx942, gfx1012 and gfx1030
So the only generally useful set is the gfx11.

Folks will be unhappily surprised when MIOpen and similar users of CK fail.

Our experiments at Solus have found the following:

gfx803, gfx900, gfx906 can all build fine, just remember to build with BUILD_DEV=OFF because compiling those architectures may produce warnings that subsequently get treated by the compiler as errors due to -Werror under BUILD_DEV=ON. We haven't run any tests on those platforms (yet), so for now we only know that they build 😅
gfx1010 is probably best to be treated as gfx900, because even though gfx1030 is theoretically the closest, there are optimizations and LLVM builtins that gfx1010 cannot use while gfx1030 can, such as __builtin_amdgcn_fdot2.

ROCm-composable_kernel: Treat gfx1010 as gfx900 for CK_BUFFER_RESOURC…

@GZGavinZhao

gfx1010 is probably best to be treated as gfx900, because even though gfx1030 is theoretically the closest, there are optimizations and LLVM builtins that gfx1010 cannot use while gfx1030 can, such as __builtin_amdgcn_fdot2.

@GZGavinZhao Its interesting that the __builtin_amdgcn_fdot2 and __builtin_amdgcn_sdot4 features are not supported, maybe we can just add them to the features list (i digged it up here) https://github.com/llvm/llvm-project/blob/8a79dc7e6f765f3f49c5dd9330fc0826d3362858/llvm/lib/Target/AMDGPU/AMDGPU.td#L1496

So if I could just get to understand that we are hitting a hardware limit with gfx1010 (RX 5700 and the likes, Navi10 processor) that it can not do what the gfx1030 (RX6900 and the likes Navi21,23) is doing or its just someones idea for business strategy that gfx1010 is too old to support or whats the case? if its a hardware limitation where did we hit a dead end? why the software is not shipped as well as for gfx1010, im just trying to understand what is going on, whos deciding the stuff?

I know there have been a tremendous amount of efforts all over the place which i was not part of, and I appreciate everyone contribution to where things have gone so far and very thankful to everyone's effort but we need to fix this issue, old GPUS are not going to the trash man (if the plan is to send them to the trash, then too bad of a waste)!

I just read some main instruction set parts for memory shader and some other registers https://www.amd.com/content/dam/amd/en/documents/radeon-tech-docs/instruction-set-architectures/rdna-shader-instruction-set-architecture.pdf both AMD-GCN-GFX10-RDNA2 and AMD-GCN-GFX10-RDNA1 seems to have an identical instructions sets, I understand there are some details i missed, but from a hardware point of view, everything seems to be existing already in gfx1010 so our problem is software, so probably just copying the gfx1030 stuff and use it for gfx1010 and maybe without modifying anything (or with modifiying few things which is a nightmare:)) we would get a better support for that Navi10 (that gddr6 monster can do crazy stuff, so people are using older software by adopting to the gfx900 arch, well its still a toy in comparison to the MI300X)

the problem is that this is not an easy job to do, it should have been considered earlier but we can maybe still do something all together, as the projects involved maybe beyond my capabilities, like:

1-LLVM
2- rocBlas
3- CK
4- MIOpen

I just hit this error while trying to add the gfx1030 features to gfx1010,

Cached asm caps differ from derived asm caps for (10, 1, 0)

which is expected, as I modified the AMDGPU Target for that arch, so now my question, which one is the cached asm caps and which one derived from where?

I passed through the error by forking the Tensile repo and IgnoringAsmCap mismatch

Could you raise an issue in Tensile for this ?
Fedora will use the upstream Tensile, not a copy.

@trixirt

Thanks @trixirt , yes, I will also add this issue to it:

below is actually the compileArgs command, and the first element is None, which i guess is the cxx_compiler(it should not be None):

printing our args as we have a None Type error at the moment [None, '-D__HIP_HCC_COMPAT_MODE__=1', '--cuda-device-only', '-x', 'hip', '-O3', '-I', '/home/bargo/projects/rocm-setup/rocBLAS/build/Tensile', '-Xoffload-linker', '--build-id', '--offload-arch=gfx1010', '--offload-arch=gfx803', '--offload-arch=gfx1030', '/home/bargo/projects/rocm-setup/rocBLAS/build/Tensile/TensileLibrary_Type_ZZ_Contraction_l_Ailk_BjlkC_Cijk_Dijk_fallback.cpp', '-c', '-o', '/home/bargo/projects/rocm-setup/rocBLAS/build/library/src/build_tmp/TENSILE/code_object_tmp/TensileLibrary_Type_ZZ_Contraction_l_Ailk_BjlkC_Cijk_Dijk_fallback.o'] launcher: [] hipFlags: ['-D__HIP_HCC_COMPAT_MODE__=1', '--cuda-device-only', '-x', 'hip', '-O3', '-I', '/home/bargo/projects/rocm-setup/rocBLAS/build/Tensile', '-Xoffload-linker', '--build-id'] archFlags: ['--offload-arch=gfx1010', '--offload-arch=gfx803', '--offload-arch=gfx1030']

I agree this issue/discussion should continue in Tensile, but just to add my 2 cents:

I just read some main instruction set parts for memory shader and some other registers https://www.amd.com/content/dam/amd/en/documents/radeon-tech-docs/instruction-set-architectures/rdna-shader-instruction-set-architecture.pdf both AMD-GCN-GFX10-RDNA2 and AMD-GCN-GFX10-RDNA1 seems to have an identical instructions sets

RDNA1 and RDNA2 are not the same instruction set. For example, consult section 6.3 of the RDNA1 and RDNA2 ISA architecture. You will see that RDNA2 has V_DOT* instructions like V_DOT2_F32_F16 that RDNA1 doesn't have. Moreover, aside from instruction differences, there are fundamentally different designs in RDNA1 and RDNA2 hardware that makes code compiled for the latter unable to run on the former. For example, the different number of registers. It's possible you may get lucky and the code you compiled against RDNA2 doesn't use the extra registers at all, so it runs fine on RDNA1, but that's pure luck and there's no way we can guarantee that.

So if I could just get to understand that we are hitting a hardware limit with gfx1010 (RX 5700 and the likes, Navi10 processor) that it can not do what the gfx1030 (RX6900 and the likes Navi21,23) is doing or its just someones idea for business strategy that gfx1010 is too old to support or whats the case? if its a hardware limitation where did we hit a dead end? why the software is not shipped as well as for gfx1010, im just trying to understand what is going on, whos deciding the stuff?

As per my reasoning above, in my understanding it is a hardware limit that gfx1010 cannot do the same as gfx1030. This makes sense anyway, because why would RDNA2 even exist in the first place if it's just more FLOPS and no new instructions/optimizations?

Therefore, any attempt at trying to compile and run code intended for RDNA2 on RDNA1 will eventually fail. The best we can do (and what it's already been done at Solus) is to patch libraries like rocBLAS, Tensile, and CK so that they can compile and run RDNA1 code on a best-effort basis. I believe Solus's support for RDNA1 hardware is complete. I assume you own a RDNA1 hardware, so feel free to grab a Solus ISO, install the pytorch package, give it a try (it should work out of the box, no environment variables or anything needed), and let me know or open an issue at the Solus repo if anything doesn't work.

@GZGavinZhao

Thanks for having a look @GZGavinZhao appreciate all your efforts.

RDNA1 and RDNA2 are not the same instruction set. For example, consult section 6.3 of the RDNA1 and RDNA2 ISA architecture. You will see that RDNA2 has V_DOT* instructions like V_DOT2_F32_F16 that RDNA1 doesn't have. Moreover, aside from instruction differences, there are fundamentally different designs in RDNA1 and RDNA2 hardware that makes code compiled for the latter unable to run on the former. For example, the different number of registers. It's possible you may get lucky and the code you compiled against RDNA2 doesn't use the extra registers at all, so it runs fine on RDNA1, but that's pure luck and there's no way we can guarantee that.

Actually regarding the above I have a feeling that we can build these instructions for the gfx1010, you referred to the VOP3P instruction which can also be built for that target, it could be all what im doing is none sense, but I like trying stuff and let me see where that takes me :) I will keep you posted either way.

Therefore, any attempt at trying to compile and run code intended for RDNA2 on RDNA1 will eventually fail. The best we can do (and what it's already been done at Solus) is to patch libraries like rocBLAS, Tensile, and CK so that they can compile and run RDNA1 code on a best-effort basis. I believe Solus's support for RDNA1 hardware is complete. I assume you own a RDNA1 hardware, so feel free to grab a Solus ISO, install the pytorch package, give it a try (it should work out of the box, no environment variables or anything needed), and let me know or open an issue at the Solus repo if anything doesn't work.

i think Solus is definitely helpful but currently running the gfx1010 on gfx900 is a downgrade.

It would be my last resort to do that nevertheless I really like what you guys are doing at Solus looks a good approach for a better GPU support,

I have plenty of GPUs as I work as a repairman for myself. And thus i ended up with plenty of Graphic cards from all generations. So i need to be fair to myself :)

i think Solus is definitely helpful but currently running the gfx1010 on gfx900 is a downgrade.

Note that only for CK, gfx1010 has the same code path as gfx900. For Tensile/rocBLAS, gfx1010 is being compiled normally, so it receives all the optimizations that gfx1010 should get by default.

I do agree that for CK specifically, we may be able to treat gfx1010 as gfx1030 in general and then workaround specific cases where VDOT instructions are used. However, for all other libraries, you shouldn't need to do anything; gfx1010 should just compile and run normally, at least that's the case for Solus with ROCm 5.7, 6.0, and 6.1.

Note that only for CK, gfx1010 has the same code path as gfx900. For Tensile/rocBLAS, gfx1010 is being compiled normally, so it receives all the optimizations that gfx1010 should get by default.

I do agree that for CK specifically, we may be able to treat gfx1010 as gfx1030 in general and then workaround specific cases where VDOT instructions are used. However, for all other libraries, you shouldn't need to do anything; gfx1010 should just compile and run normally, at least that's the case for Solus with ROCm 5.7, 6.0, and 6.1.

Yes I noticed that, and that what brought me here :) I think there is a room of improvement in general :) the repos are huge and there are a lot of stuff like 20 repos :D

Right now I'm just reading some weird stuff like this:
defm V_DOT2_F32_F16 : VOP3P_Real_gfx10_gfx11_gfx12_Triple<0x13>; this is basically in some llvm file I think this one, llvm-project/llvm/lib/Target/AMDGPU/VOP3PInstructions.td and it shows that this V_DOT2_F32_F16 Which means its already part of that VOP3P Instructions set I dont understand that language very well, very complex for my heart

Hi all, sorry for the lack of official response here. The community is doing great work to preserve compatibility with GPUs that have lost official support.

As noted in this discussion, ROCm does not support Navi 10 GPUs. However, we do have another macro in include/ck/ck.hpp which is defined for gfx1010, so I don't think it's unreasonable to also define CK_BUFFER_RESOURCE_3RD_DWORD and other macros for gfx1010 if there are appropriate values for them (and similarly for other architectures where this is not currently defined); I'll check with the internal team to see if this is something we want to do. We don't extensively test for unsupported architectures, so even if we enable support here we cannot guarantee that this will function correctly.

Regarding submitting PRs to address issues on unsupported architectures, I will also discuss this with the internal team. In general you are more than welcome to submit PRs for any improvements to the ROCm stack, but since we won't be testing for unsupported architectures I can't say how likely it is that fixes for these architectures get merged. I encourage you to document your fixes somewhere (whether here, in a PR, fork, etc.) so future users with similar issues can refer to them.

@SzczurekYT

@SzczurekYT

ROCm 6.2.x: https://drive.google.com/file/d/1RFbfYtG0B0JbtTai9iWdVomammAFL8Db/view
ROCm 6.3.x: https://drive.google.com/file/d/1lC2MEgXbE85IgnKlvVug3R6lPypy1ATl/view

The attached diff file synthesizes this issue and the community patches from rocm_sdk_bulider. The latter is a repository that has many fixes and features to introduce ROCm functionality on several architectures, including APUs. However, the data-center ones (CDNA) have limited testing and are based on older versions of ROCm. This file is designed with the latest point release in mind, which is 6.2 as of writing.

Note that this only restores compilation success for RDNA1 discrete GPUs (RX 5500/5600/5700), so having gfx1010, gfx1011, or gfx1012 as a target will work but will be fairly slow to compile. I have not figured out how to target just for these architectures, as others will as well. Feel free to report any issues regarding this file.

I've spoken with the internal team, and our stance is that PRs are encouraged, and can be accepted even for fixes on unsupported architectures as long as they don't break existing test cases. However, we might not have the resources to review these PRs in a timely manner, so I understand if you would rather provide code fixes via other means, and at the end of the day the fixes/guidance are available to future users either way.

Since the official ROCm team accepts pull requests for unsupported architectures from the community, and merging may be delayed due to time constraints, this is very good news for improving the user experience.

I have tested my patch and can assure you both composable kernels and MIOpen will compile if you use ck's -DDL_KERNELS or -DCK_ENABLE_DL_KERNELS=ON build flags, thus resolving the root issue. However, I have not run any follow-up unit tests to ensure correctness. The last time I checked with my build environment, only two failed verification as seen a few posts above without these flags. I also attached error logs where it went wrong to help guide developers to a solution.

Currently working on a PR to fix this. Code has been written, just tests pending.

CTest results for ROCm 6.3 on gfx1012. It seemed to regress significantly. I saw lots of NaNs and zeros after rerunning the failed tests.

Test project /root/rocm-build/composable-kernel/test
      Start  1: test_magic_number_division
 1/50 Test  #1: test_magic_number_division ...................   Passed    4.78 sec
      Start  2: test_space_filling_curve
 2/50 Test  #2: test_space_filling_curve .....................   Passed    0.02 sec
      Start  3: test_conv_util
 3/50 Test  #3: test_conv_util ...............................   Passed    0.02 sec
      Start  4: test_reference_conv_fwd
 4/50 Test  #4: test_reference_conv_fwd ......................   Passed    0.02 sec
      Start  5: test_gemm_fp32
 5/50 Test  #5: test_gemm_fp32 ...............................   Passed    1.24 sec
      Start  6: test_gemm_fp16
 6/50 Test  #6: test_gemm_fp16 ...............................   Passed  120.30 sec
      Start  7: test_gemm_bf16
 7/50 Test  #7: test_gemm_bf16 ...............................   Passed    0.13 sec
      Start  8: test_gemm_int8
 8/50 Test  #8: test_gemm_int8 ...............................   Passed   33.31 sec
      Start  9: test_reduce_no_index
 9/50 Test  #9: test_reduce_no_index .........................***Failed   46.65 sec
      Start 10: test_reduce_with_index
10/50 Test #10: test_reduce_with_index .......................***Failed   47.98 sec
      Start 11: test_grouped_convnd_fwd_multi_ab_interface
11/50 Test #11: test_grouped_convnd_fwd_multi_ab_interface ...   Passed    0.03 sec
      Start 12: test_grouped_convnd_bwd_weight
12/50 Test #12: test_grouped_convnd_bwd_weight ...............***Failed    8.71 sec
      Start 13: test_block_to_ctile_map
13/50 Test #13: test_block_to_ctile_map ......................   Passed    0.02 sec
      Start 14: test_softmax_rank3
14/50 Test #14: test_softmax_rank3 ...........................***Failed   45.97 sec
      Start 15: test_softmax_rank4
15/50 Test #15: test_softmax_rank4 ...........................***Failed   97.00 sec
      Start 16: test_softmax_interface
16/50 Test #16: test_softmax_interface .......................   Passed    0.05 sec
      Start 17: test_layernorm2d_fwd_fp32
17/50 Test #17: test_layernorm2d_fwd_fp32 ....................   Passed    0.71 sec
      Start 18: test_groupnorm_fwd_fp32
18/50 Test #18: test_groupnorm_fwd_fp32 ......................   Passed    0.58 sec
      Start 19: test_layernorm2d_fwd_fp16
19/50 Test #19: test_layernorm2d_fwd_fp16 ....................***Failed    0.54 sec
      Start 20: test_layernorm4d_fwd_fp16
20/50 Test #20: test_layernorm4d_fwd_fp16 ....................***Failed    0.26 sec
      Start 21: test_groupnorm_fwd_fp16
21/50 Test #21: test_groupnorm_fwd_fp16 ......................***Failed    0.57 sec
      Start 22: test_layernorm2d_bwd_data_fp32
22/50 Test #22: test_layernorm2d_bwd_data_fp32 ...............   Passed    0.44 sec
      Start 23: test_groupnorm_bwd_data_fp32
23/50 Test #23: test_groupnorm_bwd_data_fp32 .................   Passed    0.38 sec
      Start 24: test_layernorm2d_bwd_gamma_beta_fp32
24/50 Test #24: test_layernorm2d_bwd_gamma_beta_fp32 .........   Passed    0.35 sec
      Start 25: test_groupnorm_bwd_gamma_beta_fp32
25/50 Test #25: test_groupnorm_bwd_gamma_beta_fp32 ...........   Passed    0.34 sec
      Start 26: test_fp8
26/50 Test #26: test_fp8 .....................................   Passed    0.02 sec
      Start 27: test_bf8
27/50 Test #27: test_bf8 .....................................   Passed    0.02 sec
      Start 28: test_custom_type
28/50 Test #28: test_custom_type .............................***Failed    0.02 sec
      Start 29: test_type_convert_const
29/50 Test #29: test_type_convert_const ......................   Passed    0.02 sec
      Start 30: test_elementwise_layernorm_fp16
30/50 Test #30: test_elementwise_layernorm_fp16 ..............***Failed    2.84 sec
      Start 31: test_batchnorm_fwd_rank_4
31/50 Test #31: test_batchnorm_fwd_rank_4 ....................***Failed   43.66 sec
      Start 32: test_batchnorm_bwd_rank_4
32/50 Test #32: test_batchnorm_bwd_rank_4 ....................***Failed   60.01 sec
      Start 33: test_batchnorm_infer_rank_4
33/50 Test #33: test_batchnorm_infer_rank_4 ..................***Failed    9.79 sec
      Start 34: test_avg_pool3d_bwd
34/50 Test #34: test_avg_pool3d_bwd ..........................***Failed    0.60 sec
      Start 35: test_max_pool3d_bwd
35/50 Test #35: test_max_pool3d_bwd ..........................***Failed    0.80 sec
      Start 36: test_avg_pool3d_fwd
36/50 Test #36: test_avg_pool3d_fwd ..........................***Failed    2.35 sec
      Start 37: test_max_pool3d_fwd
37/50 Test #37: test_max_pool3d_fwd ..........................***Failed    5.01 sec
      Start 38: test_avg_pool2d_bwd
38/50 Test #38: test_avg_pool2d_bwd ..........................***Failed    0.25 sec
      Start 39: test_max_pool2d_bwd
39/50 Test #39: test_max_pool2d_bwd ..........................***Failed    0.27 sec
      Start 40: test_avg_pool2d_fwd
40/50 Test #40: test_avg_pool2d_fwd ..........................***Failed    0.35 sec
      Start 41: test_max_pool2d_fwd
41/50 Test #41: test_max_pool2d_fwd ..........................***Failed    0.31 sec
      Start 42: test_batched_gemm_multi_d
42/50 Test #42: test_batched_gemm_multi_d ....................   Passed    0.64 sec
      Start 43: test_conv_tensor_rearrange
43/50 Test #43: test_conv_tensor_rearrange ...................***Failed   13.57 sec
      Start 44: test_conv_tensor_rearrange_interface
44/50 Test #44: test_conv_tensor_rearrange_interface .........   Passed    0.03 sec
      Start 45: test_permute_scale
45/50 Test #45: test_permute_scale ...........................***Failed    0.40 sec
      Start 46: test_wrapper_layout
46/50 Test #46: test_wrapper_layout ..........................   Passed    0.02 sec
      Start 47: test_wrapper_tensor
47/50 Test #47: test_wrapper_tensor ..........................   Passed    0.23 sec
      Start 48: test_wrapper_copy
48/50 Test #48: test_wrapper_copy ............................   Passed    0.22 sec
      Start 49: test_wrapper_partition
49/50 Test #49: test_wrapper_partition .......................   Passed    0.02 sec
      Start 50: test_position_embedding
50/50 Test #50: test_position_embedding ......................   Passed    0.02 sec

54% tests passed, 23 tests failed out of 50

Total Test time (real) = 551.92 sec

The following tests FAILED:
	  9 - test_reduce_no_index (Failed)
	 10 - test_reduce_with_index (Failed)
	 12 - test_grouped_convnd_bwd_weight (Failed)
	 14 - test_softmax_rank3 (Failed)
	 15 - test_softmax_rank4 (Failed)
	 19 - test_layernorm2d_fwd_fp16 (Failed)
	 20 - test_layernorm4d_fwd_fp16 (Failed)
	 21 - test_groupnorm_fwd_fp16 (Failed)
	 28 - test_custom_type (Failed)
	 30 - test_elementwise_layernorm_fp16 (Failed)
	 31 - test_batchnorm_fwd_rank_4 (Failed)
	 32 - test_batchnorm_bwd_rank_4 (Failed)
	 33 - test_batchnorm_infer_rank_4 (Failed)
	 34 - test_avg_pool3d_bwd (Failed)
	 35 - test_max_pool3d_bwd (Failed)
	 36 - test_avg_pool3d_fwd (Failed)
	 37 - test_max_pool3d_fwd (Failed)
	 38 - test_avg_pool2d_bwd (Failed)
	 39 - test_max_pool2d_bwd (Failed)
	 40 - test_avg_pool2d_fwd (Failed)
	 41 - test_max_pool2d_fwd (Failed)
	 43 - test_conv_tensor_rearrange (Failed)
	 45 - test_permute_scale (Failed)
Errors while running CTest
Output from these tests are in: /root/rocm-build/composable-kernel/test/Testing/Temporary/LastTest.log
Use "--rerun-failed --output-on-failure" to re-run the failed cases verbosely.

Good grief AWS finally approved my quota limit increase so I now have access to a RDNA1 GPU in the cloud. I'm currently debugging these failures.

Everything fixed instead of one single test failure:

[ RUN      ] TestBatchNormInferRank4/3.nchw
found 9 instances
y results        out[1507712] != ref[1507712]: 0.6240114 != 1.241642
y results        out[1507716] != ref[1507716]: 0.6240114 != -0.9240708
y results        out[1507720] != ref[1507720]: 0.6240114 != 0.5014129
y results        out[1507724] != ref[1507724]: 0.6240114 != 1.268496
max err: 6.406953e+38, number of errors: 96, 0.001525879% wrong values
/home/ubuntu/composable_kernel/test/batchnorm/batchnorm_infer_rank_4.cpp:65: Failure
Value of: pass
  Actual: false
Expected: true

found 9 instances
y results        out[16576] != ref[16576]: -3.343448e+38 != -0.6865153
y results        out[16580] != ref[16580]: 0.2960358 != 0.2241838
y results        out[16584] != ref[16584]: 0.2960358 != -0.3484504
y results        out[16588] != ref[16588]: 0.2960358 != -0.1144298
max err: 5.055431e+38, number of errors: 190, 0.003019969% wrong values
/home/ubuntu/composable_kernel/test/batchnorm/batchnorm_infer_rank_4.cpp:65: Failure
Value of: pass
  Actual: false
Expected: true

0.003% of wrong values... 😅

@GZGavinZhao

@GZGavinZhao Do you have your patches available to download? I have a gfx1012 and would like to get this working on my system for use in the interim.

@mattmcadoo

@mattmcadoo The following three (apply in the given order) should suffice:

@GZGavinZhao

@GZGavinZhao Thank you! Those patches allowed me to finally compile composable_kernel 6.3.0 on my Gentoo system.

sci-libs/composable-kernel: fix compilation with AMDGPU_TARGETS="gfx1…

Navigation Menu

Repository navigation

Compilation error for navi10 (use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD') #775

26 remaining items

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Compilation error for navi10 (use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD') #775

Description

Activity

zjing14 commented on Jun 30, 2023

trixirt commented on Oct 15, 2023

GZGavinZhao commented on Jan 4, 2024

waheedi commented on Jul 19, 2024

waheedi commented on Jul 22, 2024

trixirt commented on Jul 22, 2024

waheedi commented on Jul 22, 2024

GZGavinZhao commented on Jul 22, 2024

waheedi commented on Jul 22, 2024

GZGavinZhao commented on Jul 22, 2024

waheedi commented on Jul 22, 2024

waheedi commented on Jul 22, 2024

26 remaining items

schung-amd commented on Nov 28, 2024

TheTrustedComputer commented on Nov 29, 2024

schung-amd commented on Nov 29, 2024

TheTrustedComputer commented on Dec 2, 2024

GZGavinZhao commented on Dec 3, 2024

TheTrustedComputer commented on Dec 15, 2024

GZGavinZhao commented on Dec 30, 2024

GZGavinZhao commented on Jan 17, 2025

mattmcadoo commented on Mar 13, 2025

GZGavinZhao commented on Mar 14, 2025

mattmcadoo commented on Mar 15, 2025

Metadata

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Issue actions