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

sub, elementmul, pow cuda operations

Browse source
This commit is contained in:
lucasdelimanogueira 2024-04-30 00:31:46 -03:00
parent 493f6b0b83
commit 8366245813
16 changed files with 228 additions and 266 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.

22
csrc/cpu.cpp
View file

@ -2,6 +2,7 @@
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <math.h>
void add_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data) {
@ -9,3 +10,24 @@ void add_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data) {
result_data[i] = tensor1->data[i] + tensor2->data[i];
}
}
void sub_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 elementwise_mul_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 pow_tensor_cpu(Tensor* tensor, float power) {
for (int i = 0; i < tensor->size; i++) {
tensor->data[i] = powf(tensor->data[i], power);
}
}

6
csrc/cpu.h
View file

@ -1 +1,5 @@
void add_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data);
void add_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data);
void sub_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data);
void elementwise_mul_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data);
void pow_tensor_cpu(Tensor* tensor, float power);

64
csrc/cuda.cu
View file

@ -1,6 +1,7 @@
#include "tensor.h"
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#define THREADS_PER_BLOCK 128
@ -54,3 +55,66 @@ __host__ void add_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_da
cudaDeviceSynchronize();
}
__global__ void sub_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 sub_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data) {
int number_of_blocks = (tensor1->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;
sub_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 elementwise_mul_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 elementwise_mul_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data) {
int number_of_blocks = (tensor1->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;
elementwise_mul_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 pow_tensor_cuda_kernel(float* data, float power, int size) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < size) {
data[i] = powf(data[i], power);
}
}
__host__ void pow_tensor_cuda(Tensor* tensor, float power) {
int number_of_blocks = (tensor->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;
pow_tensor_cuda_kernel<<<number_of_blocks, THREADS_PER_BLOCK>>>(tensor->data, power, tensor->size);
cudaError_t error = cudaGetLastError();
if (error != cudaSuccess) {
printf("CUDA error: %s\n", cudaGetErrorString(error));
exit(-1);
}
cudaDeviceSynchronize();
}

9
csrc/cuda.h
View file

@ -3,8 +3,17 @@
__host__ void cpu_to_cuda(Tensor* tensor);
__host__ void cuda_to_cpu(Tensor* tensor);
__global__ void add_tensor_cuda_kernel(float* data1, float* data2, float* result_data, int size);
__host__ void add_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data);
__global__ void sub_tensor_cuda_kernel(float* data1, float* data2, float* result_data, int size);
__host__ void sub_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data);
__global__ void elementwise_mul_tensor_cuda_kernel(float* data1, float* data2, float* result_data, int size);
__host__ void elementwise_mul_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data);
__global__ void pow_tensor_cuda_kernel(float* data, float power, int size);
__host__ void pow_tensor_cuda(Tensor* tensor, float power);
#endif /* CUDA_KERNEL_H_ */

232
csrc/tensor.cpp
View file

@ -20,7 +20,7 @@ extern "C" {
tensor->data = data;
tensor->shape = shape;
tensor->ndim = ndim;
tensor->device = (char*)malloc(strlen(device) + 1);
if (device != NULL) {
strcpy(tensor->device, device);
@ -101,7 +101,6 @@ extern "C" {
}
Tensor* add_tensor(Tensor* tensor1, Tensor* tensor2) {
printf("Adding tensor\n");
if (tensor1->ndim != tensor2->ndim) {
fprintf(stderr, "Tensors must have the same number of dimensions %d and %d for addition\n", tensor1->ndim, tensor2->ndim);
exit(1);
@ -126,44 +125,6 @@ extern "C" {
exit(1);
}
printf("Size: %d\n", tensor1->size);
/*printf("Data: [");
for (int i = 0; i < tensor1->size; i++) {
printf("%.2f", tensor1->data[i]);
if (i < tensor1->size - 1) {
printf(", ");
}
}
printf("]\n");*/
printf("Size: %d\n", tensor2->size);
/*printf("Data: [");
for (int i = 0; i < tensor2->size; i++) {
printf("%.2f", tensor2->data[i]);
if (i < tensor2->size - 1) {
printf(", ");
}
}
printf("]\n");*/
printf("Shapes : [");
for (int i = 0; i < tensor1->ndim; i++) {
printf("%d", tensor1->shape[i]);
if (i < tensor1->ndim - 1) {
printf(", ");
}
}
printf("]\n");
printf("Shapes : [");
for (int i = 0; i < tensor2->ndim; i++) {
printf("%d", tensor2->shape[i]);
if (i < tensor2->ndim - 1) {
printf(", ");
}
}
printf("]\n");
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 addition\n", tensor1->shape[i], tensor2->shape[i], i);
@ -190,7 +151,6 @@ extern "C" {
}
Tensor* sub_tensor(Tensor* tensor1, Tensor* tensor2) {
printf("Subtracting tensor\n");
if (tensor1->ndim != tensor2->ndim) {
fprintf(stderr, "Tensors must have the same number of dimensions %d and %d for subtraction\n", tensor1->ndim, tensor2->ndim);
exit(1);
@ -201,8 +161,13 @@ extern "C" {
exit(1);
}
char* device = tensor1->device;
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) {
@ -210,44 +175,6 @@ extern "C" {
exit(1);
}
printf("Size: %d\n", tensor1->size);
printf("Data: [");
for (int i = 0; i < tensor1->size; i++) {
printf("%.2f", tensor1->data[i]);
if (i < tensor1->size - 1) {
printf(", ");
}
}
printf("]\n");
printf("Size: %d\n", tensor2->size);
printf("Data: [");
for (int i = 0; i < tensor2->size; i++) {
printf("%.2f", tensor2->data[i]);
if (i < tensor2->size - 1) {
printf(", ");
}
}
printf("]\n");
printf("Shapes : [");
for (int i = 0; i < tensor1->ndim; i++) {
printf("%d", tensor1->shape[i]);
if (i < tensor1->ndim - 1) {
printf(", ");
}
}
printf("]\n");
printf("Shapes : [");
for (int i = 0; i < tensor2->ndim; i++) {
printf("%d", tensor2->shape[i]);
if (i < tensor2->ndim - 1) {
printf(", ");
}
}
printf("]\n");
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);
@ -256,21 +183,25 @@ extern "C" {
shape[i] = tensor1->shape[i];
}
float* result_data = (float*)malloc(tensor1->size * sizeof(float));
if (result_data == NULL) {
fprintf(stderr, "Memory allocation failed\n");
exit(1);
}
if (strcmp(tensor1->device, "cuda") == 0) {
for (int i = 0; i < tensor1->size; i++) {
result_data[i] = tensor1->data[i] - tensor2->data[i];
float* result_data;
cudaMalloc((void **)&result_data, tensor1->size * sizeof(float));
sub_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);
}
sub_tensor_cpu(tensor1, tensor2, result_data);
return create_tensor(result_data, shape, ndim, device);
}
return create_tensor(result_data, shape, ndim, device);
}
Tensor* elementwise_mul_tensor(Tensor* tensor1, Tensor* tensor2) {
printf("Elementwise multiplying tensor\n");
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);
@ -281,8 +212,13 @@ extern "C" {
exit(1);
}
char* device = tensor1->device;
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) {
@ -290,44 +226,6 @@ extern "C" {
exit(1);
}
printf("Size: %d\n", tensor1->size);
printf("Data: [");
for (int i = 0; i < tensor1->size; i++) {
printf("%.2f", tensor1->data[i]);
if (i < tensor1->size - 1) {
printf(", ");
}
}
printf("]\n");
printf("Size: %d\n", tensor2->size);
printf("Data: [");
for (int i = 0; i < tensor2->size; i++) {
printf("%.2f", tensor2->data[i]);
if (i < tensor2->size - 1) {
printf(", ");
}
}
printf("]\n");
printf("Shapes : [");
for (int i = 0; i < tensor1->ndim; i++) {
printf("%d", tensor1->shape[i]);
if (i < tensor1->ndim - 1) {
printf(", ");
}
}
printf("]\n");
printf("Shapes : [");
for (int i = 0; i < tensor2->ndim; i++) {
printf("%d", tensor2->shape[i]);
if (i < tensor2->ndim - 1) {
printf(", ");
}
}
printf("]\n");
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 element-wise multiplication\n", tensor1->shape[i], tensor2->shape[i], i);
@ -336,17 +234,22 @@ extern "C" {
shape[i] = tensor1->shape[i];
}
float* result_data = (float*)malloc(tensor1->size * sizeof(float));
if (result_data == NULL) {
fprintf(stderr, "Memory allocation failed\n");
exit(1);
}
if (strcmp(tensor1->device, "cuda") == 0) {
for (int i = 0; i < tensor1->size; i++) {
result_data[i] = tensor1->data[i] * tensor2->data[i];
float* result_data;
cudaMalloc((void **)&result_data, tensor1->size * sizeof(float));
elementwise_mul_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);
}
elementwise_mul_tensor_cpu(tensor1, tensor2, result_data);
return create_tensor(result_data, shape, ndim, device);
}
return create_tensor(result_data, shape, ndim, device);
}
Tensor* matmul_tensor(Tensor* tensor1, Tensor* tensor2) {
@ -401,50 +304,13 @@ extern "C" {
return create_tensor(result_data, shape, ndim, device);
}
Tensor* pow_tensor(Tensor* tensor, float power) {
printf("Powering tensor\n");
char* device = tensor->device;
int ndim = tensor->ndim;
int* shape = (int*)malloc(ndim * sizeof(int));
if (shape == NULL) {
fprintf(stderr, "Memory allocation failed\n");
exit(1);
void pow_tensor(Tensor* tensor, float power) {
if (strcmp(tensor->device, "cuda") == 0) {
pow_tensor_cuda(tensor, power);
}
else {
pow_tensor_cpu(tensor, power);
}
printf("Size: %d\n", tensor->size);
printf("Data: [");
for (int i = 0; i < tensor->size; i++) {
printf("%.2f", tensor->data[i]);
if (i < tensor->size - 1) {
printf(", ");
}
}
printf("]\n");
printf("Shapes : [");
for (int i = 0; i < tensor->ndim; i++) {
printf("%d", tensor->shape[i]);
if (i < tensor->ndim - 1) {
printf(", ");
}
}
printf("]\n");
for (int i = 0; i < ndim; i++) {
shape[i] = tensor->shape[i];
}
float* result_data = (float*)malloc(tensor->size * sizeof(float));
if (result_data == NULL) {
fprintf(stderr, "Memory allocation failed\n");
exit(1);
}
for (int i = 0; i < tensor->size; i++) {
result_data[i] = powf(tensor->data[i], power);
}
return create_tensor(result_data, shape, ndim, device);
}
void reshape_tensor(Tensor* tensor, int* new_shape, int new_ndim) {

2
csrc/tensor.h
View file

@ -20,7 +20,7 @@ extern "C" {
Tensor* elementwise_mul_tensor(Tensor* tensor1, Tensor* tensor2);
void reshape_tensor(Tensor* tensor, int* new_shape, int new_ndim);
Tensor* matmul_tensor(Tensor* tensor1, Tensor* tensor2);
Tensor* pow_tensor(Tensor* tensor, float power);
void pow_tensor(Tensor* tensor, float power);
void to_device(Tensor* tensor, char* device);
}

1
norch/__init__.py Normal file
View file

@ -0,0 +1 @@
from norch.tensor import Tensor

BIN
norch/__pycache__/__init__.cpython-38.pyc Normal file
View file

Binary file not shown.

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

Binary file not shown.

85
norch/tensor.py
View file

@ -1,4 +1,5 @@
import ctypes
import os
class CTensor(ctypes.Structure):
_fields_ = [
@ -11,7 +12,8 @@ class CTensor(ctypes.Structure):
]
class Tensor:
_C = ctypes.CDLL("../build/libtensor.so")
os.path.abspath(os.curdir)
_C = ctypes.CDLL(os.path.join(os.path.abspath(os.curdir), "build/libtensor.so"))
def __init__(self, data=None, device="cpu"):
@ -202,83 +204,8 @@ class Tensor:
power = ctypes.c_float(power)
Tensor._C.pow_tensor.argtypes = [ctypes.POINTER(CTensor), ctypes.c_float]
Tensor._C.pow_tensor.restype = ctypes.POINTER(CTensor)
Tensor._C.pow_tensor.restype = None
result_tensor_ptr = Tensor._C.pow_tensor(self.tensor, power)
Tensor._C.pow_tensor(self.tensor, power)
result_data = Tensor()
result_data.tensor = result_tensor_ptr
result_data.shape = self.shape.copy()
result_data.device = self.device
result_data.ndim = 2
return result_data
def matrix_sum(matrix1, matrix2):
# Check if the matrices can be multiplied
if len(matrix1[0]) != len(matrix2):
raise ValueError("Matrices cannot be multiplied. Inner dimensions must match.")
# Initialize the result matrix with zeros
result = [[0 for _ in range(len(matrix2[0]))] for _ in range(len(matrix1))]
# Perform matrix multiplication
for i in range(len(matrix1)):
for j in range(len(matrix2[0])):
result[i][j] += matrix1[i][i] + matrix2[i][j]
return result
if __name__ == "__main__":
from tensor import Tensor
import time
import random
import numpy as np
a = Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]]).to("cuda")
b = Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]]).to("cuda")
c = Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]]).to("cuda")
d = b+c
print(d)
#print(a ** 2)
"""#print(a)
N = 1000
a = Tensor([[random.uniform(0, 1) for _ in range(N)] for _ in range(N)])
b = Tensor([[random.uniform(0, 1) for _ in range(N)] for _ in range(N)])
#b = Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]])
#a = Tensor([[[[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]]]])
#b = Tensor([[[[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]]]])
ini = time.time()
c = a + b
print("\n#####2######")
fim = time.time()
print(fim-ini)
print("\n\n")
a = [[random.uniform(0, 1) for _ in range(N)] for _ in range(N)]
b = [[random.uniform(0, 1) for _ in range(N)] for _ in range(N)]
ini = time.time()
result_matrix = matrix_sum(a, b)
fim = time.time()
print(fim-ini)
print("\n\n")
ini = time.time()
a = np.random.rand(N, N)
b = np.random.rand(N, N)
result_matrix = a + b
fim = time.time()
print(fim-ini)
"""
return self

70
test.py Normal file
View file

@ -0,0 +1,70 @@
def matrix_sum(matrix1, matrix2):
# Check if the matrices can be multiplied
if len(matrix1[0]) != len(matrix2):
raise ValueError("Matrices cannot be multiplied. Inner dimensions must match.")
# Initialize the result matrix with zeros
result = [[0 for _ in range(len(matrix2[0]))] for _ in range(len(matrix1))]
# Perform matrix multiplication
for i in range(len(matrix1)):
for j in range(len(matrix2[0])):
result[i][j] += matrix1[i][i] + matrix2[i][j]
return result
if __name__ == "__main__":
import norch
a = norch.Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]])
b = norch.Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]])
import time
import random
import numpy as np
#a = Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]])#.to("cuda")
#b = Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]])#.to("cuda")
#c = Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]])#.to("cuda")
#d = b-c
print(a ** 2)
#print(a ** 2)
"""#print(a)
N = 1000
a = Tensor([[random.uniform(0, 1) for _ in range(N)] for _ in range(N)])
b = Tensor([[random.uniform(0, 1) for _ in range(N)] for _ in range(N)])
#b = Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]])
#a = Tensor([[[[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]]]])
#b = Tensor([[[[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]]]])
ini = time.time()
c = a + b
print("\n#####2######")
fim = time.time()
print(fim-ini)
print("\n\n")
a = [[random.uniform(0, 1) for _ in range(N)] for _ in range(N)]
b = [[random.uniform(0, 1) for _ in range(N)] for _ in range(N)]
ini = time.time()
result_matrix = matrix_sum(a, b)
fim = time.time()
print(fim-ini)
print("\n\n")
ini = time.time()
a = np.random.rand(N, N)
b = np.random.rand(N, N)
result_matrix = a + b
fim = time.time()
print(fim-ini)
"""

3
test.sh
View file

@ -1,5 +1,4 @@
cd build
make
cd ..
cd norch
python3 tensor.py
python3 test.py