在CUDA中,存在warp的概念,其被定义为可以在单个处理元件内同时执行相同指令的最大线程数.对于NVIDIA,目前市场上所有卡的经线尺寸均为32.
在ATI卡中,有一个类似的概念,但在这种情况下的术语是波前.经过一番狩猎后,我发现我所拥有的ATI卡的波前大小为64.
我的问题是,在运行时为OpenCL查询此SIMD宽度我该怎么办？
主机是否等待设备完成执行？例如,程序具有如下结构
// cpu code segment
// data transfer from host to device
QUESTION - WILL CPU WAIT FOR DEVICE TO FINISH TRANSFER? IF NO, IS IT POSSIBLE? IF YES, HOW?
// kernel launch
QUESTION - WILL CPU WAIT FOR DEVICE TO LET IT FINISH KERNEL EXECUTION (CONSIDERING KERNEL EXECUTION WILL TAKE NOTABLE TIME say-5 sec)? IF NO, IS IT POSSIBLE? IF YES, HOW?
// data transfer from device to host
// program terminates after printing some information
我正在使用Windows 7 64位SP1上的CUDA Toolkit 4.0和Visual Studio 2010 Professional为GTX 580开发CUDA应用程序.我的程序比典型的CUDA程序更耗费内存,我试图为每个CUDA块分配尽可能多的共享内存.但是,每次尝试为每个块使用超过32K的共享内存时,程序都会崩溃.
通过阅读官方CUDA文档,我了解到CUDA设备上每个SM有48KB的片上存储器,其计算能力为2.0或更高,而片上存储器在L1缓存和共享存储器之间分配:
相同的片上存储器用于L1和共享存储器,并且可以为每个内核调用配置多少L1和共享存储器(第F.4.1节) http://developer.download.nvidia.com /compute/DevZone/docs/html/C/doc/Fermi_Tuning_Guide.pdf
这让我怀疑在我的程序运行时只有32KB的单内存被分配为共享内存.因此我的问题是:是否可以将所有48KB的片上内存用作共享内存？
我尝试了我能想到的一切.我为nvcc指定了选项--ptxas-options =" - v -dlcm = cg",我在程序中调用了cudaDeviceSetCacheConfig()和cudaFuncSetCacheConfig(),但没有一个解决了这个问题.我甚至确保没有寄存器溢出,并且我没有意外地使用本地内存:
1> 24 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
1> ptxas info : Used 63 registers, 40000+0 bytes smem, 52 bytes cmem[0], 2540 bytes cmem[2], 8 bytes cmem[14], 72 bytes cmem[16]
虽然我可以使用32KB的共享内存,这已经给了我巨大的性能提升,但我宁愿充分利用所有快速的片上内存.任何帮助深表感谢.
更新:我在程序崩溃时启动了640个线程.512给了我比256更好的性能,所以我试图进一步增加线程数.
给出以下代码片段,使用推力生成一种带CUDA的代码字典(CUDA的C++模板库):
thrust::device_vector<float> dCodes(codes->begin(), codes->end());
thrust::device_vector<int> dCounts(counts->begin(), counts->end());
thrust::device_vector<int> newCounts(counts->size());
for (int i = 0; i < dCodes.size(); i++) {
float code = dCodes[i];
int count = thrust::count(dCodes.begin(), dCodes.end(), code);
newCounts[i] = dCounts[i] + count;
//Had we already a count in one of the last runs?
if (dCounts[i] > 0) {
newCounts[i]--;
}
//Remove
thrust::detail::normal_iterator<thrust::device_ptr<float> > newEnd = thrust::remove(dCodes.begin()+i+1, dCodes.end(), code);
int dist = thrust::distance(dCodes.begin(), newEnd);
dCodes.resize(dist);
newCounts.resize(dist);
}
codes->resize(dCodes.size());
counts->resize(newCounts.size());
thrust::copy(dCodes.begin(), dCodes.end(), codes->begin());
thrust::copy(newCounts.begin(), newCounts.end(), counts->begin());
问题是,通过使用CUDA视觉分析器,我注意到4个字节的多个副本.IMO这是由生成的
我在C#中使用ManagedCuda,我有一个问题,我无法找到答案......也许你可以帮助我.我在C++和CUDA中读到你可以声明一个变量(这是一个数组),如:
__constant__ double myVar[X]; (这是为了容纳一个X元素数组)
然后使用它来设置主机代码的值:
cudaMemcpyToSymbol(myVar, &arrayFromHost[0], sizeof(arrayFromHost) * numElements,
size_t(0),cudaMemcpyHostToDevice);
所以现在你可以使用类似的东西:
__global__ void myFunction(double *res)
{
*res = myVar[0] + 2.5;
}
使用myVar从主机设置的值...
但是在ManagedCuda中,我似乎无法做到这一点......我怎么能这样做？
(或__device__变量......我不知道......它将是一个变量,它将在第一次运行时接收一个数组(具有未知数量的元素),从那时起,该函数将引用它值,但该变量永远不会改变)
现在我只声明一个CudaDeviceVariable,我不再触摸它,但在我的内核上我总是要发送DevicePointer,我觉得这在阅读时更难理解......
现在它看起来像这样:
myKernel.Run(staticData.DevicePointer, moreData.DevicePointer,
evenMoreData.DevicePointer, numberOfElementsWhichNeverChange,
moreStaticData.DevicePointer, myResults.DevicePointer)
我想跳过具有永不改变的数据的3个参数,并将其设置在另一个函数中,setData.Run(numElements, staticData, moreStaticData);
并在我的*.cu文件中的其他函数中使用常量或设备变量.
有人可以对此发表评论,
我想做一个矢量点积.我的浮点矢量是[2080:2131]和[2112:2163],每个都包含52个元素.
a[52] = {2080 2081 2082 ... ... 2129 2130 2131};
b[52] = {2112 2113 2114 ... ... 2161 2162 2163};
for (int i = 0; i < 52; i++)
{
sum += a[i]*b[i];
}
我的内核的全长(52元素)的结果总和为234038032,而matlab的结果为234038038.对于产品的1到9元素总和,我的内核结果与matlab结果一致.对于10个元素的总和,它偏离1并逐渐增加.结果是可重复的.我检查了所有元素,发现没有问题.