Merge pull request #42 from lucasdelimanogueira/tmp

Cosine and sine operations with autograd

norch/autograd/functions.py

@@ -114,4 +114,21 @@ class DivisionBackward:
         grad_x = gradient / y
         grad_y = -1 * gradient * (x / (y * y))
         return [grad_x, grad_y]
+    
+class SinBackward:
+    def __init__(self, x):
+        self.input = [x]
+
+    def backward(self, gradient):
+        x = self.input[0]
+        return [gradient * x.cos()]
+    
+class CosBackward:
+    def __init__(self, x):
+        self.input = [x]
+
+    def backward(self, gradient):
+        x = self.input[0]
+        return [-gradient * x.sin()]
+
 

norch/csrc/cpu.cpp

@@ -189,3 +189,15 @@ void assign_tensor_cpu(Tensor* tensor, float* result_data) {
         result_data[i] = tensor->data[i];
     }
 }
+
+void sin_tensor_cpu(Tensor* tensor, float* result_data) {
+    for (int i = 0; i < tensor->size; i++) {
+        result_data[i] = sinf(tensor->data[i]);
+    }
+}
+
+void cos_tensor_cpu(Tensor* tensor, float* result_data) {
+    for (int i = 0; i < tensor->size; i++) {
+        result_data[i] = cosf(tensor->data[i]);
+    }
+}

norch/csrc/cpu.h

@@ -24,5 +24,7 @@ void transpose_2D_tensor_cpu(Tensor* tensor, float* result_data);
 void transpose_3D_tensor_cpu(Tensor* tensor, float* result_data);
 void transpose_axes_cpu(Tensor* tensor, float* result_data, int axis1, int axis2, int* new_shape);
 void assign_tensor_cpu(Tensor* tensor, float* result_data);
+void sin_tensor_cpu(Tensor* tensor, float* result_data);
+void cos_tensor_cpu(Tensor* tensor, float* result_data);
 
 #endif /* CPU_H */

norch/csrc/cuda.cu

@@ -482,4 +482,50 @@ __host__ void assign_tensor_cuda(Tensor* tensor, float* result_data) {
     cudaDeviceSynchronize();
 }
 
+__global__ void sin_tensor_cuda_kernel(float* data, float* result_data, int size) {
+    
+    int i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < size) {
+        result_data[i] = sinf(data[i]);
+    }
+}
+
+__host__ void sin_tensor_cuda(Tensor* tensor, float* result_data) {
+    
+    int number_of_blocks = (tensor->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;
+    sin_tensor_cuda_kernel<<<number_of_blocks, THREADS_PER_BLOCK>>>(tensor->data, result_data, tensor->size);
+
+    cudaError_t error = cudaGetLastError();
+    if (error != cudaSuccess) {
+        printf("CUDA error: %s\n", cudaGetErrorString(error));
+        exit(-1);
+    }
+
+    cudaDeviceSynchronize();
+}
+
+__global__ void cos_tensor_cuda_kernel(float* data, float* result_data, int size) {
+    
+    int i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < size) {
+        result_data[i] = cosf(data[i]);
+    }
+}
+
+__host__ void cos_tensor_cuda(Tensor* tensor, float* result_data) {
+    
+    int number_of_blocks = (tensor->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;
+    cos_tensor_cuda_kernel<<<number_of_blocks, THREADS_PER_BLOCK>>>(tensor->data, result_data, tensor->size);
+
+    cudaError_t error = cudaGetLastError();
+    if (error != cudaSuccess) {
+        printf("CUDA error: %s\n", cudaGetErrorString(error));
+        exit(-1);
+    }
+
+    cudaDeviceSynchronize();
+}
+
+
+
 

norch/csrc/cuda.h

@@ -52,6 +52,11 @@
     __global__ void assign_tensor_cuda_kernel(float* data, float* result_data, int size);
     __host__ void assign_tensor_cuda(Tensor* tensor, float* result_data);
 
+    __global__ void sin_tensor_cuda_kernel(float* data, float* result_data, int size);
+    __host__ void sin_tensor_cuda(Tensor* tensor, float* result_data);
+
+     __global__ void cos_tensor_cuda_kernel(float* data, float* result_data, int size);
+    __host__ void cos_tensor_cuda(Tensor* tensor, float* result_data);
 
 
 

norch/csrc/tensor.cpp

@@ -105,6 +105,7 @@ extern "C" {
             }
             shape[i] = tensor1->shape[i];
         }        
+        
         if (strcmp(tensor1->device, "cuda") == 0) {
 
             float* result_data;
@@ -854,6 +855,80 @@ extern "C" {
         }
     }
 
+    Tensor* sin_tensor(Tensor* tensor) {
+        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 = tensor->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++) {
+            shape[i] = tensor->shape[i];
+        }
+
+        if (strcmp(tensor->device, "cuda") == 0) {
+
+            float* result_data;
+            cudaMalloc((void **)&result_data, tensor->size * sizeof(float));
+            sin_tensor_cuda(tensor, result_data);
+            return create_tensor(result_data, shape, ndim, device);
+        } 
+        else {
+            float* result_data = (float*)malloc(tensor->size * sizeof(float));
+            if (result_data == NULL) {
+                fprintf(stderr, "Memory allocation failed\n");
+                exit(1);
+            }
+            sin_tensor_cpu(tensor, result_data);
+            return create_tensor(result_data, shape, ndim, device);
+        }
+    }
+
+    Tensor* cos_tensor(Tensor* tensor) {
+        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 = tensor->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++) {
+            shape[i] = tensor->shape[i];
+        }
+
+        if (strcmp(tensor->device, "cuda") == 0) {
+
+            float* result_data;
+            cudaMalloc((void **)&result_data, tensor->size * sizeof(float));
+            cos_tensor_cuda(tensor, result_data);
+            return create_tensor(result_data, shape, ndim, device);
+        } 
+        else {
+            float* result_data = (float*)malloc(tensor->size * sizeof(float));
+            if (result_data == NULL) {
+                fprintf(stderr, "Memory allocation failed\n");
+                exit(1);
+            }
+            cos_tensor_cpu(tensor, result_data);
+            return create_tensor(result_data, shape, ndim, device);
+        }
+    }
+
     Tensor* transpose_tensor(Tensor* tensor) {
         char* device = (char*)malloc(strlen(tensor->device) + 1);
         if (device != NULL) {

norch/csrc/tensor.h

@@ -31,6 +31,8 @@ extern "C" {
     void to_device(Tensor* tensor, char* device);
     Tensor* ones_like_tensor(Tensor* tensor);
     Tensor* zeros_like_tensor(Tensor* tensor);
+    Tensor* sin_tensor(Tensor* tensor);
+    Tensor* cos_tensor(Tensor* tensor);
     Tensor* transpose_tensor(Tensor* tensor);
     Tensor* transpose_axes_tensor(Tensor* tensor, int axis1, int axis2);
     void make_contiguous(Tensor* tensor);

norch/tensor.py

@@ -402,6 +402,7 @@ class Tensor:
             result_data.shape = [other.shape[0], self.shape[1], other.shape[2]]
             result_data.ndim = 3
             result_data.device = self.device
+            result_data.numel = 1
             for s in result_data.shape:
                 result_data.numel *= s
         else:
@@ -422,6 +423,7 @@ class Tensor:
             result_data.shape = [self.shape[0], other.shape[1]]
             result_data.ndim = 2
             result_data.device = self.device
+            result_data.numel = 1
             for s in result_data.shape:
                 result_data.numel *= s
 
@@ -571,6 +573,46 @@ class Tensor:
 
         return result_data
     
+    def sin(self):
+        Tensor._C.sin_tensor.argtypes = [ctypes.POINTER(CTensor)]
+        Tensor._C.sin_tensor.restype = ctypes.POINTER(CTensor)
+
+        result_tensor_ptr = Tensor._C.sin_tensor(self.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
+
+        result_data.requires_grad = self.requires_grad
+        if result_data.requires_grad:
+            result_data.grad_fn = SinBackward(self)
+        
+        return result_data
+    
+    def cos(self):
+        Tensor._C.cos_tensor.argtypes = [ctypes.POINTER(CTensor)]
+        Tensor._C.cos_tensor.restype = ctypes.POINTER(CTensor)
+
+        result_tensor_ptr = Tensor._C.cos_tensor(self.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
+
+        result_data.requires_grad = self.requires_grad
+        if result_data.requires_grad:
+            result_data.grad_fn = CosBackward(self)
+        
+        return result_data
+    
+
+    
     def transpose(self, axis1, axis2):
         if axis1 < 0:
             axis1 = self.ndim + axis1
@@ -596,7 +638,6 @@ class Tensor:
             result_data.grad_fn = TransposeBackward(self, axis1, axis2)
 
         return result_data
-
     
     @property
     def T(self):
@@ -611,7 +652,7 @@ class Tensor:
         result_data.ndim = self.ndim
         result_data.device = self.device
         result_data.numel = self.numel
-        
+
         result_data.requires_grad = self.requires_grad
         if result_data.requires_grad:
             result_data.grad_fn = TBackward(self)

tests/__init__.py

@@ -1 +1,3 @@
-from .test_operations import *
+from .test_operations import *
+from .test_autograd import *
+from .test_nn import *

tests/test_autograd.py

@@ -128,7 +128,6 @@ class TestTensorAutograd(unittest.TestCase):
 
         self.assertTrue(utils.compare_torch(norch_tensor_grad_power_st, torch_tensor_grad_power_st))
     
-    
     def test_power_tensor_scalar(self):
         """
         Test autograd from tensor raised to scalar: tensor ** scalar
@@ -207,6 +206,39 @@ class TestTensorAutograd(unittest.TestCase):
         self.assertTrue(utils.compare_torch(norch_tensor1_grad_elemwise_mul, torch_tensor1_grad_elemwise_mul))
         self.assertTrue(utils.compare_torch(norch_tensor2_grad_elemwise_mul, torch_tensor2_grad_elemwise_mul))
 
+    def test_sin_tensor(self):
+        """
+        Test autograd from sin operation: sin(tensor)
+        """
+        norch_sin_tensor = norch.Tensor([[[2, 3.21], [4, 2.1]], [[6, 7], [8, 9]]], requires_grad=True).to(self.device)
+        norch_result_sin_tensor = (norch_sin_tensor.sin()).sum()
+        norch_result_sin_tensor.backward()
+        torch_result_sin_tensor_grad = utils.to_torch(norch_sin_tensor.grad)
+
+        torch_sin_tensor = torch.tensor([[[2, 3.21], [4, 2.1]], [[6, 7], [8, 9]]], requires_grad=True).to(self.device)
+        torch_expected_sin_tensor = (torch.sin(torch_sin_tensor)).sum()
+        torch_expected_sin_tensor.backward()
+        torch_expected_sin_tensor_grad = torch_sin_tensor.grad
+
+        self.assertTrue(utils.compare_torch(torch_result_sin_tensor_grad, torch_expected_sin_tensor_grad))
+    
+    def test_cos_tensor(self):
+        """
+        Test autograd from cosine operation: cos(tensor)
+        """
+        norch_cos_tensor = norch.Tensor([[[2, 3.21], [4, 2.1]], [[6, 7], [8, 9]]], requires_grad=True).to(self.device)
+        norch_result_cos_tensor = (norch_cos_tensor.sin()).sum()
+        norch_result_cos_tensor.backward()
+        torch_result_cos_tensor_grad = utils.to_torch(norch_cos_tensor.grad)
+
+        torch_cos_tensor = torch.tensor([[[2, 3.21], [4, 2.1]], [[6, 7], [8, 9]]], requires_grad=True).to(self.device)
+        torch_expected_cos_tensor = (torch.sin(torch_cos_tensor)).sum()
+        torch_expected_cos_tensor.backward()
+        torch_expected_cos_tensor_grad = torch_cos_tensor.grad
+
+        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_nn.py

@@ -4,7 +4,7 @@ from norch import utils
 import torch
 import os
 
-class TestNNModule(unittest.TestCase):
+class TestNNModuleLoss(unittest.TestCase):
 
     def setUp(self):
         self.device = os.environ.get('device')
@@ -40,9 +40,16 @@ class TestNNModule(unittest.TestCase):
         labels_torch = torch.tensor([4, 3, 2.1, 1])
         loss_torch_expected = loss_fn_torch(predictions_torch, labels_torch)        
         
-        print(loss_torch_result, loss_torch_expected)
         self.assertTrue(utils.compare_torch(loss_torch_result, loss_torch_expected))
 
+
+class TestNNModuleActivationFn(unittest.TestCase):
+
+    def setUp(self):
+        self.device = os.environ.get('device')
+        if self.device is None or self.device != 'cuda':
+            self.device = 'cpu'
+
     def test_sigmoid_activation(self):
         """
         Test Sigmoid activation function
@@ -78,7 +85,4 @@ class TestNNModule(unittest.TestCase):
         x = torch.tensor([0, 0, 0])
         sigmoid_torch_expected = sigmoid_fn_torch.forward(x)
 
-        self.assertTrue(utils.compare_torch(sigmoid_torch_result, sigmoid_torch_expected))
-
-
-        
+        self.assertTrue(utils.compare_torch(sigmoid_torch_result, sigmoid_torch_expected))

tests/test_operations.py

@@ -267,6 +267,58 @@ class TestTensorOperations(unittest.TestCase):
 
         self.assertTrue(utils.compare_torch(torch_result, torch_expected))
 
+    def test_tensor_sin(self):
+        """
+        Test sine function on tensor
+        """
+        norch_tensor = norch.Tensor([[[0, 30], [45, 60]], [[90, 120], [135, 180]]]).to(self.device)
+        norch_result = norch_tensor.sin()
+        torch_result = utils.to_torch(norch_result)
+
+        torch_tensor = torch.tensor([[[0, 30], [45, 60]], [[90, 120], [135, 180]]]).to(self.device)
+        torch_expected = torch.sin(torch_tensor)
+
+        self.assertTrue(utils.compare_torch(torch_result, torch_expected))
+
+    def test_tensor_cos(self):
+        """
+        Test cosine function on tensor
+        """
+        norch_tensor = norch.Tensor([[[0, 30], [45, 60]], [[90, 120], [135, 180]]]).to(self.device)
+        norch_result = norch_tensor.cos()
+        torch_result = utils.to_torch(norch_result)
+
+        torch_tensor = torch.tensor([[[0, 30], [45, 60]], [[90, 120], [135, 180]]]).to(self.device)
+        torch_expected = torch.cos(torch_tensor)
+
+        self.assertTrue(utils.compare_torch(torch_result, torch_expected))
+
+
+    def test_zeros_like(self):
+        """
+        Test creating a tensor of zeros with the same shape as another tensor.
+        """
+        norch_tensor = norch.Tensor([[[1, 2], [3, -4]], [[5, 6], [7, 8]]]).to(self.device)
+        norch_zeros = norch_tensor.zeros_like()
+        torch_zeros_result = utils.to_torch(norch_zeros)
+
+        torch_tensor_expected = torch.tensor([[[1, 2], [3, -4]], [[5, 6], [7, 8]]]).to(self.device)
+        torch_zeros_expected = torch.zeros_like(torch_tensor_expected)
+
+        self.assertTrue(utils.compare_torch(torch_zeros_result, torch_zeros_expected))
+
+    def test_ones_like(self):
+        """
+        Test creating a tensor of ones with the same shape as another tensor.
+        """
+        norch_tensor = norch.Tensor([[[1, 2], [3, -4]], [[5, 6], [7, 8]]]).to(self.device)
+        norch_ones = norch_tensor.ones_like()
+        torch_ones_result = utils.to_torch(norch_ones)
+
+        torch_tensor_expected = torch.tensor([[[1, 2], [3, -4]], [[5, 6], [7, 8]]]).to(self.device)
+        torch_ones_expected = torch.ones_like(torch_tensor_expected)
+
+        self.assertTrue(utils.compare_torch(torch_ones_result, torch_ones_expected))
 
 
 if __name__ == '__main__':