排序和并发：
设备运行时的内核启动顺序遵循CUDA Stream排序语义。 在一个线程块内，所有内核启动到同一个流中都会按顺序执行。 由于同一线程块中的多个线程启动到同一个流中，流内的排序取决于该块内的线程调度，这可以通过同步原语来控制，例如__syncthreads()
请注意，因为流由线程块内的所有线程共享，所以隐式NULL流也被共享。 如果线程块中的多个线程启动到隐式流中，那么这些启动将按顺序执行。 如果需要并发性，则应使用显式命名流。
动态并行可以使并发在程序中更容易表达; 但是，设备运行时不会在CUDA执行模型中引入新的并发保证。 不能保证设备上任何数量的不同线程块之间的并发执行。
缺乏并发性保证扩展到父线程块及其子网格。 当一个父线程块启动一个子网格时，子节点不能保证开始执行，直到父节点线程块达到显式同步点（例如cudaDeviceSynchronize（））。
虽然并发性通常很容易实现，但它可能会随着设备配置，应用程序工作负载和运行时间调度而变化。 因此依赖不同线程块之间的任何并发是不安全的。
设备管理:
设备运行时没有多GPU支持; 设备运行时间只能在当前正在执行的设备上运行。 但是，允许查询系统中任何支持CUDA的设备的属性。
内存模型：
父网格和子网格共享相同的全局和常量内存存储，但具有不同的本地和共享内存。
一致性：
全局内存：
父子网格对全球记忆具有连贯性，并且父子之间的一致性保证较弱。 当其内存视图与父线程完全一致时，子网格的执行有两点：当父网格调用子网格时，以及父网格中的同步API调用发出信号时子网格完成时线。
在子网格调用之前，父线程中的所有全局内存操作对子网格都是可见的。子网格的所有内存操作在父网格已完成子网格同步后对父级可见。
在以下示例中，执行child_launch的子网格只能保证在子网格启动之前查看对数据的修改。 由于父级的线程0正在执行启动，所以子级将与父级的线程0所见的内存一致。 由于第一个__syncthreads（）调用，子网格将看到数据[0] = 0，数据[1] = 1，...，数据[255] = 255（不带__syncthreads（）调用，只有数据[0] 将保证被子网格看到）。 当子网格返回时，线程0将保证看到它的子网格中的线程所做的修改。 这些修改仅在第二个__syncthreads（）调用后才可用于父网格的其他线程：
__global__ void child_launch(int *data) {
data[threadIdx.x] = data[threadIdx.x] + 1;
}
__global__ void parent_launch(int *data) {
data[threadIdx.x] = threadIdx.x;
__syncthreads();
if (threadIdx.x == 0) {
child_launch << < 1, 256 >> >(data);
cudaDeviceSynchronize();
}
__syncthreads();
}
void host_launch(int *data) {
parent_launch << < 1, 256 >> >(data);
}零拷贝内存:
零拷贝系统内存对全局内存具有相同的一致性和一致性保证，并遵循上面详述的语义。 内核可能不会分配或释放零拷贝内存，但可能使用指向从主机程序传入的零拷贝的指针。
常量内存：
常量是不可改变的，并且可能不会从设备中修改，即使在父级和子级启动之间也是如此。 也就是说，必须在启动之前从主机设置所有__constant__变量的值。 所有子内核都从它们各自的父级中自动继承常量内存。
从内核线程中获取常量内存对象的地址与所有CUDA程序具有相同的语义，并且自然支持将该指针从父对象传递给子对象或从子对象传递给父对象。
共享和本地内存：
共享和本地内存分别对于线程块或线程是私有的，并且在父级和子级之间不可见或不一致。 如果这些位置中的某个位置中的对象被引用到其所属的范围之外并且可能导致错误，则行为未定义。
NVIDIA编译器会尝试警告它是否可以检测到指向本地或共享内存的指针作为内核启动参数传递。 在运行时，程序员可以使用__isGlobal（）内部函数确定指针是否引用全局内存，因此可以安全地传递给子启动。
请注意，对cudaMemcpy Async（）或cudaMemset Async（）的调用可能会在设备上调用新的子内核以保留流语义。 因此，将共享或本地内存指针传递给这些API是非法的，并且会返回错误。
本地内存:
本地内存是执行线程的私有存储空间，并且在该线程之外不可见。 启动子内核时，将指针传递给本地内存作为启动参数是非法的。 取消引用来自子级的这种本地存储器指针的结果将是未定义的。
例如，以下是非法的，如果x_array被child_launch访问，则具有未定义的行为：
int x_array[10]; // Creates x_array in parent's local memory
child_launch<<< 1, 1 >>>(x_array);程序员有时很难意识到编译器将变量放入本地内存的时间。 作为一般规则，传递给子内核的所有存储应该从全局内存堆中明确分配，或者使用cudaMalloc（），new（）或者在全局范围内声明__device__存储。 例如：
// Correct - "value" is global storage
__device__ int value;
__device__ void x() {
value = 5;
child << < 1, 1 >> >(&value);
}
// Invalid - "value" is local storage
__device__ void y() {
int value = 5;
child << < 1, 1 >> >(&value);
}纹理内存：
写入纹理映射的全局内存区域与纹理访问不相关。 在调用子网格和子网格完成时，强制执行纹理内存的一致性。 这意味着在子内核启动之前写入内存会反映在子级的纹理内存访问中。 同样，由孩子写入内存将反映在父级的纹理内存访问中，但仅在父级和子级完成后进行同步之后。 父级和子级同时访问可能会导致数据不一致。
