我正在尝试使用CUDA开发一个小程序,但由于它是SLOW,我做了一些测试并用Google搜索了一下.我发现虽然单个变量默认存储在本地线程内存中,但数组通常不存在.我想这就是为什么它花了这么多时间来执行.现在我想知道:因为本地线程内存至少应该是16KB,因为我的数组就像52个字符长,有没有办法(语法请:))将它们存储在本地内存中？
不应该是这样的:
__global__ my_kernel(int a)
{
__local__ unsigned char p[50];
}
我只发现本地内存比寄存器内存慢,每个线程两种类型.
共享内存应该很快,但它比[线程]的本地内存更快？
我想要做的是一种中值滤波器,但具有给定的百分位数而不是中位数.因此,我需要获取列表的块,对它们进行排序,然后选择一个合适的列表.但我无法开始对共享内存列表进行排序或出现问题.只要复制到本地内存,我会失去很多性能吗？
我在NVIDIA文档(http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#features-and-technical-specifications,table#12)中读到每个线程的本地内存量我的GPU是512 Ko(GTX 580,计算能力2.0).
我尝试用CUDA 6.5检查Linux上的这个限制是不成功的.
这是我使用的代码(它的唯一目的是测试本地内存限制,它不会进行任何有用的计算):
#include <iostream>
#include <stdio.h>
#define MEMSIZE 65000 // 65000 -> out of memory, 60000 -> ok
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=false)
{
if (code != cudaSuccess)
{
fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
if( abort )
exit(code);
}
}
inline void gpuCheckKernelExecutionError( const char *file, int line)
{
gpuAssert( cudaPeekAtLastError(), file, line);
gpuAssert( cudaDeviceSynchronize(), file, line);
}
__global__ void kernel_test_private(char *output)
{
int c = …我读了一些引用本地内存的CUDA文档.(这主要是早期文档.)设备属性报告本地mem大小(每个线程)."本地"记忆是什么意思？什么是'本地'记忆？"本地"记忆在哪里？如何访问"本地"内存？这是__device__记忆,不是吗？
设备属性还报告:全局,共享和常量mem大小.这些陈述是否正确:
全局记忆是__device__记忆.它具有网格范围和网格(内核)的生命周期.
常量内存是__device__ __constant__内存.它具有网格范围和网格(内核)的生命周期.
共享内存是__device__ __shared__内存.它具有单个块范围和该块(线程)的生命周期.
我在想共享内存是SM内存.即只有那个SM才能直接访问的内存.资源相当有限.SM是不是一次分配了一堆块？这是否意味着SM可以交错执行不同的块(或不是)？即运行阻止*A*线程直到它们停止.然后运行block*B*线程直到它们停止.然后再次换回阻止*A*线程.或者SM是否为块*A*运行一组线程,直到它们停止.然后交换另一组块*A*线程.此交换继续,直到块*A*用尽.然后,只有这样才能在块*B*上开始工作.因为共享记忆,我问.如果单个SM从2个不同的块交换代码,那么SM如何快速交换/分出共享内存块？(我认为后来的senerio是真的,并且没有交换共享内存空间.块*A*运行直到完成,然后块*B*开始执行.注意:块*A*可能是不同的内核比块*B*.)
我在CUDA上写了一个简单的函数.它将图像调整为双倍缩放.对于1920*1080的图像,此功能需要~20ms才能完成.我尝试了一些不同的方法来优化该功能.而我发现可能是本地记忆是关键原因.
我尝试了三种不同的方法来获取图像.
他们都不能给我带来一点改善.
然后我使用nvvp找出原因.在上述所有三个条件下,本地内存开销约为95%.
所以我转向我的代码,了解nvcc如何使用内存.然后我发现一个简单的函数就像这样:
__global__ void performDoubleImage(float* outData, size_t step, const int cols, const int rows)
{
int x = threadIdx.x + blockIdx.x * blockDim.x;
if (x >= cols)
return;
int y = threadIdx.y + blockIdx.y * blockDim.y;
if (y >= rows)
return;
((float*)((size_t)outData+y*step))[x] = tex2D(texRef, x/2, y/2);
}
需要80字节的堆栈帧(它们在本地内存中).
而另一个功能是这样的:
__global__ void performFinalDoubleImage(const PtrStepSz<float> in, PtrStepSz<float> out)
{
out(out.rows-1, out.cols-1) = out(in.rows-1, in.cols-1);
}
还需要88字节的堆栈帧.
问题是,为什么我的函数在这个简单的任务中使用了如此多的本地内存和寄存器？为什么OpenCV中的函数可以通过不使用本地内存来执行相同的功能(这是由nvvp测试,本地内存负载是ZERO)？
我的代码是在调试模式下编译的.我的卡是GT650(192 SP/SM,2 SM).