电脑怎样才能用共享gpu内存,笔记本共享gpu内存

电脑如何共享gpu内存，笔记本如何共享gpu内存？

CUDA的内存大致可以分为两类

板载内存主要包括全局内存、局部内存、常量内存、纹理内存等。片上存储器主要包括寄存器和共享存储器。不同类型的内存各有特点，但片上内存通常比板载显存快，寄存器是各种内存中最快的。本文重点介绍共享内存的基础知识和应用实例。

01

检查视频卡上的共享内存信息。

CUDA提供了两个函数，cudaGetDeviceCount和cudaGetDeviceProperties，分别用于获取CUDA设备数量和CUDA设备属性。通过调用这两个函数，可以方便地获得共享内存信息和其他CUDA设备信息

//显示CUDA设备信息void sho _ GPU _ info(void){ int device count；//获取CUDA设备总数cudaGetDeviceCount(device count)；//分别获取每个CUDA设备的信息(int I=0；ideviceCountI) {//定义用于存储信息的结构cudaDeviceProp devProp//将第I个CUDA设备的信息写入结构cudagetdeviceproperties(devprop，I)；Std:cout '使用GPU设备' I' : '开发工具。名称STD 3360: endlStd:cout '设备的总全局内存' devprop.totalglobalmem/1024/1024' MB ' STD 3360: endl；STD : cout ' SM SM:' devprop的数量。多处理器计数标准 endlStd:cout '每个线程块的共享内存大小' devprop.sharedmemperblock/1024.0' KB ' STD : endl；Std:cout '每个线程块的最大线程数' devprop . maxthreadsperblock STD : endl；Std:cout '设备上线程块中可用的32位寄存器的数量' devprop . regsperblock STD : endl；Std:cout '每个EM的最大线程数' devprop。MaxThreadSpermultipProcessor STD 33603360 endl；Std:cout '每个EM的最大线程束数' devprop.maxthreadspermultprocessor/32STD : endl；Std:cout '设备上的多处理器数量' devprop。多处理器计数标准ENDLSTD : cout '=====================================================' STD : endl；}}

运行上述函数来获取共享内存信息和其他设备信息。如下图所示，我可以为每个线程块使用的最大共享内存是48 KB。

02

共享内存的特征

共享内存的主要特点是“共享”，即同一个线程块中的所有线程都可以读写这个内存块，所以“共享”是针对同一个线程块中的所有线程的。一旦定义了共享内存并指定了大小，系统将为所有线程块分配相同大小的共享内存。例如，如果定义了一个8字节的无符号字符共享内存，那么所有线程块都将被分配一个8字节的无符号字符共享内存。如下图所示

共享内存是在CUDA内核函数中定义的，通常有两种方式静态方式和动态方式。

_ _全球

__ shared_memory_kernel(uchar inputs, int ro, int col){ int x = threadIdx.x + blockDim.x blockIdx.x; //col int y = threadIdx.y + blockDim.y blockIdx.y; //ro if (x < col && y < ro) { __shared__ uchar s[8]; //定义的指定大小为8 bytes，每个线程块都被分配8 bytes的共享内存 . . . }}

__global__ shared_memory_kernel(uchar inputs, int ro, int col){ int x = threadIdx.x + blockDim.x blockIdx.x; //col int y = threadIdx.y + blockDim.y blockIdx.y; //ro if (x < col && y < ro) { extern __shared__ uchar s[]; //定义的时候不指定大小 . . . }}void shared_memory_test(void){ . . . //传入的第1个参数block_num为线程块总数 //第2个参数thread_num为每个线程块包含的线程数 //第3个参数8为共享内存大小，所以动态共享内存大小通过第3个参数传入 shared_memory_kernel<<<block_num, thread_num, 8>>>(inputs, ro, col); . . .}

需要注意动态定义共享内存时，调用核函数传入的数值必须以字节byte为单位，所以如果定义的共享内存不是byte类型，数值必须乘以类型占用的字节数。比如要动态定义长度为8的float类型共享内存，那么传入核函数的数值为8sizeof(float)。

shared_memory_kernel<<<block_num, thread_num, 8 sizeof(float)>>>(inputs, ro, col);

03

共享内存的应用例子

前文我们讲的数组元素规约求和算法，使用CUDA全局内存来存储数据

CUDA加速——基于规约思想的数组元素求和

我们知道全局内存属于板载显存，而共享内存属于片上内存，共享内存的读写速度比全局内存快得多。在前文代码的核函数中有个for循环需要多次读写全局内存，全局内存本身就很慢，而且如果不是连续访问会更慢，本文我们尝试使用共享内存来代替全局内存实现前文讲的规约求和算法。

由于前文的规约算法是在不同线程块分别进行的，而共享内存又具有线程块内共享的特性，故共享内存正好适合此应用场景。

前文的规约结构

本文使用共享内存的规约结构

下面我们比较使用共享内存的核函数与前文使用全局内存的核函数

//使用全局内存__global__ void cal_sum_ker0(float Para, float blocksum_cuda){ int tid = blockIdx.x blockDim.x + threadIdx.x; if(tid < N) { for (int index = 1; index < blockDim.x; index = (index2)) { if (threadIdx.x % (index2) == 0) { Para[tid] += Para[tid + index]; } __syncthreads(); //同步，以防止归约过程中某个线程运行速度过快导致计算错误 } if(threadIdx.x == 0) //整个数组相加完成后，将共享内存数组0号元素的值赋给全局内存数组0号元素 blocksum_cuda[blockIdx.x] = Para[tid]; }}//使用共享内存//blockIdx.x为线程块的ID号//blockDim.x每个线程块中包含的线程总个数//threadIdx.x为每个线程块中的线程ID号__global__ void cal_sum_ker(float Para, float blocksum_cuda){ int tid = blockIdx.x blockDim.x + threadIdx.x; if(tid < N) { //动态方式定义float型共享内存 extern __shared__ float s_Para[]; //线程块中的每个线程负责把其对应的数据从全局内存加载到共享内存 s_Para[threadIdx.x] = Para[tid]; __syncthreads(); //块内线程同步，等待线程块内所有线程加载数据完毕 for (int index = 1; index < blockDim.x; index = (index2)) { if (threadIdx.x % (index2) == 0) { //在for循环中使用共享内存实现规约，避免频繁读写全局内存 s_Para[threadIdx.x] += s_Para[threadIdx.x + index]; } __syncthreads(); //块内线程同步，以防止归约过程中某个线程运行速度过快导致计算错误 } if(threadIdx.x == 0) //将共享内存数组0号元素的值赋给全局内存数组 blocksum_cuda[blockIdx.x] = s_Para[threadIdx.x]; }}

接着在测试代码中分别调用上方两个核函数。调用时指定共享内存的长度为每个线程块包含的线程数

void cal_sum_test() { Timer_Us2 timer; //定义CPU端数组 float test_d = (float )malloc(N sizeof(float)); for (long long i = 0; i < N; i++) { test_d[i] = 0.5; } dim3 sumblock(512);//设置每个线程块有512个线程 dim3 sumgrid(((N%sumblock.x) ? (N/sumblock.x + 1) : (N/sumblock.x))); float test_d_cuda; float blocksum_cuda; float blocksum_host = (float )malloc(sizeof(float) sumgrid.x); cudaMalloc((void )&test_d_cuda, sizeof(float) N); cudaMalloc((void )&blocksum_cuda, sizeof(float) sumgrid.x); timer.start_timer(); //将数据从CPU端拷贝到GPU端 cudaMemcpy(test_d_cuda, test_d, sizeof(float) N, cudaMemcpyHostToDevice); //调用使用全局内存规约的核函数 cal_sum_ker0 << < sumgrid, sumblock>> > (test_d_cuda, blocksum_cuda); //将所有线程块的规约结果从GPU端拷贝到CPU端 cudaMemcpy(blocksum_host, blocksum_cuda, sizeof(float) sumgrid.x, cudaMemcpyDeviceToHost); //在CPU端对所有线程块的规约求和结果做串行求和 double sum = 0.0; for(int i = 0; i < sumgrid.x; i++) { sum += blocksum_host[i]; } timer.s_timer("GPU time (global memory):"); cout << " GPU result (global memory) = " << sum << endl; //显示GPU端结果 ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// timer.start_timer(); cudaMemcpy(test_d_cuda, test_d, sizeof(float) N, cudaMemcpyHostToDevice); //调用使用共享内存规约的核函数，sumblock.x为每个线程块包含的线程数，sumblock.x sizeof(float)就是传入的共享内存字节数 cal_sum_ker << < sumgrid, sumblock, sumblock.x sizeof(float) >> > (test_d_cuda, blocksum_cuda); cudaMemcpy(blocksum_host, blocksum_cuda, sizeof(float) sumgrid.x, cudaMemcpyDeviceToHost); sum = 0.0; for(int i = 0; i < sumgrid.x; i++) { sum += blocksum_host[i]; } timer.s_timer("GPU time (shared memory):"); cout << " GPU result (shared memory) = " << sum << endl; //显示GPU端结果 cudaFree(test_d_cuda); cudaFree(blocksum_cuda); free(blocksum_host); free(test_d);}

运行结果如下，可以看到使用共享内存之后，耗时减少了，这是因为共享内存的读写效率比全局内存高。

内存能用

电脑怎样才能用共享gpu内存,笔记本共享gpu内存，转载请注明出处。