aparler/PyNorch
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

fix pow, div, log operations

Browse source
This commit is contained in:
lucasdelimanogueira 2024-05-06 10:18:12 -03:00
parent ca4a0f0dbc
commit 065c6a0f29
15 changed files with 160 additions and 51 deletions
Download patch file Download diff file Expand all files Collapse all files

BIN
build/cpu.o
View file

Binary file not shown.

BIN
build/cuda.cu.o
View file

Binary file not shown.

BIN
build/libtensor.so
View file

Binary file not shown.

BIN
build/tensor.o
View file

Binary file not shown.

BIN
norch/__pycache__/tensor.cpython-38.pyc
View file

Binary file not shown.

BIN
norch/autograd/__pycache__/functions.cpython-38.pyc
View file

Binary file not shown.

7
norch/autograd/functions.py
View file

@ -64,13 +64,12 @@ class PowBackward:
class LogBackward:
def __init__(self, input_tensor):
self.input_tensor = input_tensor
def __init__(self, x):
self.input = [x]
def backward(self, gradient):
input_tensor = self.input_tensor
grad_input = gradient / input_tensor
grad_input = gradient / self.input[0]
return [grad_input]

7
norch/csrc/cpu.cpp
View file

@ -47,6 +47,13 @@ void tensor_div_scalar_cpu(Tensor* tensor, float scalar, float* result_data) {
}
}
void tensor_div_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];
}
}
void matmul_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data) {
for (int i = 0; i < tensor1->shape[0]; i++) {

1
norch/csrc/cpu.h
View file

@ -9,6 +9,7 @@ void sub_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data);
void elementwise_mul_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data);
void scalar_div_tensor_cpu(float scalar, Tensor* tensor, float* result_data);
void tensor_div_scalar_cpu(Tensor* tensor, float scalar, float* result_data);
void tensor_div_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data);
void matmul_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data);
void broadcasted_batched_matmul_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data);
void batched_matmul_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data);

22
norch/csrc/cuda.cu
View file

@ -217,6 +217,28 @@ __host__ void tensor_div_scalar_cuda(Tensor* tensor, float scalar, float* result
cudaDeviceSynchronize();
}
__global__ void tensor_div_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];
}
}
__host__ void tensor_div_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data) {
int number_of_blocks = (tensor1->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;
tensor_div_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 matmul_tensor_cuda_kernel(float* data1, float* data2, float* result_data, int rows1, int cols1, int cols2) {
int row = blockIdx.y * blockDim.y + threadIdx.y;

3
norch/csrc/cuda.h
View file

@ -25,6 +25,9 @@
__global__ void tensor_div_scalar_cuda_kernel(float* data, float scalar, float* result_data, int size);
__host__ void tensor_div_scalar_cuda(Tensor* tensor, float scalar, float* result_data);
__global__ void tensor_div_tensor_cuda_kernel(float* data1, float* data2, float* result_data, int size);
__host__ void tensor_div_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data);
__global__ void matmul_tensor_cuda_kernel(float* data1, float* data2, float* result_data, int rows1, int cols1, int cols2);
__host__ void matmul_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data);

51
norch/csrc/tensor.cpp
View file

@ -411,6 +411,57 @@ extern "C" {
}
}
Tensor* tensor_div_tensor(Tensor* tensor1, Tensor* tensor2) {
if (tensor1->ndim != tensor2->ndim) {
fprintf(stderr, "Tensors must have the same number of dimensions %d and %d for element-wise multiplication\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 subtraction\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));
tensor_div_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);
}
tensor_div_tensor_cpu(tensor1, tensor2, result_data);
return create_tensor(result_data, shape, ndim, device);
}
}
Tensor* matmul_tensor(Tensor* tensor1, Tensor* tensor2) {
//MxN @ NxP = MxP
// Check if tensors have compatible shapes for matrix multiplication

1
norch/csrc/tensor.h
View file

@ -22,6 +22,7 @@ extern "C" {
Tensor* scalar_mul_tensor(Tensor* tensor, float scalar);
Tensor* scalar_div_tensor(float scalar, Tensor* tensor);
Tensor* tensor_div_scalar(Tensor* tensor, float scalar);
Tensor* tensor_div_tensor(Tensor* tensor1, Tensor* tensor2);
Tensor* reshape_tensor(Tensor* tensor, int* new_shape, int new_ndim);
Tensor* matmul_tensor(Tensor* tensor1, Tensor* tensor2);
Tensor* tensor_pow_scalar(Tensor* tensor, float exponent);

114
norch/tensor.py
View file

@ -153,7 +153,8 @@ class Tensor:
if self.grad_fn is not None: # not a leaf
grads = self.grad_fn.backward(gradient)
for tensor, grad in zip(self.grad_fn.input, grads):
tensor.backward(grad)
if isinstance(tensor, Tensor):
tensor.backward(grad)
def zero_grad(self):
self.grad = None
@ -200,6 +201,9 @@ class Tensor:
return self.__str__()
def __add__(self, other):
if isinstance(other, (int, float)):
other = other * self.ones_like()
if self.shape != other.shape:
raise ValueError("Tensors must have the same shape for addition")
@ -221,6 +225,9 @@ class Tensor:
return result_data
def __sub__(self, other):
if isinstance(other, (int, float)):
other = other * self.ones_like()
if self.shape != other.shape:
raise ValueError("Tensors must have the same shape for subtraction")
@ -345,51 +352,11 @@ class Tensor:
return result_data
def __pow__(self, other):
if isinstance(other, (int, float)):
other = float(other)
Tensor._C.tensor_pow_scalar.argtypes = [ctypes.POINTER(CTensor), ctypes.c_float]
Tensor._C.tensor_pow_scalar.restype = ctypes.POINTER(CTensor)
result_tensor_ptr = Tensor._C.tensor_pow_scalar(self.tensor, ctypes.c_float(other))
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.requires_grad = self.requires_grad
if result_data.requires_grad:
result_data.grad_fn = PowBackward(self, other)
elif isinstance(self, (int, float)):
self = float(self)
Tensor._C.scalar_pow_tensor.argtypes = [ctypes.c_float, ctypes.POINTER(CTensor)]
Tensor._C.scalar_pow_tensor.restype = ctypes.POINTER(CTensor)
result_tensor_ptr = Tensor._C.scalar_pow_tensor(ctypes.c_float(self), other.tensor)
result_data = Tensor()
result_data.tensor = result_tensor_ptr
result_data.shape = other.shape.copy()
result_data.ndim = other.ndim
result_data.device = other.device
result_data.requires_grad = other.requires_grad
if result_data.requires_grad:
result_data.grad_fn = PowBackward(other, self)
else:
raise ValueError("Tensor to tensor power is not supported")
return result_data
def __truediv__(self, other):
other = float(other)
Tensor._C.tensor_div_scalar.argtypes = [ctypes.POINTER(CTensor), ctypes.c_float]
Tensor._C.tensor_div_scalar.restype = ctypes.POINTER(CTensor)
Tensor._C.tensor_pow_scalar.argtypes = [ctypes.POINTER(CTensor), ctypes.c_float]
Tensor._C.tensor_pow_scalar.restype = ctypes.POINTER(CTensor)
result_tensor_ptr = Tensor._C.tensor_div_scalar(other.tensor, ctypes.c_float(self))
result_tensor_ptr = Tensor._C.tensor_pow_scalar(self.tensor, ctypes.c_float(other))
result_data = Tensor()
result_data.tensor = result_tensor_ptr
@ -399,7 +366,62 @@ class Tensor:
result_data.requires_grad = self.requires_grad
if result_data.requires_grad:
result_data.grad_fn = DivisionBackward(self, other)
result_data.grad_fn = PowBackward(self, other)
return result_data
def __rpow__(self, other):
other = float(other)
Tensor._C.scalar_pow_tensor.argtypes = [ctypes.c_float, ctypes.POINTER(CTensor)]
Tensor._C.scalar_pow_tensor.restype = ctypes.POINTER(CTensor)
result_tensor_ptr = Tensor._C.scalar_pow_tensor(ctypes.c_float(other), 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.requires_grad = self.requires_grad
if result_data.requires_grad:
result_data.grad_fn = PowBackward(other, self)
return result_data
def __truediv__(self, other):
if isinstance(other, (int, float)):
other = float(other)
Tensor._C.tensor_div_scalar.argtypes = [ctypes.POINTER(CTensor), ctypes.c_float]
Tensor._C.tensor_div_scalar.restype = ctypes.POINTER(CTensor)
result_tensor_ptr = Tensor._C.tensor_div_scalar(self.tensor, ctypes.c_float(other))
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.requires_grad = self.requires_grad
if result_data.requires_grad:
result_data.grad_fn = DivisionBackward(self, other)
elif isinstance(self, Tensor) and isinstance(other, Tensor):
Tensor._C.tensor_div_tensor.argtypes = [ctypes.POINTER(CTensor), ctypes.POINTER(CTensor)]
Tensor._C.tensor_div_tensor.restype = ctypes.POINTER(CTensor)
result_tensor_ptr = Tensor._C.tensor_div_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.requires_grad = self.requires_grad or other.requires_grad
if result_data.requires_grad:
result_data.grad_fn = DivisionBackward(self, other)
return result_data
@ -418,7 +440,7 @@ class Tensor:
result_data.shape = self.shape.copy()
result_data.ndim = self.ndim
result_data.device = self.device
result_data.requires_grad = self.requires_grad
if result_data.requires_grad:
result_data.grad_fn = DivisionBackward(other, self)

5
test.py
View file

@ -75,7 +75,10 @@ if __name__ == "__main__":
[[19, 20], [21, 22], [23, 24]],
[[25, 26], [27, 28], [29, 30]]], requires_grad=True)
print(tensor1 / 1)
result = 2 ** tensor1
result = result.sum()
result.backward()
print(tensor1.grad)
exit()
# Reshape tensor1 to 2x3x5