About amd gpu Issue #527

[Lookforworld]

@anthonix

Hi, I'm here.

Is it possible that this error is due to the Rocm version or torch?

[anthonix]

Which ROCm version are you running? It should be 6.1.1

[Lookforworld]

Yes!

Package: rocm-libs
Version: 6.1.1.60101-90~22.04
Priority: optional
Section: devel
Maintainer: ROCm Dev Support <rocm-dev.support@amd.com>
Installed-Size: 13.3 kB
Depends: hipblas (= 2.1.0.60101-90~22.04), hipblaslt (= 0.7.0.60101-90~22.04), hipfft (= 1.0.14.60101-90~22.04), hipsolver (= 2.1.1.60101-90~22.04), hipsparse (= 3.0.1.60101-90~22.04), hiptensor (= 1.2.0.60101-90~22.04), miopen-hip (= 3.1.0.60101-90~22.04), half (= 1.12.0.60101-90~22.04), rccl (= 2.18.6.60101-90~22.04), rocalution (= 3.1.1.60101-90~22.04), rocblas (= 4.1.0.60101-90~22.04), rocfft (= 1.0.27.60101-90~22.04), rocrand (= 3.0.1.60101-90~22.04), hiprand (= 2.10.16.60101-90~22.04), rocsolver (= 3.25.0.60101-90~22.04), rocsparse (= 3.1.2.60101-90~22.04), rocm-core (= 6.1.1.60101-90~22.04), hipsparselt (= 0.1.0.60101-90~22.04), composablekernel-dev (= 1.1.0.60101-90~22.04), hipblas-dev (= 2.1.0.60101-90~22.04), hipblaslt-dev (= 0.7.0.60101-90~22.04), hipcub-dev (= 3.1.0.60101-90~22.04), hipfft-dev (= 1.0.14.60101-90~22.04), hipsolver-dev (= 2.1.1.60101-90~22.04), hipsparse-dev (= 3.0.1.60101-90~22.04), hiptensor-dev (= 1.2.0.60101-90~22.04), miopen-hip-dev (= 3.1.0.60101-90~22.04), rccl-dev (= 2.18.6.60101-90~22.04), rocalution-dev (= 3.1.1.60101-90~22.04), rocblas-dev (= 4.1.0.60101-90~22.04), rocfft-dev (= 1.0.27.60101-90~22.04), rocprim-dev (= 3.1.0.60101-90~22.04), rocrand-dev (= 3.0.1.60101-90~22.04), hiprand-dev (= 2.10.16.60101-90~22.04), rocsolver-dev (= 3.25.0.60101-90~22.04), rocsparse-dev (= 3.1.2.60101-90~22.04), rocthrust-dev (= 3.0.1.60101-90~22.04), rocwmma-dev (= 1.4.0.60101-90~22.04), hipsparselt-dev (= 0.1.0.60101-90~22.04)
Homepage: https://github.com/RadeonOpenCompute/ROCm
Download-Size: 1,060 B
APT-Sources: https://repo.radeon.com/rocm/apt/6.1.1 jammy/main amd64 Packages
Description: Radeon Open Compute (ROCm) Runtime software stack

[Lookforworld]

@anthonix
I checked the train_gpt2.hip file and the error was given on the last line of this function. However, it is not specified that there is a problem with the specific operator.

void fused_residual_forward5(floatX* residual, floatX* normed, floatX* mean, floatX* rstd,
                             const floatX* inp1, const floatX* inp2,
                             const floatX* weight, const floatX* bias,
                             int N, int C) {
    const int block_size = 256;
    int block_y = block_size / WARP_SIZE;
    const int grid_size = CEIL_DIV(N, block_y);
    size_t smem = (2 + block_y) * C * sizeof(floatX);

    // in order to use more than 48 KiB of smem, need to call hipFuncSetAttribute
    // this may fail, in which case we fall back to the smem free implementation.
    cudaCheck(hipGetLastError());
    auto status = hipFuncSetAttribute(fused_residual_forward_kernel5, hipFuncAttributeMaxDynamicSharedMemorySize, smem);
    hipGetLastError();
    if(status == hipSuccess) {
        fused_residual_forward_kernel5<<<grid_size, dim3(WARP_SIZE, block_y), smem>>>(residual, normed, mean, rstd, inp1, inp2,
                                                                               weight, bias, N, C);
    } else {
        residual_forward(residual, inp1, inp2, N*C);
        layernorm_forward(normed, mean, rstd, residual, weight, bias, N, 1, C);
    }
    cudaCheck(hipGetLastError());
}