求一个完整的代码

[fannyqiq]

求哪位大侠给一个完整的代码，要求很简单，实现：
两个矩阵相加。（矩阵就4×4的吧）
包括主机端和设备端的空间开辟、主机端两个矩阵的初始化、主机端与设备端的数据拷贝、Kernel、计算结果传回主机端、输出结果。
本人菜鸟。

[fannyqiq]

咋就没有人呢，不难的啊，自己写的就是各种小问题。求高手们不惜指教啊~

[cai_niao_yi_zhi]

//////////////////主机端代码///////////////// // includes, system #include <stdlib.h> #include <stdio.h> #include <math.h> // includes, kernels #include <VecAdd_kernel.cu> int main() { const unsigned int N = 10; //内存中分配向量空间，并赋初值 float h_A[N], h_B[N], h_C[N]; for(int i=0; i<N; i++) { h_A[i] = i*1.0f; h_B[i] = (N-i)*1.0f; } //显存中分配向量空间 size_t size = N * sizeof(float); float *d_A; cudaMalloc((void**)&d_A, size); float *d_B; cudaMalloc((void**)&d_B, size); float *d_C; cudaMalloc((void**)&d_C, size); for(int i=0; i<N; i++) d_C[i] = 0.0f; //从内存向显存拷贝向量 cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice); cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice); //启动kernel int threadsPerBlock = 5; int blocksPerGrid = (N + threadsPerBlock - 1)/threadsPerBlock; VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N); //从显存向内存拷回结果 cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost); //释放显存空间 cudaFree(d_A); cudaFree(d_B); cudaFree(d_C); //打印结果 for(int i=0; i<N; i++) { printf("%5.0f\n", h_C[i]); } return 0; } ///////////////////设备端代码/////////////////// #ifndef _VECADD_KERNEL_H_ #define _VECADD_KERNEL_H_ __global__ void VecAdd(float *A, float *B, float *C, unsigned int N) { int i = blockDim.x * blockIdx.x + threadIdx.x; if(i < N) C[i] = A[i] + B[i]; } #endif // #ifndef _VECADD_KERNEL_H_

[fannyqiq]

谢谢！ 请问设备端代码开始的 #ifndef _VECADD_KERNEL_H_ #define _VECADD_KERNEL_H_ 是什么用途？

[fannyqiq]

要求的是2维的矩阵相加，跟一维的向量在分配空间、BlocksPerGrid和ThreadsPerBlock的确定细节上还是有区别的。还是谢谢你~

[fannyqiq]

要是可以，就麻烦再给个矩阵的呗

[cai_niao_yi_zhi]

#ifndef _VECADD_KERNEL_H_ #define _VECADD_KERNEL_H_ 是为了防止多次包含！ 2维矩阵的我得思考下，我也是初学者！

[fannyqiq]

好复杂的样子，我看看，谢谢！

[fannyqiq]

对了 ，for(int i=0; i<N; i++) d_C[i] = 0.0f;好像不对的，设备端不需要初始化，只要开辟空间就好了~

[fannyqiq]

求矩阵相加代码

[cai_niao_yi_zhi]

这是一个二维矩阵相乘的代码，你看看，原理是一样的： // includes, system #include <stdlib.h> #include <stdio.h> #include <string.h> #include <math.h> #define BLOCK_SIZE 2 typedef struct { size_t width; size_t height; size_t pitch; float *element; }Matrix; ///////////////////////////设备端代码////////////////////////////// __global__ void MatMulKernel(Matrix A, Matrix B, Matrix C) { float Cvalue = 0; unsigned int row = blockIdx.y * blockDim.y + threadIdx.y; unsigned int col = blockIdx.x * blockDim.x + threadIdx.x; if(row<C.height && col<C.width) { for(int i=0; i<A.width; i++) { Cvalue += A.element[row*C.pitch/sizeof(float)+i]*B.element[i*B.pitch/sizeof(float)+col]; C.element[row*C.pitch/sizeof(float)+col] = Cvalue; } } } ///////////////////////////主机端代码////////////////////////////// //定义矩阵相乘函数 void MatMul(const Matrix A, const Matrix B, Matrix C) { //设备端分配内存 Matrix d_A; d_A.width = A.width; d_A.height = A.height; cudaMallocPitch((void**)&d_A.element, &d_A.pitch, A.width*sizeof(float), A.height); cudaMemcpy2D(d_A.element, d_A.pitch, A.element, A.pitch, A.width*sizeof(float), A.height, cudaMemcpyHostToDevice); Matrix d_B; d_B.width = B.width; d_B.height = B.height; cudaMallocPitch((void**)&d_B.element, &d_B.pitch, B.width*sizeof(float), B.height); cudaMemcpy2D(d_B.element, d_B.pitch, B.element, B.pitch, B.width*sizeof(float), B.height, cudaMemcpyHostToDevice); Matrix d_C; d_C.width = C.width; d_C.height = C.height; cudaMallocPitch((void**)&d_C.element, &d_C.pitch, C.width*sizeof(float), C.height); //定义网格形状 dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE, 1); dim3 dimGrid((C.width-1+dimBlock.x)/dimBlock.x, (C.height-1+dimBlock.y)/dimBlock.y, 1); //调用内核函数 MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C); //将结果从显存拷贝到内存 cudaMemcpy2D(C.element, C.pitch, d_C.element, d_C.pitch, C.width*sizeof(float), C.height, cudaMemcpyDeviceToHost); //打印结果 for(unsigned int i=0; i<C.height; i++) { for(unsigned int j=0; j<C.width; j++) printf("%8.2f", C.element[i*C.width+j]); printf("\n"); } //释放设备端内存 cudaFree(d_A.element); cudaFree(d_B.element); cudaFree(d_C.element); } int main() { //定义三个矩阵结构体 Matrix h_A, h_B, h_C; //h_A矩阵结构体初始化 h_A.width = 4; h_A.height = 6; h_A.element = (float*)malloc(h_A.width*h_A.height*sizeof(float)); h_A.pitch = h_A.width*sizeof(float); for(unsigned int i=0; i<h_A.height; i++) for(unsigned int j=0; j<h_A.width; j++) h_A.element[i*h_A.width+j] = float(i*h_A.width+j); //h_B矩阵结构体初始化 h_B.width = 6; h_B.height = 4; h_B.element = (float*)malloc(h_B.width*h_B.height*sizeof(float)); h_B.pitch = h_B.width*sizeof(float); for(unsigned int i=0; i<h_B.height; i++) for(unsigned int j=0; j<h_B.width; j++) h_B.element[i*h_B.width+j] = float(i*h_B.width+j); //h_C矩阵结构体初始化 h_C.width = h_B.width; h_C.height = h_A.height; h_C.element = (float*)malloc(h_C.width*h_C.height*sizeof(float)); h_C.pitch = h_C.width*sizeof(float); //调用矩阵相乘函数进行计算，并输出结果 MatMul(h_A, h_B, h_C); //释放主机端内存 free(h_A.element); free(h_B.element); free(h_C.element); return 0; }

[cai_niao_yi_zhi]

引用 7 楼 fannyqiq 的回复:

对了 ，for(int i=0; i<N; i++) d_C[i] = 0.0f;好像不对的，设备端不需要初始化，只要开辟空间就好了~

嗯 这里确实错了，主机端不能访问设备端内存，而且也不用初始化。呵呵，多谢指出错误啊！

[cai_niao_yi_zhi]

引用 9 楼 cai_niao_yi_zhi 的回复:

这是一个二维矩阵相乘的代码，你看看，原理是一样的： // includes, system #include <stdlib.h> #include <stdio.h> #include <string.h> #include <math.h> #define BLOCK_SIZE 2 typedef struct { size_t widt……

设备端代码有点错误，应该是这样： ///////////////////////////设备端代码////////////////////////////// __global__ void MatMulKernel(Matrix A, Matrix B, Matrix C) { float Cvalue = 0; unsigned int row = blockIdx.y * blockDim.y + threadIdx.y; unsigned int col = blockIdx.x * blockDim.x + threadIdx.x; float *rowA = (float*)((char*)A.element + row*A.pitch); float *rowC = (float*)((char*)C.element + row*C.pitch); if(row<C.height && col<C.width) { for(unsigned int i=0; i<A.width; i++) { float *rowB = (float*)((char*)B.element + i*B.pitch); Cvalue += rowA[i]*rowB[col]; rowC[col] = Cvalue; } } }