diff --git a/build/cpu.o b/build/cpu.o index 9b48aff..4f63638 100644 Binary files a/build/cpu.o and b/build/cpu.o differ diff --git a/build/cuda.cu.o b/build/cuda.cu.o index 3ddee86..43c68ba 100644 Binary files a/build/cuda.cu.o and b/build/cuda.cu.o differ diff --git a/build/libtensor.so b/build/libtensor.so index 856e985..5dfcec6 100755 Binary files a/build/libtensor.so and b/build/libtensor.so differ diff --git a/build/tensor.o b/build/tensor.o index 1e68803..ce0b93c 100644 Binary files a/build/tensor.o and b/build/tensor.o differ diff --git a/norch/__pycache__/tensor.cpython-38.pyc b/norch/__pycache__/tensor.cpython-38.pyc index b9d5e9b..21215c7 100644 Binary files a/norch/__pycache__/tensor.cpython-38.pyc and b/norch/__pycache__/tensor.cpython-38.pyc differ diff --git a/norch/autograd/__pycache__/functions.cpython-38.pyc b/norch/autograd/__pycache__/functions.cpython-38.pyc index 503e242..73c0c12 100644 Binary files a/norch/autograd/__pycache__/functions.cpython-38.pyc and b/norch/autograd/__pycache__/functions.cpython-38.pyc differ diff --git a/norch/autograd/functions.py b/norch/autograd/functions.py index 89c1a8d..dec3f54 100644 --- a/norch/autograd/functions.py +++ b/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()] + diff --git a/norch/csrc/cpu.cpp b/norch/csrc/cpu.cpp index ff5995d..6cb8071 100644 --- a/norch/csrc/cpu.cpp +++ b/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]); + } +} diff --git a/norch/csrc/cpu.h b/norch/csrc/cpu.h index 6e0e2a3..9e5dc62 100644 --- a/norch/csrc/cpu.h +++ b/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 */ diff --git a/norch/csrc/cuda.cu b/norch/csrc/cuda.cu index 73efdd2..beae140 100644 --- a/norch/csrc/cuda.cu +++ b/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<<>>(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<<>>(tensor->data, result_data, tensor->size); + + cudaError_t error = cudaGetLastError(); + if (error != cudaSuccess) { + printf("CUDA error: %s\n", cudaGetErrorString(error)); + exit(-1); + } + + cudaDeviceSynchronize(); +} + + + diff --git a/norch/csrc/cuda.h b/norch/csrc/cuda.h index d21beef..c92cfb1 100644 --- a/norch/csrc/cuda.h +++ b/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); diff --git a/norch/csrc/tensor.cpp b/norch/csrc/tensor.cpp index 1c46245..837e720 100644 --- a/norch/csrc/tensor.cpp +++ b/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) { diff --git a/norch/csrc/tensor.h b/norch/csrc/tensor.h index 90b4d0d..d45901f 100644 --- a/norch/csrc/tensor.h +++ b/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); diff --git a/norch/tensor.py b/norch/tensor.py index bd5bf18..3d309f9 100644 --- a/norch/tensor.py +++ b/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) diff --git a/tests/__init__.py b/tests/__init__.py index fef67f5..4750076 100644 --- a/tests/__init__.py +++ b/tests/__init__.py @@ -1 +1,3 @@ -from .test_operations import * \ No newline at end of file +from .test_operations import * +from .test_autograd import * +from .test_nn import * \ No newline at end of file diff --git a/tests/test_autograd.py b/tests/test_autograd.py index 2ac0d92..d70cddd 100644 --- a/tests/test_autograd.py +++ b/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) diff --git a/tests/test_nn.py b/tests/test_nn.py index 483eac3..1bf12d3 100644 --- a/tests/test_nn.py +++ b/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)) - - - \ No newline at end of file + self.assertTrue(utils.compare_torch(sigmoid_torch_result, sigmoid_torch_expected)) \ No newline at end of file diff --git a/tests/test_operations.py b/tests/test_operations.py index 6ded830..ed72ec6 100644 --- a/tests/test_operations.py +++ b/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__':