# Copyright (c) Huawei Technologies Co., Ltd. 2025. All rights reserved.

#

# Permission is hereby granted, free of charge, to any person obtaining a copy

# of this software and associated documentation files (the "Software"), to deal

# in the Software without restriction, including without limitation the rights

# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell

# copies of the Software, and to permit persons to whom the Software is

# furnished to do so, subject to the following conditions:

#

# The above copyright notice and this permission notice shall be included in

# all copies or substantial portions of the Software.

#

# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR

# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,

# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE

# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER

# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,

# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN

# THE SOFTWARE.



import triton

import triton.language as tl

import triton.language.extra.cann.libdevice as libdevice

import torch

import pytest

import test_common

import os



PERF_TEST_ENABLE = os.getenv('PERF_TEST_ENABLE', 'False').lower() == 'true'



def torch_uint_as_float(x):

    return x.view(torch.int32).view(torch.float32)



@triton.jit

def triton_uint_as_float(in_ptr0, out_ptr0, XBLOCK: tl.constexpr, XBLOCK_SUB: tl.constexpr):

    offset = tl.program_id(0) * XBLOCK

    base1 = tl.arange(0, XBLOCK_SUB)

    loops1: tl.constexpr = (XBLOCK + XBLOCK_SUB - 1) // XBLOCK_SUB

    for loop1 in range(loops1):

        x0 = offset + (loop1 * XBLOCK_SUB) + base1

        tmp0 = tl.load(in_ptr0 + (x0), None)

        tmp1 = libdevice.uint_as_float(tmp0)

        tl.store(out_ptr0 + (x0), tmp1, None)



default_param_list = test_common.make_default_param_list(['uint32'])



full_param_list = test_common.make_full_param_list(['uint32'])



@pytest.mark.parametrize(

    'param_list',

    default_param_list if not PERF_TEST_ENABLE else full_param_list,

)

def test_common_case(param_list):

    dtype, shape, ncore, xblock, xblock_sub = param_list

    np_x0 = test_common.generate_numpy(shape, dtype)

    x0 = torch.from_numpy(np_x0).to(eval('torch.' + dtype)).npu()

    y_ref = torch_uint_as_float(x0)

    y_cal = torch.zeros(shape, dtype = torch.float32).npu()

    if PERF_TEST_ENABLE:

        test_common.run_with_profiler(

            lambda: triton_uint_as_float[ncore, 1, 1](x0, y_cal, xblock, xblock_sub, force_simt_only=True),

            shape,

            'uint_as_float'

        )

    else:

        triton_uint_as_float[ncore, 1, 1](x0, y_cal, xblock, xblock_sub, force_simt_only=True)

    test_common.validate_cmp('float32', y_cal, y_ref)



@pytest.mark.parametrize('param_list',

                        [

                            ['uint32', (1, 16), 1, 16, 16],

                        ]

                        )

def test_special_case(param_list):

    dtype, shape, ncore, xblock, xblock_sub = param_list

    np_x0 = test_common.generate_numpy(shape, dtype)

    x0 = torch.from_numpy(np_x0).to(eval('torch.' + dtype)).npu()

    x0 = x0.view(torch.int32)

    x0[0, 0] = 0

    x0[0, 1] = 0xffff

    x0 = x0.view(torch.uint32)

    y_ref = torch_uint_as_float(x0)

    y_cal = torch.zeros(shape, dtype = torch.float32).npu()

    triton_uint_as_float[ncore, 1, 1](x0, y_cal, xblock, xblock_sub, force_simt_only=True)

    test_common.validate_cmp('float32', y_cal, y_ref)