中国农村建设投资有限公司网站首页,360免费建站模板,四川建设网证书查询平台官网,网站维护要多久时间1. 知识点
nthreads的取值#xff0c;不能大于block能取值的最大值。一般可以直接给512、256#xff0c;性能就是比较不错的 (input_size block_size - 1) / block_size;是向上取整 对于一维数组时#xff0c;采用只定义layout的x维度#xff0c;若处理的是二维#xff…1. 知识点
nthreads的取值不能大于block能取值的最大值。一般可以直接给512、256性能就是比较不错的 (input_size block_size - 1) / block_size;是向上取整 对于一维数组时采用只定义layout的x维度若处理的是二维则可以考虑定义x、y维度例如处理的是图像关于把数据视作一维时索引的计算 以下是通用的计算公式 Pseudo code:
position 0
for i in range(6):position * dims[i]position indexs[i]例如当只使用x维度时实际上dims [1, 1, gd, 1, 1, bd]indexs [0, 0, bi, 0, 0, ti] 因为0和1的存在上面的循环则可以简化为idx threadIdx.x blockIdx.x * blockDim.x即idx ti bi * bd
2. main.cpp文件
#include cuda_runtime.h
#include stdio.h#define checkRuntime(op) __check_cuda_runtime((op), #op, __FILE__, __LINE__)bool __check_cuda_runtime(cudaError_t code, const char* op, const char* file, int line){if(code ! cudaSuccess){ const char* err_name cudaGetErrorName(code); const char* err_message cudaGetErrorString(code); printf(runtime error %s:%d %s failed. \n code %s, message %s\n, file, line, op, err_name, err_message); return false;}return true;
}void vector_add(const float* a, const float* b, float* c, int ndata);int main(){const int size 3;float vector_a[size] {2, 3, 2};float vector_b[size] {5, 3, 3};float vector_c[size] {0};float* vector_a_device nullptr;float* vector_b_device nullptr;float* vector_c_device nullptr;checkRuntime(cudaMalloc(vector_a_device, size * sizeof(float)));checkRuntime(cudaMalloc(vector_b_device, size * sizeof(float)));checkRuntime(cudaMalloc(vector_c_device, size * sizeof(float)));checkRuntime(cudaMemcpy(vector_a_device, vector_a, size * sizeof(float), cudaMemcpyHostToDevice));checkRuntime(cudaMemcpy(vector_b_device, vector_b, size * sizeof(float), cudaMemcpyHostToDevice));vector_add(vector_a_device, vector_b_device, vector_c_device, size);checkRuntime(cudaMemcpy(vector_c, vector_c_device, size * sizeof(float), cudaMemcpyDeviceToHost));for(int i 0; i size; i){printf(vector_c[%d] %f\n, i, vector_c[i]);}checkRuntime(cudaFree(vector_a_device));checkRuntime(cudaFree(vector_b_device));checkRuntime(cudaFree(vector_c_device));return 0;
}先定义三个数组: a, b, c 再用cudaMalloc()在GPU上开辟三个内存在GPU上让a b 并且让结果存储进c上再把c的内存从GPU上放到Host上输出
3. 案例.cu文件
#include stdio.h
#include cuda_runtime.h__global__ void vector_add_kernel(const float* a, const float* b, float* c, int ndata){int idx threadIdx.x blockIdx.x * blockDim.x;if(idx ndata) return;/* dims indexsgridDim.z blockIdx.zgridDim.y blockIdx.ygridDim.x blockIdx.xblockDim.z threadIdx.zblockDim.y threadIdx.yblockDim.x threadIdx.xPseudo code:position 0for i in 6:position * dims[i]position indexs[i]*/c[idx] a[idx] b[idx];
}void vector_add(const float* a, const float* b, float* c, int ndata){const int nthreads 512;int block_size ndata nthreads ? ndata : nthreads; // 如果ndata nthreads 那block_size ndata就够了int grid_size (ndata block_size - 1) / block_size; // 其含义是我需要多少个blocks可以处理完所有的任务printf(block_size %d, grid_size %d\n, block_size, grid_size);vector_add_kernelgrid_size, block_size, 0, nullptr(a, b, c, ndata);// 在核函数执行结束后通过cudaPeekAtLastError获取得到的代码来知道是否出现错误// cudaPeekAtLastError和cudaGetLastError都可以获取得到错误代码// cudaGetLastError是获取错误代码并清除掉也就是再一次执行cudaGetLastError获取的会是success// 而cudaPeekAtLastError是获取当前错误但是再一次执行cudaPeekAtLastError或者cudaGetLastErro拿到的还是那个错cudaError_t code cudaPeekAtLastError();if(code ! cudaSuccess){ const char* err_name cudaGetErrorName(code); const char* err_message cudaGetErrorString(code); printf(kernel error %s:%d test_print_kernel failed. \n code %s, message %s\n, __FILE__, __LINE__, err_name, err_message); }
}两个注意的点 像这个案例他就三个数相加其实启动三个线程就足够了但是一般block给的是512 256所以要设定一下如果数组的长度小于256/512, 就直接用数组的长度的线程数就好。这里就是3个线程 如果线程索引大于了数组的长度就直接返回了不然就访问了不知道在哪里的内存了
