小编Dou*_*oug的帖子

后续问答来自：CUDA：从内核调用__device__函数

我正在尝试加快排序操作。简化的伪版本如下：

// some costly swap operation
__device__ swap(float* ptrA, float* ptrB){
  float saveData;         // swap some 
  saveData= *Adata;       //   big complex
  *Adata= *Bdata          //     data chunk
  *Bdata= saveData;
}

// a rather simple sort operation
__global__ sort(float data[]){
  for (i=0; i<limit: i++){
  find left swap point
  find right swap point
  swap<<<1,1>>>(left, right);
  }
}

（注意：这个简单的版本没有在块中显示还原技术。）这个想法是很容易（快速）识别交换点。交换操作成本高（缓慢）。因此，使用一个块来查找/标识交换点。使用其他块执行交换操作。即并行进行实际交换。这听起来像一个不错的计划。但是，如果编译器内联设备调用，则不会发生并行交换。有没有办法告诉编译器不要内联设备调用？

跟进:CUDA:停止所有其他线程

如果发生"坏情况",我正在寻找退出内核的方法.编程手册说NVCC不支持异常处理.我想知道是否有用户定义的cud​​a-error-code.换句话说,如果发生"坏",则以此用户错误代码终止.我怀疑是否有一个,所以我的另一个想法是造成一个.

如果"坏"发生,则除以零.但我不确定一个线程是否进行零除,是否足以使整个内核崩溃,或只是崩溃？

有没有更好的方法来终止内核？

我想在CUDA中实现这个原子函数:

__device__ float lowest;   // global var
__device__ int  lowIdx;    // global var
float realNum;   // thread reg var
int index;       // thread reg var

if(realNum < lowest) {
 lowest= realNum;  // the new lowest
 lowIdx= index;    // update the 'low' index
}

我不相信我能用任何原子功能做到这一点.我需要锁定几个全局内存loc以获取一些指令.我可以用PTXAS(汇编)代码实现这个吗？

我需要一些帮助来追踪push :: fill给我的编译错误.

代码没有错:

line 9   #include <thrust/device_vector.h>                // needed for other thrust stuff
line 10  #include <thrust/fill.h>                        // not needed (same err w/ or w/o)
.
.
line 389 thrust::device_vector<int> junk(20);           // any ol array
line 390 thrust::fill(junk.begin(), junk.end(), (int)0); // the problem line

如果我注释掉'fill'行,它将编译正常.我使用的任何(raw_pointer,device_vector,..)thrust :: fill命令都会生成此错误.其他THRUST cmds编译正常.错误是几页长.我不确定错误是什么.'for_each'被提到了很多.带有'compute_capability'的最后一个没有意义,因为它不在我的任何代码中.我尝试添加一堆'#include'文件,但它没有任何帮助.

任何想法,将不胜感激.

In file included from /usr/local/cuda/include/thrust/detail/backend/cuda/for_each.h:54:0,
                 from /usr/local/cuda/include/thrust/detail/backend/for_each.inl:23,
                 from /usr/local/cuda/include/thrust/detail/backend/for_each.h:55,
                 from /usr/local/cuda/include/thrust/detail/for_each.inl:22,
                 from /usr/local/cuda/include/thrust/for_each.h:91,
                 from /usr/local/cuda/include/thrust/detail/swap_ranges.inl:25,
                 from /usr/local/cuda/include/thrust/detail/swap.inl:33,
                 from /usr/local/cuda/include/thrust/swap.h:127,
                 from /usr/local/cuda/include/thrust/detail/contiguous_storage.inl:20,
                 from /usr/local/cuda/include/thrust/detail/contiguous_storage.h:92,
                 from /usr/local/cuda/include/thrust/detail/uninitialized_array.h:26,
                 from /usr/local/cuda/include/thrust/detail/backend/cuda/copy_cross_space.inl:20, …

我读了一些引用本地内存的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*.)

我需要一些帮助来理解Ron Farber代码的行为:http://www.drdobbs.com/parallel/cuda-supercomputing-for-the-masses-part/208801731？ pgno = 2

我不明白共享内存的使用如何在非共享内存版本上提供更快的性能.即如果我添加一些索引计算步骤并使用添加另一个Rd/Wr循环来访问共享内存,这怎么能比单独使用全局内存更快？在任何一种情况下,相同的数字或Rd/Wr循环访问全局mem.每个内核实例只能访问一次数据.数据仍然使用全局内存进/出.内核实例的数量是相同的.寄存器计数看起来是一样的.如何添加更多处理步骤使其更快.(我们不会减去任何流程步骤.)基本上我们正在做更多的工作,而且它的工作更快.

共享内存访问速度比全局快得多,但它不是零(或负数).我错过了什么？

'慢'代码:

__global__ void reverseArrayBlock(int *d_out, int *d_in) {
int inOffset  = blockDim.x * blockIdx.x;
int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);
int in  = inOffset + threadIdx.x;
int out = outOffset + (blockDim.x - 1 - threadIdx.x);
d_out[out] = d_in[in];
}

"快速"代码:

__global__ void reverseArrayBlock(int *d_out, int *d_in) {
extern __shared__ int s_data[];

int inOffset  = blockDim.x * blockIdx.x;
int in  = inOffset + threadIdx.x;

// Load …