▶ 動態並行。算法
● 動態並行直接從 GPU 上建立工做，能夠減小主機和設備間數據傳輸，在設備線程中調整配置。有數據依賴的並行工做能夠在內核運行時生成，並利用 GPU 的硬件調度和負載均衡。動態並行要求算法和程序要提早改進，消除遞歸、不規則的循環、結構或其餘不適合並行的狀況。api
● 動態並行的經典圖緩存
● 主機中 Runtime API 提供了跟蹤運行核、流與事件的函數，對主機進程中的全部線程來講 CUDA 對象都是可共享的，可是主機調用的各核函數之間是相互獨立的，CUDA 對象不能共享（重疊讀寫）。一樣的狀況也存在於設備中建立子內核函數的時候。併發
● 父線程中的 Runtime API 操做對於該線程所在的線程塊是可見的。意思就是在同一個線程塊中能夠由任意一個線程來調用子內核、調整流和事件，其做用等價。負載均衡
● 動態並行隱式的完成了父線程和子線程之間的同步，要求全部的子線程都結束後父線程才能結束。若 「父線程所在的線程塊中全部的線程」 都在 「子線程結束前」 所有結束了，則子線程隱式的強制結束。異步
● 父線程格與子線程格共享全局內存和常量內存，可是共享內存和局部內存私有。函數
● 父線程與子線程在兩個時間節點上共享的全局內存具備一致性：父線程格調用子線程格的時候；子線程格計算完成後，在父線程格中調用同步函數的時候。意思就是，父線程格先對全局內存進行操做，而後調用子線程格，則這些操做對子線程格來講都是可見的；子線程格對全局內存進行操做，而後在父線程格中進行同步，則這些操做對父線程格來講都是可見的。ui
● 動態並行與父 - 子內核之間的全局內存同步的代碼舉例。this
1 // 子內核 2 __global__ void child_launch(int *data) 3 { 4 data[threadIdx.x] = data[threadIdx.x] + 1; 5 } 6 7 // 父內核 8 __global__ void parent_launch(int *data) 9 { 10 data[threadIdx.x] = threadIdx.x; 11 12 __syncthreads(); // 同步全部父線程對全局內存的讀寫 13 14 if (threadIdx.x == 0)// 使用一個線程來啓動子內核 15 { 16 child_launch << < 1, 256 >> >(data); 17 // 調用子內核時，隱式保證了父線程對全局內存讀寫（data[0] = 0）對子內核可見， 18 // 但不能保證父內核中其餘線程的全局內存讀寫可見（由於調用子線程時父內核中其餘線程的全局內存讀寫不必定都完成了） 19 // 這裏多虧調用子內核以前使用了 __syncthreads();，保證父內核中全部線程的全局內存讀寫在調用子內核以前都已經完成，保證了子內核可見 20 cudaDeviceSynchronize();// 退出子內核時使用同步，保證父線程對子內核的全局內存讀寫可見 21 } 22 __syncthreads();// 父內核同步，保證父內核中全部線程對子內核的全局內存讀寫均可見 23 } 24 25 void host_launch(int *data) 26 { 27 parent_launch << < 1, 256 >> >(data); 28 }
● 零拷貝內存與全局內存具備相同的一致性，且不能在設備代碼中申請或釋放。spa
● 常量內存不能被設備修改，保證了設備之間高度的一致性。全部的常量內存都應該調用核函數以前由主機讀寫完成，調用核函數時常量內存就被自動繼承。主機與設備之間、設備與設備之間常量內存的指針能夠平凡傳遞。
● 共享內存被線程塊私有，局部內存被線程私有，二者均不能在父內核與子內核之間共享。把共享內存指針或局部內存指針傳遞給子內核的時候編譯器會發出警告；能夠使用函數 __isGlobal() 來檢測一個指針是否指向全局內存，防止將一個共享內存指針或局部內存指針傳遞給子內核；
1 // sm_20_intrinsics.h 2 // "ptr" 指向全局內空間則返回 1；指向共享、局部、常量內存空間則返回 0 3 __SM_20_INTRINSICS_DECL__ unsigned int __isGlobal(const void *ptr) 4 { 5 unsigned int ret; 6 asm volatile ("{ \n\t" 7 " .reg .pred p; \n\t" 8 " isspacep.global p, %1; \n\t" 9 " selp.u32 %0, 1, 0, p; \n\t" 10 #if (defined(_MSC_VER) && defined(_WIN64)) || defined(__LP64__) || defined(__CUDACC_RTC__) 11 "} \n\t" : "=r"(ret) : "l"(ptr)); 12 #else 13 "} \n\t" : "=r"(ret) : "r"(ptr)); 14 #endif 15 16 return ret; 17 }
● 須要在父內核與子內核之間傳遞的內存能夠在全局做用域裏顯式的聲明（__device__ int array[256];），防止在父內核中將局部內存地址傳入子內核。
● 紋理內存存與全局內存具備相同的一致性，注意使用同步函數來得到父內核與子內核之間的一致。
● 父內核啓用子內核的過程是異步的，這與主機啓用內核的狀況相同。
● 子內核集成父內核的配置參數，即函數 cudaDeviceGetCacheConfig() 和 cudaDeviceGetLimit() 調整的緩存、共享內存及其餘參數。從主機中調用內核時使用的設置優先於全局默認設置，且不能在設備中調用子內核的時候更改這些設置。
● 流能夠在指定的線程塊中的任意線程使用，但流句柄不能在線程塊之間、父子內核之間傳遞。不一樣的流中啓動的內核能夠並行運行，但不保證併發，這點在主機和設備之間、設備父子內核之間局成立。
● 主機端 NULL 流的跨流屏障語義對設備穿件的流不適用。設備內的流不能使用函數 cudaStreamCreate() 來建立，而是要使用函數 cudaStreamCreateWithFlags() 來建立，並傳入標誌 cudaStreamNonBlocking。
● 設備流中的同步不能使用函數 cudaStreamSynchronize() 或 cudaStreamQuery()，而要使用函數cudaDeviceSynchronize()。
● 主機端的 NULL 流隱式地添加了標誌 cudaStreamNonBlocking， NULL 流中啓動的的內核不會依賴於其餘流中掛起的工做（work launched into the NULL stream will not insert an implicit dependency on pending work in any other streams）。
● 自內核僅支持 CUDA 事件的流內同步功能，意思就是函數 cudaStreamWaitEvent()是可用的，但函數 cudaEventSynchronize()，cudaEventElapsedTime()，cudaEventQuery() 不可用。並且在子線程塊中建立事件時，還須要向函數 cudaEventCreateWithFlags() 傳入標誌 cudaEventDisableTiming。
● 事件在建立它的線程塊中的全部線程之間共享，不能傳遞給另外一個內核或線程塊。事件句柄不能保證在塊之間惟一，因此使用事件句柄前要先建立。
● (?) It is up to the program to perform sufficient additional inter-thread synchronization, for example via a call to __syncthreads(), if the calling thread is intended to synchronize with child grids invoked from other threads.
● (?) The cudaDeviceSynchronize() function does not imply intra-block synchronization. In particular, without explicit synchronization via a __syncthreads() directive the calling thread can make no assumptions about what work has been launched by any thread other than itself. For example if multiple threads within a block are each launching work and synchronization is desired for all this work at once (perhaps because of event-based dependencies), it is up to the program to guarantee that this work is submitted by all threads before calling cudaDeviceSynchronize().
● (?) Because the implementation is permitted to synchronize on launches from any thread in the block, it is quite possible that simultaneous calls to cudaDeviceSynchronize() by multiple threads will drain all work in the first call and then have no effect for the later calls.
● 一個內核只能在一個設備上啓用，父內核中不能使用函數 cudaSetDevice() 或 cudaGetDevicePropertites()，但能夠使用函數 cudaDeviceGetAttribute() 來訪問其餘設備的屬性。
● 在文件做用域內聲明的 __device__ 和 __constant__ 變量能夠被全部內核進行讀寫或讀。
● 只能在主機端建立或銷燬紋理或表面對象，設備中不能。只有頂層內核（直接被主機調用的內核）中才能使用紋理和表面（The device runtime does not support legacy module-scope textures and surfaces within a kernel launched from the device）。
● 主機與內核之間、設備中父子內核之間都可使用靜態或動態的方法調用共享內存，可是數據傳遞須要藉助全局內存來實現。全局內存地址能夠直接用 & 算符來獲取。
● 內核中常量內存能夠直接引用，且不能更改其內容，因此內核中不支持函數 cudaMemcpyToSymbol() 或 cudaGetSymbolAddress()。
● 內核中也能夠使用函數 cudaGetLastError() 來捕獲調用內核的錯誤，注意同時啓用多個內核的時候可能會發生多個錯誤，可是該函數返回值中只保存了最後一個。
● 調用內核的 <<< >>> 算符實際上在 PTX 中被解釋爲函數 cudaGetParameterBuffer() 和函數 cudaLaunchDevice()，（在 cuda_cevice_runtime_api.h中的定義稍有不一樣）。
1 extern __device__ cudaError_t cudaGetParameterBuffer(void **params); 2 extern __device__ cudaError_t cudaLaunchDevice(void *kernel, void *params, dim3 gridDim, dim3 blockDim, unsigned int sharedMemSize = 0, cudaStream_t stream = 0);
● 給出了全部能在設備代碼中使用的 Runtime API 函數。
● PTX 階段內核調用的低層次實現，詳細說明函數 cudaGetParameterBuffer() 和函數 cudaLaunchDevice() 的細節。
■ PTX 階段 cudaLaunchDevice() 的兩種實現。
1 // .address_size == 64 2 .extern.func(.param.b32 func_retval0) cudaLaunchDevice 3 ( 4 .param.b64 func, 5 .param.b64 parameterBuffer, 6 .param.align 4.b8 gridDimension[12], 7 .param.align 4.b8 blockDimension[12], 8 .param.b32 sharedMemSize, 9 .param.b64 stream 10 ); 11 12 // .address_size == 32 13 .extern.func(.param.b32 func_retval0) cudaLaunchDevice 14 ( 15 .param.b32 func, 16 .param.b32 parameterBuffer, 17 .param.align 4.b8 gridDimension[12], 18 .param.align 4.b8 blockDimension[12], 19 .param.b32 sharedMemSize, 20 .param.b32 stream 21 );
■ PTX 階段 cudaGetParameterBuffer() 的兩種實現。
1 // .address_size == 64 2 .extern.func(.param.b64 func_retval0) cudaGetParameterBuffer 3 ( 4 .param.b64 alignment, 5 .param.b64 size 6 ); 7 8 // .address_size == 32 9 .extern.func(.param.b32 func_retval0) cudaGetParameterBuffer 10 ( 11 .param.b32 alignment, 12 .param.b32 size 13 );
■ 使用的兩個函數的聲明。動態並行中必需要有頭文件 cuda_device_runtime_api.h，不過其已經在 Runtime 環境中被包含了（cuda_runtime.h）。
1 extern "C" __device__ cudaError_t cudaLaunchDevice(void *func, void *parameterBuffer, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize, cudaStream_t stream); 2 extern "C" __device__ void *cudaGetParameterBuffer(size_t alignment, size_t size); 3 4 // cuda_device_runtime_api.h 5 static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaLaunchDevice(void *func, void *parameterBuffer, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize, cudaStream_t stream) 6 { 7 return cudaLaunchDevice_ptsz(func, parameterBuffer, gridDimension, blockDimension, sharedMemSize, stream); 8 } 9 static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaLaunchDeviceV2(void *parameterBuffer, cudaStream_t stream) 10 { 11 return cudaLaunchDeviceV2_ptsz(parameterBuffer, stream); 12 } 13 extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaLaunchDevice(void *func, void *parameterBuffer, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize, cudaStream_t stream); 14 extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaLaunchDeviceV2(void *parameterBuffer, cudaStream_t stream); 15 16 extern __device__ __cudart_builtin__ void * CUDARTAPI cudaGetParameterBuffer(size_t alignment, size_t size);
■ 函數 cudaGetParameterBuffer() 的第一個參數是數據緩衝區的對齊值，默認爲 64 Byte，以保證各類類型的數據都能容納。(?) Parameter reordering in the parameter buffer is prohibited, and each individual parameter placed in the parameter buffer is required to be aligned. That is, each parameter must be placed at the nth byte in the parameter buffer, where n is the smallest multiple of the parameter size that is greater than the offset of the last byte taken by the preceding parameter. The maximum size of the parameter buffer is 4KB.
● 編譯使用動態並行的程序須要顯式連接庫文件（Windows：cudadevrt.lib，Linux MacOS：libcudadevrt.a）。
1 # 直接編譯和鏈接 2 $ nvcc hello_world.cu -o hello.exe -arch=sm_35 -rdc=true -lcudadevrt 3 4 # 先編譯後鏈接 5 $ nvcc hello_world.cu -o hello_world.o -arch=sm_35 -dc 6 $ nvcc hello_world.o -o hello.exe -arch=sm_35 -rdc=true -lcudadevrt