小编Ale*_*lex的帖子

使用 CUDA SDK 5.5，我可以用来复制数据：

• 来自主机：cudaMemcpy();如果内存固定则使用 GPU-DMA
• 来自主机：memcpy();或者cudaMemcpy();如果内存未固定则使用 CPU 内核
• 来自 gpu：for() { dst[i] = src[i]; }或memcpy();使用 GPU 核心
• 从 GPU：使用 GPU-DMA ???

如何在GPU-CUDA 代码的内核函数中使用 GPU-DMA 来复制数据？

Linux中的所有驱动程序都在相同的上下文(内核空间的地址空间)中运行,或者每个驱动程序在不同的环境中运行(类似于不同进程在用户空间的不同地址空间中的工作方式)？

我想使用 .map 将设备内存（NIC）映射到内核空间内存区域ioremap_wc()。然后我想从内核空间到用户空间的重新映射内存区域，我可以使用2个功能是：vm_insert_page()与remap_pfn_range()

POSIXmmap(3)通常使用第二种：remap_pfn_range()

vm_insert_page()和之间有什么区别remap_pfn_range()，我什么时候需要使用vm_insert_page()而不是remap_pfn_range()？

正如我们从之前的回答中所知道的,它是否在处理器x86/x86_64中指示LFENCE？我们不能使用SFENCE而不是MFENCE顺序一致性.

这里的答案表明MFENCE= SFENCE+ LFENCE,即LFENCE没有我们不能提供顺序一致性的东西.

LFENCE 无法重新排序:

SFENCE
LFENCE
MOV reg, [addr]

- 到 - >

MOV reg, [addr]
SFENCE
LFENCE

例如重新排序MOV [addr], reg LFENCE- > LFENCE MOV [addr], reg由机制提供- 存储缓冲区,它重新排序存储 - 负载以提高性能,并且因为LFENCE它不会阻止它.并SFENCE 禁用此机制.

什么机制禁用LFENCE无法重新排序(x86没有机制 - Invalidate-Queue)？

并且只是在理论上或者在现实中重新排序SFENCE MOV reg, [addr]- > MOV reg, [addr] SFENCE可能吗？如果可能,实际上,什么机制,它是如何工作的？

超线程是否允许使用 L1 缓存在两个线程之间交换数据，这两个线程在单个物理内核上同时执行，但在两个虚拟内核中？

前提是两者都属于同一个进程，即在同一个地址空间中。

第 85 页 (2-55) -英特尔® 64 位和 IA-32 架构优化参考手册：http : //www.intel.com/content/dam/www/public/us/en/documents/manuals/64-ia- 32-architectures-optimization-manual.pdf

2.5.9 英特尔® 微架构代号 Nehalem 中的超线程技术支持

...

更深的缓冲和增强的资源共享/分区策略：

• HT 操作的复制资源：寄存器状态、重命名的返回堆栈缓冲区、大页面 ITLB。

• HT 操作的分区资源：加载缓冲区、存储缓冲区、重新排序缓冲区、小页面 ITLB 在两个逻辑处理器之间静态分配。

• HT 操作期间竞争性共享的资源：保留站、缓存层次结构、填充缓冲区、DTLB0 和 STLB。

• HT操作时交替：前端操作一般在两个逻辑处理器之间交替，以保证公平性。

• HT 未知资源：执行单元。

我需要使用VirtualAlloc/VirtualAllocEx？

一个例子,我发现的一个案例 - 如果我分配了4 GB的虚拟内存,那么如果我不使用所有这些,那么我不会花费物理内存,如果我调整我的数组,我不需要做新的分配旧数据并将其复制到新数组.

struct T_custom_allocator; // which using VirtualAllocEx()
std::vector<int, T_custom_allocator> vec;
vec.reserve(4*1024*1024*1024);  // allocated virtual memory (physical memory is not used)
vec.resize(16384); // allocated 16KB of physical memory
// ...
vec.resize(32768); // allocated 32KB of physical memory 
                   // (no need to copy of first 16 KB of data)

如果我使用标准分配器,我需要在调整大小时复制数据:

std::vector<int> vec;
vec.resize(16384); // allocated 16KB of physical memory
// ...
vec.resize(32768); // allocated 32KB of physical memory 
                   // and need to …

正如安东尼·威廉斯说:

some_atomic.load(std :: memory_order_acquire)只是简单地加载到一个简单的加载指令,some_atomic.store(std :: memory_order_release)直接转到一个简单的存储指令.

众所周知,在x86上,操作load()和store()内存屏障memory_order_consume, memory_order_acquire, memory_order_release, memory_order_acq_rel不需要处理器指令.

但在ARMv8我们知道,这里都是内存壁垒load()和store(): http://channel9.msdn.com/Shows/Going+Deep/Cpp-and-Beyond-2012-Herb-Sutter-atomic-Weapons-1-of- 2 http://channel9.msdn.com/Shows/Going+Deep/Cpp-and-Beyond-2012-Herb-Sutter-atomic-Weapons-2-of-2

关于CPU的不同架构:http://g.oswego.edu/dl/jmm/cookbook.html

接下来,但对于x86上的CAS操作,具有不同内存屏障的这两行在反汇编代码(MSVS2012 x86_64)中是相同的:

    a.compare_exchange_weak(temp, 4, std::memory_order_seq_cst, std::memory_order_seq_cst);
000000013FE71A2D  mov         ebx,dword ptr [temp]  
000000013FE71A31  mov         eax,ebx  
000000013FE71A33  mov         ecx,4  
000000013FE71A38  lock cmpxchg dword ptr [temp],ecx  

    a.compare_exchange_weak(temp, 5, std::memory_order_relaxed, std::memory_order_relaxed);
000000013FE71A4D  mov         ecx,5  
000000013FE71A52  mov         eax,ebx  
000000013FE71A54  lock cmpxchg dword ptr [temp],ecx  

由GCC 4.8.1 …

我使用CUDA 6.5和4 x GPU开普勒.

我使用多线程,CUDA运行时API和从不同CPU线程访问CUDA上下文(通过使用OpenMP - 但它并不重要).

1. 当我打电话cudaDeviceSynchronize();时,它会等待内核只在最新调用选择的当前CUDA上下文中完成cudaSetDevice(),还是在所有CUDA上下文中完成？

2. 如果它将等待内核在所有CUDA上下文中完成,那么它将等待当前CPU线程中使用的所有CUDA上下文(例如CPU thread_0将等待GPU:0和1)或通常所有CUDA上下文(CPU) thread_0会等待GPU:0,1,2和3)？

以下代码:

// For using OpenMP requires to set:
// MSVS option: -Xcompiler "/openmp"
// GCC option: –Xcompiler –fopenmp
#include <omp.h>

int main() {

    // execute two threads with different: omp_get_thread_num() = 0 and 1
    #pragma omp parallel num_threads(2)
    {
        int omp_threadId = omp_get_thread_num();

        // CPU thread 0
        if(omp_threadId == 0) {

            cudaSetDevice(0);
            kernel_0<<<...>>>(...);
            cudaSetDevice(1);
            kernel_1<<<...>>>(...);

            cudaDeviceSynchronize(); // …

为什么我应该使用CUDA驱动程序API,在哪些情况下我不能使用CUDA Runtime API(比Driver API更方便)？

在CUDA 6.5的文档中写道:http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#ixzz3PIXMTktb

5.2.3.多处理器级别

...

• 8L用于计算能力3.x的设备,因为多处理器在一个时钟周期内每次发生一对指令,一次四个经线,如Compute Capability 3.x中所述.

这是否意味着GPU Kepler CC3.0处理器不仅是流水线架构,还是超标量？

1. 流水线 - 这两个序列并行执行(一次执行不同的操作):

   • LOAD [addr1] - > ADD - > STORE [addr1] - > NOP
   • NOP - > LOAD [addr2] - > ADD - > STORE [addr2]

2. 超标量 - 这两个序列并行执行(一次执行相同的操作):

   • LOAD [reg1] - > ADD - > STORE [reg1]
   • LOAD [reg2] - > ADD - > STORE [reg2]