Mathis/CUDANet
1
0
Fork 0
mirror of https://github.com/lordmathis/CUDANet.git synced 2025-11-05 17:34:21 +00:00
Code Issues Packages Projects Releases Wiki Activity

Fix matmul and max reduce memcheck errors

Browse Source
This commit is contained in:
Matúš Námešný
2024-04-21 22:11:02 +02:00
parent 9a6152469a
commit bdbb3f978e
4 changed files with 16 additions and 41 deletions
Download Patch File Download Diff File Expand all files Collapse all files

35
src/kernels/matmul.cu
View File

@@ -3,7 +3,6 @@
using namespace CUDANet;
__global__ void Kernels::mat_vec_mul(
const float* __restrict__ d_matrix,
const float* __restrict__ d_vector,
@@ -13,32 +12,17 @@ __global__ void Kernels::mat_vec_mul(
) {
int tid = blockDim.x * blockIdx.x + threadIdx.x;
__shared__ float shared[BLOCK_SIZE];
if (tid < h) {
float temp = 0.0f;
float temp = 0.0f;
#pragma unroll
for (unsigned int i = 0; i < (w + BLOCK_SIZE - 1) / BLOCK_SIZE; i++) {
if (i * BLOCK_SIZE + threadIdx.x < w) {
shared[threadIdx.x] = d_vector[i * BLOCK_SIZE + threadIdx.x];
} else {
shared[threadIdx.x] = 0.0f;
for (unsigned int j = 0; j < w; j++) {
temp += d_matrix[tid * w + j] * d_vector[j];
}
__syncthreads();
#pragma unroll
for (unsigned int j = 0; j < BLOCK_SIZE; j++) {
temp += d_matrix[tid * w + i * BLOCK_SIZE + j] * shared[j];
}
__syncthreads();
d_output[tid] = temp;
}
d_output[tid] = temp;
}
__global__ void Kernels::vec_vec_add(
const float* __restrict__ d_vector1,
const float* __restrict__ d_vector2,
@@ -52,7 +36,6 @@ __global__ void Kernels::vec_vec_add(
d_output[tid] = d_vector1[tid] + d_vector2[tid];
}
__global__ void Kernels::vec_scalar_sub(
const float* __restrict__ d_src,
float* __restrict__ d_out,
@@ -66,21 +49,19 @@ __global__ void Kernels::vec_scalar_sub(
d_out[tid] = d_src[tid] - d_scalar[0];
}
__global__ void Kernels::vec_scalar_div(
const float* __restrict__ d_src,
float* __restrict__ d_out,
const float* __restrict__ d_scalar,
const unsigned int len
) {
int tid = blockDim.x * blockIdx.x + threadIdx.x;
int tid = blockDim.x * blockIdx.x + threadIdx.x;
if (tid >= len) {
return;
}
d_out[tid] = d_src[tid] / d_scalar[0];
}
__global__ void Kernels::vec_exp(
const float* __restrict__ src,
float* __restrict__ dst,
@@ -107,7 +88,7 @@ __global__ void Kernels::max_reduce(
shared_max[threadIdx.x] = d_vector[i];
} else {
shared_max[threadIdx.x] = -INFINITY;
}
}
__syncthreads();
@@ -129,7 +110,7 @@ __global__ void Kernels::sum_reduce(
const unsigned int len
) {
__shared__ float partial_sum[BLOCK_SIZE];
int i = blockIdx.x * blockDim.x + threadIdx.x;
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < len) {
partial_sum[threadIdx.x] = d_vector[i];

10
src/utils/vector.cu
View File

@@ -27,17 +27,7 @@ void Utils::clear(float* d_vec, const unsigned int length) {
void Utils::max(float* d_vec, float* d_max, const unsigned int length) {
const int grid_size = (length + BLOCK_SIZE - 1) / BLOCK_SIZE;
std::cout << "grid_size: " << grid_size << ", length: " << length << std::endl;
CUDA_CHECK(cudaGetLastError());
Kernels::max_reduce<<<grid_size, BLOCK_SIZE>>>(d_vec, d_max, length);
std::cout << "input: " << std::endl;
print_vec(d_vec, length);
std::cout << "max: " << std::endl;
print_vec(d_max, length);
CUDA_CHECK(cudaGetLastError());
int remaining = grid_size;