cuda笔记-一个Block多线程求卷积
发布日期:2021-06-30 10:46:04
浏览次数:2
分类:技术文章
本文共 3482 字,大约阅读时间需要 11 分钟。
最近在学cuda,发现自己数学方面的知识不太够,C语言的知识也有待加强。
这里记录个笔记对矩阵求卷积。
逻辑是这样的:
1. 先CUDA生成一个16*16的矩阵;
2. 将这16*16的矩阵,外面包一层0,也就变成18*18的矩阵。
3. 然后再开18*18个线程,进行矩阵的卷积
程序运行截图如下:
源码如下:
#include "cuda_runtime.h"#include "device_launch_parameters.h"#include "curand.h"#include "curand_kernel.h"#include#include using namespace std;#pragma comment(lib, "cudart.lib")#pragma comment(lib, "curand.lib")#define N 16__global__ void Matrix_convolution(float *a, float *b) { int x = threadIdx.x; int y = threadIdx.y; __shared__ float shared[(N + 2) * (N + 2)]; shared[y * (N + 2) + x] = a[y * (N + 2) + x]; __syncthreads(); int convolution[9] = { 1, 1, 1, 1, -8, 1, 1, 1, 1 }; //对shared进行卷积 int pos = y * (N + 2) + x; //第一行和最后一行不要卷积 if (pos < N + 2 || (pos >(N + 2) * (N + 1))) { return; } //最左边和最右边一行不要卷积 if (pos % (N + 2) == 0 || pos % (N + 2) == N + 1) { return; } //卷积的9个值 float a00 = shared[(y * (N + 2) + x) - (N + 2) - 1]; float a01 = shared[(y * (N + 2) + x) - (N + 2)]; float a02 = shared[(y * (N + 2) + x) - (N + 2) + 1]; float a10 = shared[(y * (N + 2) + x) - 1]; float a11 = shared[(y * (N + 2) + x)]; float a12 = shared[(y * (N + 2) + x) + 1]; float a20 = shared[(y * (N + 2) + x) + (N + 2) - 1]; float a21 = shared[(y * (N + 2) + x) + (N + 2)]; float a22 = shared[(y * (N + 2) + x) + (N + 2) + 1]; float ret = convolution[0] * a00 + convolution[1] * a01 + convolution[2] * a02 + convolution[3] * a10 + convolution[4] * a11 + convolution[5] * a12 + convolution[6] * a20 + convolution[7] * a21 + convolution[8] * a22; //目前在多少行 int rowCount = (y * (N + 2) + x) / (N + 2); int posOffset = (N + 2) + (rowCount - 1) * 2 + 1; b[y * (N + 1) + x - posOffset] = ret;}void Matrix_init_gpu(float *a) { curandGenerator_t gen; curandCreateGenerator(&gen, CURAND_RNG_PSEUDO_MRG32K3A); curandSetPseudoRandomGeneratorSeed(gen, 11ULL); curandGenerateUniform(gen, a, N * N);}int main() { float *p_original_gpu; float *p_original_cpu; float *p_final_cpu; float *p_final_gpu; float *p_ret_gpu; float *p_ret_cpu; p_original_cpu = (float*)malloc(N * N * sizeof(float)); p_final_cpu = (float*)malloc((N + 2) * (N + 2) * sizeof(float)); p_ret_cpu = (float*)malloc(N * N * sizeof(float)); cudaMalloc((void**)&p_original_gpu, N * N * sizeof(float)); cudaMalloc((void**)&p_final_gpu, (N + 2) * (N + 2) * sizeof(float)); cudaMalloc((void**)&p_ret_gpu, N * N * sizeof(float)); Matrix_init_gpu(p_original_gpu); cudaMemcpy(p_original_cpu, p_original_gpu, N * N * sizeof(float), cudaMemcpyDeviceToHost); for(int i = 0; i < N * N; i++){ if (i % N == 0) printf("\n"); cout << p_original_cpu[i] << " "; } //开始填充数据 for (int i = 0; i < (N + 2) * (N + 2); i++) { p_final_cpu[i] = 0; } int pos = 0; for (int i = N + 2; i < (N + 2) * (N + 1); i++) { if (i % (N + 2) != 0 && i % (N + 2) != N + 1) { p_final_cpu[i] = p_original_cpu[pos++]; } } cout << "\n\n填充数据:" << endl; for (int i = 0; i < (N + 2) * (N + 2); i++) { if (i % (N + 2) == 0) printf("\n"); cout << p_final_cpu[i] << " "; } cout << "\n\n最后结果:" << endl; cudaMemcpy(p_final_gpu, p_final_cpu, (N + 2) * (N + 2) * sizeof(float), cudaMemcpyHostToDevice); Matrix_convolution << <1, (N + 2) * (N + 2) >> >(p_final_gpu, p_ret_gpu); cudaMemcpy(p_ret_cpu, p_ret_gpu, N * N * sizeof(float), cudaMemcpyDeviceToHost); for (int i = 0; i < N * N; i++) { if (i % N == 0) printf("\n"); cout << p_ret_cpu[i] << " "; } cudaFree(p_original_gpu); cudaFree(p_final_gpu); cudaFree(p_ret_gpu); free(p_original_cpu); free(p_final_cpu); free(p_ret_cpu); getchar(); return 0;}
转载地址:https://it1995.blog.csdn.net/article/details/111219482 如侵犯您的版权,请留言回复原文章的地址,我们会给您删除此文章,给您带来不便请您谅解!
发表评论
最新留言
路过按个爪印,很不错,赞一个!
[***.219.124.196]2024年04月16日 10时28分45秒
关于作者
喝酒易醉,品茶养心,人生如梦,品茶悟道,何以解忧?唯有杜康!
-- 愿君每日到此一游!
推荐文章
数据库索引 & 为什么要对数据库建立索引 / 数据库建立索引为什么会加快查询速度
2019-04-30
IEEE与APA引用格式
2019-04-30
research gap
2019-04-30
pytorch训练cifar10数据集查看各个种类图片的准确率
2019-04-30
Python鼠标点击图片,获取点击点的像素坐标
2019-04-30
路径规划(一) —— 环境描述(Grid Map & Feature Map) & 全局路径规划(最优路径规划(Dijkstra&A*star) & 概率路径规划(PRM&RRT))
2019-04-30
RRT算法(快速拓展随机树)的Python实现
2019-04-30
D*算法
2019-04-30
强化学习(四) —— Actor-Critic演员评论家 & code
2019-04-30
RESTful API
2019-04-30
优化算法(四)——粒子群优化算法(PSO)
2019-04-30
数据在Oracle中的存储
2019-04-30
轨迹规划 trajectory planning
2019-04-30
AGV自动导引运输车
2019-04-30
Trie树(字典树)
2019-04-30