max min and equal operations

norch/autograd/functions.py

@@ -198,5 +198,23 @@ class CosBackward:
     def backward(self, gradient):
         x = self.input[0]
         return [-gradient * x.sin()]
+    
+class MaxBackward:
+    def __init__(self, x, axis=None, keepdim=False):
+        self.input = [x]
+        self.axis = axis
+        self.keepdim = keepdim
+
+    def backward(self, gradient):
+        pass
+
+class MinBackward:
+    def __init__(self, x, axis=None, keepdim=False):
+        self.input = [x]
+        self.axis = axis
+        self.keepdim = keepdim
+
+    def backward(self, gradient):
+        pass
 
 

norch/csrc/cpu.cpp

@@ -235,8 +235,76 @@ void sum_tensor_cpu(Tensor* tensor, float* result_data, int size, int* result_sh
     }
 }
 
+void max_tensor_cpu(Tensor* tensor, float* result_data, int size, int* result_shape, int axis) {
+    if (axis == -1) {
+        float max_value = -INFINITY;
+        for (int i = 0; i < tensor->size; i++) {
+            max_value = fmax(max_value, tensor->data[i]);
+        }
+        *result_data = max_value;
+    } else {
+        for (int i = 0; i < size; i++) {
+            result_data[i] = -INFINITY;
+        }
+        if (axis < 0 || axis >= tensor->ndim) {
+            printf("Invalid axis");
+            return;
+        }
+        
+        int axis_stride = tensor->strides[axis];
 
+        for (int i = 0; i < tensor->shape[axis]; i++) {
+            for (int j = 0; j < size; j++) {
+                int index = 0;
+                int remainder = j;
+                for (int k = tensor->ndim - 2; k >= 0; k--) {
+                    index += (remainder % result_shape[k]) * tensor->strides[k < axis ? k : k + 1];     
+                    remainder /= result_shape[k];
+                }
+                result_data[j] = fmax(result_data[j], tensor->data[index + i * axis_stride]);
+            }
+        }
+    }
+}
 
+void min_tensor_cpu(Tensor* tensor, float* result_data, int size, int* result_shape, int axis) {
+    if (axis == -1) {
+        float min_value = INFINITY;
+        for (int i = 0; i < tensor->size; i++) {
+            min_value = fmin(min_value, tensor->data[i]);
+        }
+        *result_data = min_value;
+    } else {
+        for (int i = 0; i < size; i++) {
+            result_data[i] = INFINITY;
+        }
+        if (axis < 0 || axis >= tensor->ndim) {
+            printf("Invalid axis");
+            return;
+        }
+        
+        int axis_stride = tensor->strides[axis];
+
+        for (int i = 0; i < tensor->shape[axis]; i++) {
+            for (int j = 0; j < size; j++) {
+                int index = 0;
+                int remainder = j;
+                for (int k = tensor->ndim - 2; k >= 0; k--) {
+                    index += (remainder % result_shape[k]) * tensor->strides[k < axis ? k : k + 1];     
+                    remainder /= result_shape[k];
+                }
+                result_data[j] = fmin(result_data[j], tensor->data[index + i * axis_stride]);
+            }
+        }
+    }
+}
+
+void equal_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data) {
+    
+    for (int i = 0; i < tensor1->size; i++) {
+        result_data[i] = (tensor1->data[i] == tensor2->data[i]) ? 1.0f : 0.0f;
+    }
+}
 
 
 void ones_like_tensor_cpu(Tensor* tensor, float* result_data) {

norch/csrc/cpu.h

@@ -6,6 +6,8 @@
 void add_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data);
 void add_broadcasted_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data, int* broadcasted_shape, int broadcasted_size);
 void sum_tensor_cpu(Tensor* tensor, float* result_data, int size, int* shape, int axis);
+void max_tensor_cpu(Tensor* tensor, float* result_data, int size, int* result_shape, int axis);
+void min_tensor_cpu(Tensor* tensor, float* result_data, int size, int* result_shape, int axis);
 void sub_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data);
 void sub_broadcasted_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data, int* broadcasted_shape, int broadcasted_size);
 void elementwise_mul_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data);
@@ -19,6 +21,7 @@ void scalar_pow_tensor_cpu(float base, Tensor* tensor, float* result_data);
 void tensor_pow_scalar_cpu(Tensor* tensor, float exponent, float* result_data);
 void log_tensor_cpu(Tensor* tensor, float* result_data);
 void scalar_mul_tensor_cpu(Tensor* tensor, float scalar, float* result_data);
+void equal_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data);
 void ones_like_tensor_cpu(Tensor* tensor, float* result_data);
 void zeros_like_tensor_cpu(Tensor* tensor, float* result_data);
 void transpose_1D_tensor_cpu(Tensor* tensor, float* result_data);

norch/csrc/cuda.cu

@@ -542,6 +542,28 @@ __host__ void log_tensor_cuda(Tensor* tensor, float* result_data) {
     cudaDeviceSynchronize();
 }
 
+__global__ void equal_tensor_cuda_kernel(float* data1, float* data2, float* result_data, int size) {
+    
+    int i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < size) {
+        result_data[i] = (data1[i] == data2[i]) ? 1.0f : 0.0f;
+    }
+}
+
+__host__ void equal_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data) {
+    
+    int number_of_blocks = (tensor1->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;
+    equal_tensor_cuda_kernel<<<number_of_blocks, THREADS_PER_BLOCK>>>(tensor1->data, tensor2->data, result_data, tensor1->size);
+
+    cudaError_t error = cudaGetLastError();
+    if (error != cudaSuccess) {
+        printf("CUDA error: %s\n", cudaGetErrorString(error));
+        exit(-1);
+    }
+
+    cudaDeviceSynchronize();
+}
+
 
 __global__ void ones_like_tensor_cuda_kernel(float* data, float* result_data, int size) {
     

norch/csrc/cuda.h

@@ -49,6 +49,9 @@
     __global__ void log_tensor_cuda_kernel(float* data, float* result_data, int size);
     __host__ void log_tensor_cuda(Tensor* tensor, float* result_data);
 
+    __global__ void equal_tensor_cuda_kernel(float* data1, float* data2, float* result_data, int size);
+    __host__ void equal_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data);
+
     __global__ void ones_like_tensor_cuda_kernel(float* data, float* result_data, int size);
     __host__ void ones_like_tensor_cuda(Tensor* tensor, float* result_data);
 

norch/csrc/tensor.cpp

@@ -173,7 +173,6 @@ extern "C" {
     }
 
     Tensor* sum_tensor(Tensor* tensor, int axis, bool keepdim) {
-
         char* device = (char*)malloc(strlen(tensor->device) + 1);
         if (device != NULL) {
             strcpy(device, tensor->device);
@@ -211,16 +210,145 @@ extern "C" {
 
             float* result_data;
             cudaMalloc((void**)&result_data, size * sizeof(float));
+            cudaMemset(result_data, 0, size * sizeof(float));
             sum_tensor_cuda(tensor, result_data);
             return create_tensor(result_data, shape, ndim, device);
         } 
+        else {
+            float* result_data = (float*)calloc(size, sizeof(float));
+            if (result_data == NULL) {
+                fprintf(stderr, "Memory allocation failed\n");
+                exit(1);
+            }
+            sum_tensor_cpu(tensor, result_data, size, shape, axis);
+
+            if (keepdim) {
+                shape = (int*) malloc((tensor->ndim) * sizeof(int));
+                for (int i = 0; i < tensor->ndim; i++) {
+                    shape[i] = tensor->shape[i];
+                }
+                shape[axis] = 1;
+                ndim = tensor->ndim;
+                
+            }
+
+            return create_tensor(result_data, shape, ndim, device);
+        }     
+    }
+
+    Tensor* max_tensor(Tensor* tensor, int axis, bool keepdim) {
+        char* device = (char*)malloc(strlen(tensor->device) + 1);
+        if (device != NULL) {
+            strcpy(device, tensor->device);
+        } else {
+            fprintf(stderr, "Memory allocation failed\n");
+            exit(-1);
+        }
+        int ndim;
+        int* shape;
+        if (axis == -1) {
+            shape = (int*) malloc(sizeof(int));
+            if (shape == NULL) {
+                fprintf(stderr, "Memory allocation failed\n");
+                exit(1);
+            }
+            shape[0] = 1;
+            ndim = 1;
+        } else {
+            shape = (int*) malloc((tensor->ndim - 1) * sizeof(int));
+            for (int i = 0, j = 0; i < tensor->ndim; ++i) {
+                if (i != axis) {
+                    shape[j++] = tensor->shape[i];
+                }
+            }
+            ndim = tensor->ndim - 1;
+            
+        }
+
+        int size = 1;
+        for (int i = 0; i < ndim; i++) {
+            size *= shape[i];
+        }
+  
+        if (strcmp(tensor->device, "cuda") == 0) {
+
+            float* result_data;
+            cudaMalloc((void**)&result_data, size * sizeof(float));
+            //cudaMemset(result_data, -INFINITY, size * sizeof(float));
+            //max_tensor_cuda(tensor, result_data);
+            return create_tensor(result_data, shape, ndim, device);
+        } 
         else {
             float* result_data = (float*)malloc(size * sizeof(float));
             if (result_data == NULL) {
                 fprintf(stderr, "Memory allocation failed\n");
                 exit(1);
             }
-            sum_tensor_cpu(tensor, result_data, size, shape, axis);
+            max_tensor_cpu(tensor, result_data, size, shape, axis);
+
+            if (keepdim) {
+                shape = (int*) malloc((tensor->ndim) * sizeof(int));
+                for (int i = 0; i < tensor->ndim; i++) {
+                    shape[i] = tensor->shape[i];
+                }
+                shape[axis] = 1;
+                ndim = tensor->ndim;
+                
+            }
+
+            return create_tensor(result_data, shape, ndim, device);
+        }     
+    }
+
+    Tensor* min_tensor(Tensor* tensor, int axis, bool keepdim) {
+        char* device = (char*)malloc(strlen(tensor->device) + 1);
+        if (device != NULL) {
+            strcpy(device, tensor->device);
+        } else {
+            fprintf(stderr, "Memory allocation failed\n");
+            exit(-1);
+        }
+        int ndim;
+        int* shape;
+        if (axis == -1) {
+            shape = (int*) malloc(sizeof(int));
+            if (shape == NULL) {
+                fprintf(stderr, "Memory allocation failed\n");
+                exit(1);
+            }
+            shape[0] = 1;
+            ndim = 1;
+        } else {
+            shape = (int*) malloc((tensor->ndim - 1) * sizeof(int));
+            for (int i = 0, j = 0; i < tensor->ndim; ++i) {
+                if (i != axis) {
+                    shape[j++] = tensor->shape[i];
+                }
+            }
+            ndim = tensor->ndim - 1;
+            
+        }
+
+        int size = 1;
+        for (int i = 0; i < ndim; i++) {
+            size *= shape[i];
+        }
+  
+        if (strcmp(tensor->device, "cuda") == 0) {
+
+            float* result_data;
+            cudaMalloc((void**)&result_data, size * sizeof(float));
+            //cudaMemset(result_data, INFINITY, size * sizeof(float));
+            //min_tensor_cuda(tensor, result_data);
+            return create_tensor(result_data, shape, ndim, device);
+        } 
+        else {
+            float* result_data = (float*)malloc(size * sizeof(float));
+            if (result_data == NULL) {
+                fprintf(stderr, "Memory allocation failed\n");
+                exit(1);
+            }
+            min_tensor_cpu(tensor, result_data, size, shape, axis);
 
             if (keepdim) {
                 shape = (int*) malloc((tensor->ndim) * sizeof(int));
@@ -905,6 +1033,58 @@ extern "C" {
         }
     }
 
+    Tensor* equal_tensor(Tensor* tensor1, Tensor* tensor2) {
+        if (tensor1->ndim != tensor2->ndim) {
+            fprintf(stderr, "Tensors must have the same number of dimensions %d and %d for equal\n", tensor1->ndim, tensor2->ndim);
+            exit(1);
+        }
+
+        if (strcmp(tensor1->device, tensor2->device) != 0) {
+            fprintf(stderr, "Tensors must be on the same device: %s and %s\n", tensor1->device, tensor2->device);
+            exit(1);
+        }
+
+        char* device = (char*)malloc(strlen(tensor1->device) + 1);
+        if (device != NULL) {
+            strcpy(device, tensor1->device);
+        } else {
+            fprintf(stderr, "Memory allocation failed\n");
+            exit(-1);
+        }
+        int ndim = tensor1->ndim;
+        int* shape = (int*)malloc(ndim * sizeof(int));
+        if (shape == NULL) {
+            fprintf(stderr, "Memory allocation failed\n");
+            exit(1);
+        }
+
+        for (int i = 0; i < ndim; i++) {
+            if (tensor1->shape[i] != tensor2->shape[i]) {
+                fprintf(stderr, "Tensors must have the same shape %d and %d at index %d for equal\n", tensor1->shape[i], tensor2->shape[i], i);
+                exit(1);
+            }
+            shape[i] = tensor1->shape[i];
+        }        
+        
+        if (strcmp(tensor1->device, "cuda") == 0) {
+
+            float* result_data;
+            cudaMalloc((void **)&result_data, tensor1->size * sizeof(float));
+            equal_tensor_cuda(tensor1, tensor2, result_data);
+            return create_tensor(result_data, shape, ndim, device);
+        } 
+        else {
+            float* result_data = (float*)malloc(tensor1->size * sizeof(float));
+            if (result_data == NULL) {
+                fprintf(stderr, "Memory allocation failed\n");
+                exit(1);
+            }
+            equal_tensor_cpu(tensor1, tensor2, result_data);
+            return create_tensor(result_data, shape, ndim, device);
+        }     
+    }
+
+
     Tensor* ones_like_tensor(Tensor* tensor) {
         char* device = (char*)malloc(strlen(tensor->device) + 1);
         if (device != NULL) {

norch/csrc/tensor.h

@@ -17,6 +17,8 @@ extern "C" {
     float get_item(Tensor* tensor, int* indices);
     Tensor* add_tensor(Tensor* tensor1, Tensor* tensor2);
     Tensor* sum_tensor(Tensor* tensor, int axis, bool keepdims);
+    Tensor* max_tensor(Tensor* tensor, int axis, bool keepdim);
+    Tensor* min_tensor(Tensor* tensor, int axis, bool keepdim);
     Tensor* sub_tensor(Tensor* tensor1, Tensor* tensor2);
     Tensor* elementwise_mul_tensor(Tensor* tensor1, Tensor* tensor2);
     Tensor* scalar_mul_tensor(Tensor* tensor, float scalar);
@@ -28,6 +30,7 @@ extern "C" {
     Tensor* tensor_pow_scalar(Tensor* tensor, float exponent);
     Tensor* scalar_pow_tensor(float base, Tensor* tensor);
     Tensor* log_tensor(Tensor* tensor);
+    Tensor* equal_tensor(Tensor* tensor1, Tensor* tensor2);
     void to_device(Tensor* tensor, char* device);
     Tensor* ones_like_tensor(Tensor* tensor);
     Tensor* zeros_like_tensor(Tensor* tensor);

norch/nn/__init__.py

@@ -1,3 +1,4 @@
 from .modules import *
 from .activation import *
-from .loss import *
+from .loss import *
+from .functional import *

norch/nn/activation.py

@@ -1,4 +1,5 @@
 from .module import Module
+from . import functional as F
 import math
 
 class Activation(Module):
@@ -17,4 +18,4 @@ class Sigmoid(Activation):
         super().__init__()
 
     def forward(self, x):
-        return 1.0 / (1.0 + (math.e) ** (-x)) 
+        return F.sigmoid(x)

norch/nn/functional.py

@@ -0,0 +1,5 @@
+import math
+
+def sigmoid(x):
+    return 1.0 / (1.0 + (math.e) ** (-x)) 
+

norch/tensor.py

@@ -641,6 +641,29 @@ class Tensor:
         return result_data
 
     
+    def equal(self, other):
+
+        if not isinstance(other, Tensor):
+            return False
+
+        if self.shape != other.shape:
+            return False
+        
+        Tensor._C.equal_tensor.argtypes = [ctypes.POINTER(CTensor), ctypes.POINTER(CTensor)]
+        Tensor._C.equal_tensor.restype = ctypes.POINTER(CTensor)
+
+        result_tensor_ptr = Tensor._C.equal_tensor(self.tensor, other.tensor)
+
+        result_data = Tensor()
+        result_data.tensor = result_tensor_ptr
+        result_data.shape = self.shape.copy()
+        result_data.ndim = self.ndim
+        result_data.device = self.device
+        result_data.numel = self.numel
+        
+        return result_data
+
+
     def log(self):
         Tensor._C.log_tensor.argtypes = [ctypes.POINTER(CTensor)]
         Tensor._C.log_tensor.restype = ctypes.POINTER(CTensor)
@@ -694,6 +717,76 @@ class Tensor:
             result_data.grad_fn = SumBackward(self, axis, keepdim=keepdim)
 
         return result_data
+    
+    def max(self, axis=-1, keepdim=False):
+        Tensor._C.max_tensor.argtypes = [ctypes.POINTER(CTensor), ctypes.c_int, ctypes.c_bool]
+        Tensor._C.max_tensor.restype = ctypes.POINTER(CTensor)
+
+        result_tensor_ptr = Tensor._C.max_tensor(self.tensor, axis, keepdim)
+
+        result_data = Tensor()
+        result_data.tensor = result_tensor_ptr
+
+        if axis == -1:
+            if keepdim:
+                result_data.ndim = self.ndim
+                result_data.shape = [1] * self.ndim
+
+            else:
+                result_data.shape = [1]
+                result_data.ndim = 1
+        else:
+            if keepdim:
+                result_data.shape = self.shape[:axis] + [1] + self.shape[axis+1:]
+            else:
+                result_data.shape = self.shape[:axis] + self.shape[axis+1:]
+            result_data.ndim = len(result_data.shape)
+
+        result_data.device = self.device
+        result_data.numel = 1
+        for s in result_data.shape:
+            result_data.numel *= s
+
+        result_data.requires_grad = self.requires_grad
+        if result_data.requires_grad:
+            result_data.grad_fn = MaxBackward(self, axis, keepdim=keepdim)
+
+        return result_data
+    
+    def min(self, axis=-1, keepdim=False):
+        Tensor._C.min_tensor.argtypes = [ctypes.POINTER(CTensor), ctypes.c_int, ctypes.c_bool]
+        Tensor._C.min_tensor.restype = ctypes.POINTER(CTensor)
+
+        result_tensor_ptr = Tensor._C.min_tensor(self.tensor, axis, keepdim)
+
+        result_data = Tensor()
+        result_data.tensor = result_tensor_ptr
+
+        if axis == -1:
+            if keepdim:
+                result_data.ndim = self.ndim
+                result_data.shape = [1] * self.ndim
+
+            else:
+                result_data.shape = [1]
+                result_data.ndim = 1
+        else:
+            if keepdim:
+                result_data.shape = self.shape[:axis] + [1] + self.shape[axis+1:]
+            else:
+                result_data.shape = self.shape[:axis] + self.shape[axis+1:]
+            result_data.ndim = len(result_data.shape)
+
+        result_data.device = self.device
+        result_data.numel = 1
+        for s in result_data.shape:
+            result_data.numel *= s
+
+        result_data.requires_grad = self.requires_grad
+        if result_data.requires_grad:
+            result_data.grad_fn = MinBackward(self, axis, keepdim=keepdim)
+
+        return result_data
 
     
     def sin(self):

test.py

@@ -6,6 +6,12 @@ import norch.optim as optim
 import random
 random.seed(1)
 
+a = norch.Tensor([1, 2, 3])
+b = norch.Tensor([1, 2, 4])
+print(a == b)
+
+"""
+
 to_tensor = lambda x: norch.Tensor(x)
 reshape = lambda x: x.reshape([-1, 784])
 transform = lambda x: reshape(to_tensor(x))
@@ -40,7 +46,6 @@ criterion = nn.MSELoss()
 optimizer = optim.SGD(model.parameters(), lr=0.001)
 loss_list = []
 
-
 for epoch in range(epochs):    
     for idx, batch in enumerate(train_loader):
         x, target = batch
@@ -72,4 +77,4 @@ for epoch in range(epochs):
         print('\n\n')
 
     print(f'Epoch [{epoch + 1}/{epochs}], Loss: {loss[0]:.4f}')
-    loss_list.append(loss[0])
+    loss_list.append(loss[0])"""

tests/test_autograd.py

@@ -56,7 +56,44 @@ class TestTensorAutograd(unittest.TestCase):
         
         self.assertTrue(utils.compare_torch(norch_tensor1_grad, torch_tensor1_grad))
         self.assertTrue(utils.compare_torch(norch_tensor2_grad, torch_tensor2_grad))
+    
+    def test_max_axis(self):
+        """
+        Test autograd from max specifying axis
+        """
+        norch_tensor = norch.Tensor([[[10, 10], [-4, -4]], [[5., 6], [7, 8]]], requires_grad=True).to(self.device)
+        norch_result = norch_tensor.max(axis=1).sum()
 
+        norch_result.backward()
+        norch_tensor_grad = utils.to_torch(norch_tensor.grad).to(self.device)
+
+        torch_tensor = torch.tensor([[[10, 10], [-4, -4]], [[5., 6], [7, 8]]], requires_grad=True).to(self.device)
+        
+        torch_result = torch_tensor.max(axis=1).sum()
+        torch_result.backward()
+        torch_tensor_grad = torch_tensor.grad
+        
+        self.assertTrue(utils.compare_torch(norch_tensor_grad, torch_tensor_grad))
+    
+
+    def test_min_axis(self):
+        """
+        Test autograd from min specifying axis
+        """
+        norch_tensor = norch.Tensor([[[10, 10], [-4, -4]], [[5., 6], [7, 8]]], requires_grad=True).to(self.device)
+        norch_result = norch_tensor.min(axis=1).sum()
+
+        norch_result.backward()
+        norch_tensor_grad = utils.to_torch(norch_tensor.grad).to(self.device)
+
+        torch_tensor = torch.tensor([[[10, 10], [-4, -4]], [[5., 6], [7, 8]]], requires_grad=True).to(self.device)
+        
+        torch_result = torch_tensor.min(axis=1).sum()
+        torch_result.backward()
+        torch_tensor_grad = torch_tensor.grad
+        
+        self.assertTrue(utils.compare_torch(norch_tensor_grad, torch_tensor_grad))
+    
 
     def test_broadcasted_addition_autograd(self):
         """
@@ -411,7 +448,6 @@ class TestTensorAutograd(unittest.TestCase):
 
         self.assertTrue(utils.compare_torch(torch_result_cos_tensor_grad, torch_expected_cos_tensor_grad))
     
-    
     def test_reshape(self):
         """
         Test autograd from reshaping a tensor: tensor.reshape(shape)

tests/test_operations.py

@@ -260,9 +260,90 @@ class TestTensorOperations(unittest.TestCase):
 
         torch_tensor = torch.tensor([[[1, 2], [3, -4]], [[5, 6], [7, 8]]]).to(self.device)
         torch_expected = torch.sum(torch_tensor, dim=1, keepdim=True)
+        
+        self.assertTrue(utils.compare_torch(torch_result, torch_expected))
 
-        print(torch_result, '\n', torch_expected)
+    def test_max(self):
+        """
+        Test max of a tensor: tensor.max()
+        """
+        norch_tensor = norch.Tensor([[[1, 2], [3, -4]], [[5, 6], [7, 8]]]).to(self.device)
+        norch_result = norch_tensor.max()
+        torch_result = utils.to_torch(norch_result).to(self.device)
 
+        torch_tensor = torch.tensor([[[1, 2], [3, -4]], [[5, 6], [7, 8]]]).to(self.device)
+        torch_expected = torch.max(torch_tensor)
+
+        print(torch_result, torch_expected)
+
+        self.assertTrue(utils.compare_torch(torch_result, torch_expected))
+
+    def test_max_axis(self):
+        """
+        Test max of a tensor along a specific axis without keeping the dimensions
+        """
+        norch_tensor = norch.Tensor([[[1, 2], [3, -4]], [[5, 6], [7, 8]]]).to(self.device)
+        norch_result = norch_tensor.max(axis=1)
+        torch_result = utils.to_torch(norch_result).to(self.device)
+
+        torch_tensor = torch.tensor([[[1, 2], [3, -4]], [[5, 6], [7, 8]]]).to(self.device)
+        torch_expected, _ = torch.max(torch_tensor, dim=1)
+
+        self.assertTrue(utils.compare_torch(torch_result, torch_expected))
+
+
+    def test_max_axis_keepdim(self):
+        """
+        Test max of a tensor along a specific axis with keepdim=True
+        """
+        norch_tensor = norch.Tensor([[[1, 2], [3, -4]], [[5, 6], [7, 8]]]).to(self.device)
+        norch_result = norch_tensor.max(axis=1, keepdim=True)
+        torch_result = utils.to_torch(norch_result).to(self.device)
+
+        torch_tensor = torch.tensor([[[1, 2], [3, -4]], [[5, 6], [7, 8]]]).to(self.device)
+        torch_expected, _ = torch.max(torch_tensor, dim=1, keepdim=True)
+        
+        self.assertTrue(utils.compare_torch(torch_result, torch_expected))
+
+    def test_min(self):
+        """
+        Test min of a tensor: tensor.min()
+        """
+        norch_tensor = norch.Tensor([[[1, 2], [3, -4]], [[5, 6], [7, 8]]]).to(self.device)
+        norch_result = norch_tensor.min()
+        torch_result = utils.to_torch(norch_result).to(self.device)
+
+        torch_tensor = torch.tensor([[[1, 2], [3, -4]], [[5, 6], [7, 8]]]).to(self.device)
+        torch_expected = torch.min(torch_tensor)
+
+        print(torch_result, torch_expected)
+
+        self.assertTrue(utils.compare_torch(torch_result, torch_expected))
+
+    def test_min_axis(self):
+        """
+        Test min of a tensor along a specific axis without keeping the dimensions
+        """
+        norch_tensor = norch.Tensor([[[1, 2], [3, -4]], [[5, 6], [7, 8]]]).to(self.device)
+        norch_result = norch_tensor.min(axis=1)
+        torch_result = utils.to_torch(norch_result).to(self.device)
+
+        torch_tensor = torch.tensor([[[1, 2], [3, -4]], [[5, 6], [7, 8]]]).to(self.device)
+        torch_expected, _ = torch.min(torch_tensor, dim=1)
+
+        self.assertTrue(utils.compare_torch(torch_result, torch_expected))
+
+
+    def test_min_axis_keepdim(self):
+        """
+        Test min of a tensor along a specific axis with keepdim=True
+        """
+        norch_tensor = norch.Tensor([[[1, 2], [3, -4]], [[5, 6], [7, 8]]]).to(self.device)
+        norch_result = norch_tensor.min(axis=1, keepdim=True)
+        torch_result = utils.to_torch(norch_result).to(self.device)
+
+        torch_tensor = torch.tensor([[[1, 2], [3, -4]], [[5, 6], [7, 8]]]).to(self.device)
+        torch_expected, _ = torch.min(torch_tensor, dim=1, keepdim=True)
         
         self.assertTrue(utils.compare_torch(torch_result, torch_expected))
 
@@ -450,6 +531,22 @@ class TestTensorOperations(unittest.TestCase):
 
         self.assertTrue(utils.compare_torch(torch_result, torch_expected))
 
+    def test_equal(self):
+        """
+        Test equal two tensors: tensor1 == tensor2
+        """
+        norch_tensor1 = norch.Tensor([[[1, 2], [3, -4]], [[5, 1], [7, 8]]]).to(self.device)
+        norch_tensor2 = norch.Tensor([[[1, 1], [1, 1]], [[1, 1], [1, 1]]]).to(self.device)
+        norch_result = norch_tensor1.equal(norch_tensor2)
+        torch_result = utils.to_torch(norch_result).to(self.device)
+
+        torch_tensor1 = torch.tensor([[[1, 2], [3, -4]], [[5, 1], [7, 8]]]).to(self.device)
+        torch_tensor2 = torch.tensor([[[1, 1], [1, 1]], [[1, 1], [1, 1]]]).to(self.device)
+        torch_expected = (torch_tensor1 == torch_tensor2).float()
+
+        self.assertTrue(utils.compare_torch(torch_result, torch_expected))
+
+
 
     def test_zeros_like(self):
         """