diff --git a/build/cpu.o b/build/cpu.o index b7ee569..b6465ab 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 b3b4837..51d5b7b 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 e23e7be..1f13ba8 100755 Binary files a/build/libtensor.so and b/build/libtensor.so differ diff --git a/build/tensor.o b/build/tensor.o index 34580b5..1d75a22 100644 Binary files a/build/tensor.o and b/build/tensor.o differ diff --git a/csrc/cpu.cpp b/csrc/cpu.cpp index cadf732..93623e9 100644 --- a/csrc/cpu.cpp +++ b/csrc/cpu.cpp @@ -25,9 +25,9 @@ void elementwise_mul_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_ } } -void pow_tensor_cpu(Tensor* tensor, float power) { +void pow_tensor_cpu(Tensor* tensor, float power, float* result_data) { for (int i = 0; i < tensor->size; i++) { - tensor->data[i] = powf(tensor->data[i], power); + result_data[i] = powf(tensor->data[i], power); } } diff --git a/csrc/cpu.h b/csrc/cpu.h index 99b88fe..20c0f2b 100644 --- a/csrc/cpu.h +++ b/csrc/cpu.h @@ -2,4 +2,4 @@ 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); \ No newline at end of file +void pow_tensor_cpu(Tensor* tensor, float power, float* result_data); \ No newline at end of file diff --git a/csrc/cuda.cu b/csrc/cuda.cu index 4664b40..d433677 100644 --- a/csrc/cuda.cu +++ b/csrc/cuda.cu @@ -97,17 +97,17 @@ __host__ void elementwise_mul_tensor_cuda(Tensor* tensor1, Tensor* tensor2, floa cudaDeviceSynchronize(); } -__global__ void pow_tensor_cuda_kernel(float* data, float power, int size) { +__global__ void pow_tensor_cuda_kernel(float* data, float power, float* result_data, int size) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < size) { - data[i] = powf(data[i], power); + result_data[i] = powf(data[i], power); } } -__host__ void pow_tensor_cuda(Tensor* tensor, float power) { +__host__ void pow_tensor_cuda(Tensor* tensor, float power, float* result_data) { int number_of_blocks = (tensor->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; - pow_tensor_cuda_kernel<<>>(tensor->data, power, tensor->size); + pow_tensor_cuda_kernel<<>>(tensor->data, power, result_data, tensor->size); cudaError_t error = cudaGetLastError(); if (error != cudaSuccess) { diff --git a/csrc/cuda.h b/csrc/cuda.h index ea30b04..c1a1404 100644 --- a/csrc/cuda.h +++ b/csrc/cuda.h @@ -13,7 +13,7 @@ __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); + __global__ void pow_tensor_cuda_kernel(float* data, float power, float* result_data, int size); + __host__ void pow_tensor_cuda(Tensor* tensor, float power, float* result_data); #endif /* CUDA_KERNEL_H_ */ diff --git a/csrc/tensor.cpp b/csrc/tensor.cpp index 5f2b26a..05baef0 100644 --- a/csrc/tensor.cpp +++ b/csrc/tensor.cpp @@ -304,12 +304,40 @@ extern "C" { return create_tensor(result_data, shape, ndim, device); } - void pow_tensor(Tensor* tensor, float power) { + Tensor* pow_tensor(Tensor* tensor, float power) { + 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) { - pow_tensor_cuda(tensor, power); + + float* result_data; + cudaMalloc((void **)&result_data, tensor->size * sizeof(float)); + pow_tensor_cuda(tensor, power, result_data); + return create_tensor(result_data, shape, ndim, device); } else { - pow_tensor_cpu(tensor, power); + float* result_data = (float*)malloc(tensor->size * sizeof(float)); + if (result_data == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + pow_tensor_cpu(tensor, power, result_data); + return create_tensor(result_data, shape, ndim, device); } } diff --git a/csrc/tensor.h b/csrc/tensor.h index c87f2af..a075fb7 100644 --- a/csrc/tensor.h +++ b/csrc/tensor.h @@ -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); - void pow_tensor(Tensor* tensor, float power); + Tensor* pow_tensor(Tensor* tensor, float power); void to_device(Tensor* tensor, char* device); } diff --git a/norch/__pycache__/tensor.cpython-38.pyc b/norch/__pycache__/tensor.cpython-38.pyc index 688c842..3f68c9e 100644 Binary files a/norch/__pycache__/tensor.cpython-38.pyc and b/norch/__pycache__/tensor.cpython-38.pyc differ diff --git a/norch/tensor.py b/norch/tensor.py index 789c788..97c3ec2 100644 --- a/norch/tensor.py +++ b/norch/tensor.py @@ -204,8 +204,14 @@ class Tensor: power = ctypes.c_float(power) Tensor._C.pow_tensor.argtypes = [ctypes.POINTER(CTensor), ctypes.c_float] - Tensor._C.pow_tensor.restype = None + Tensor._C.pow_tensor.restype = ctypes.POINTER(CTensor) - Tensor._C.pow_tensor(self.tensor, power) + result_tensor_ptr = Tensor._C.pow_tensor(self.tensor, power) - return self \ No newline at end of file + 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 + + return result_data \ No newline at end of file diff --git a/test.py b/test.py index ed5aff9..32946a6 100644 --- a/test.py +++ b/test.py @@ -28,7 +28,9 @@ if __name__ == "__main__": #c = Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]])#.to("cuda") #d = b-c - print(a ** 2) + + b = a ** 2 + print(b) #print(a ** 2) """#print(a)