Library/ManagedCuda/CudaKernel.cs
/// <summary>
/// Launches CUDA functions on multiple devices where thread blocks can cooperate and synchronize as they execute
/// <para/>
/// Invokes kernels as specified in the \p launchParamsList array where each element
/// of the array specifies all the parameters required to perform a single kernel launch.
/// These kernels can cooperate and synchronize as they execute. The size of the array is
/// specified by \p numDevices.
/// <para/>
/// No two kernels can be launched on the same device. All the devices targeted by this
/// multi-device launch must be identical. All devices must have a non-zero value for the
/// device attribute ::CU_DEVICE_ATTRIBUTE_COOPERATIVE_MULTI_DEVICE_LAUNCH.
/// <para/>
/// All kernels launched must be identical with respect to the compiled code. Note that
/// any __device__, __constant__ or __managed__ variables present in the module that owns
/// the kernel launched on each device, are independently instantiated on every device.
/// It is the application's responsiblity to ensure these variables are initialized and
/// used appropriately.
/// <para/>
/// The size of the grids as specified in blocks, the size of the blocks themselves
/// and the amount of shared memory used by each thread block must also match across
/// all launched kernels.
/// <para/>
/// The streams used to launch these kernels must have been created via either ::cuStreamCreate
/// or ::cuStreamCreateWithPriority. The NULL stream or ::CU_STREAM_LEGACY or ::CU_STREAM_PER_THREAD
/// cannot be used.
/// <para/>
/// The total number of blocks launched per kernel cannot exceed the maximum number of blocks
/// per multiprocessor as returned by ::cuOccupancyMaxActiveBlocksPerMultiprocessor (or
/// ::cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags) times the number of multiprocessors
/// as specified by the device attribute ::CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT. Since the
/// total number of blocks launched per device has to match across all devices, the maximum
/// number of blocks that can be launched per device will be limited by the device with the
/// least number of multiprocessors.
/// <para/>
/// The kernels cannot make use of CUDA dynamic parallelism.
/// <para/>
/// By default, the kernel won't begin execution on any GPU until all prior work in all the specified
/// streams has completed. This behavior can be overridden by specifying the flag
/// ::CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_PRE_LAUNCH_SYNC. When this flag is specified, each kernel
/// will only wait for prior work in the stream corresponding to that GPU to complete before it begins
/// execution.
/// <para/>
/// Similarly, by default, any subsequent work pushed in any of the specified streams will not begin
/// execution until the kernels on all GPUs have completed. This behavior can be overridden by specifying
/// the flag ::CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_POST_LAUNCH_SYNC. When this flag is specified,
/// any subsequent work pushed in any of the specified streams will only wait for the kernel launched
/// on the GPU corresponding to that stream to complete before it begins execution.
/// <para/>
/// Calling ::cuLaunchCooperativeKernelMultiDevice() sets persistent function state that is
/// the same as function state set through ::cuLaunchKernel API when called individually for each
/// element in \p launchParamsList.
/// <para/>
/// When kernels are launched via ::cuLaunchCooperativeKernelMultiDevice(), the previous
/// block shape, shared size and parameter info associated with each ::CUDA_LAUNCH_PARAMS::function
/// in \p launchParamsList is overwritten.
/// <para/>
/// Note that to use ::cuLaunchCooperativeKernelMultiDevice(), the kernels must either have
/// been compiled with toolchain version 3.2 or later so that it will
/// contain kernel parameter information, or have no kernel parameters.
/// If either of these conditions is not met, then ::cuLaunchCooperativeKernelMultiDevice() will
/// return ::CUDA_ERROR_INVALID_IMAGE.
/// </summary>
/// <param name="launchParameterList">List of launch parameters, one per device</param>
/// <param name="flags">Flags to control launch behavior</param>
Launches CUDA functions on multiple devices where thread blocks can cooperate and synchronize as they execute

Invokes kernels as specified in the \p launchParamsList array where each element
of the array specifies all the parameters required to perform a single kernel launch.
These kernels can cooperate and synchronize as they execute. The size of the array is
specified by \p numDevices.

No two kernels can be launched on the same device. All the devices targeted by this
multi-device launch must be identical. All devices must have a non-zero value for the
device attribute ::CU_DEVICE_ATTRIBUTE_COOPERATIVE_MULTI_DEVICE_LAUNCH.

All kernels launched must be identical with respect to the compiled code. Note that
any __device__, __constant__ or __managed__ variables present in the module that owns
the kernel launched on each device, are independently instantiated on every device.
It is the application's responsiblity to ensure these variables are initialized and
used appropriately.

The size of the grids as specified in blocks, the size of the blocks themselves
and the amount of shared memory used by each thread block must also match across
all launched kernels.

The streams used to launch these kernels must have been created via either ::cuStreamCreate
or ::cuStreamCreateWithPriority. The NULL stream or ::CU_STREAM_LEGACY or ::CU_STREAM_PER_THREAD
cannot be used.



The kernels cannot make use of CUDA dynamic parallelism.

By default, the kernel won't begin execution on any GPU until all prior work in all the specified
streams has completed. This behavior can be overridden by specifying the flag
::CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_PRE_LAUNCH_SYNC. When this flag is specified, each kernel
will only wait for prior work in the stream corresponding to that GPU to complete before it begins
execution.

Similarly, by default, any subsequent work pushed in any of the specified streams will not begin
execution until the kernels on all GPUs have completed. This behavior can be overridden by specifying
the flag ::CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_POST_LAUNCH_SYNC. When this flag is specified,
any subsequent work pushed in any of the specified streams will only wait for the kernel launched
on the GPU corresponding to that stream to complete before it begins execution.

Calling ::cuLaunchCooperativeKernelMultiDevice() sets persistent function state that is
the same as function state set through ::cuLaunchKernel API when called individually for each
element in \p launchParamsList.



Note that to use ::cuLaunchCooperativeKernelMultiDevice(), the kernels must either have
been compiled with toolchain version 3.2 or later so that it will
contain kernel parameter information, or have no kernel parameters.
If either of these conditions is not met, then ::cuLaunchCooperativeKernelMultiDevice() will
return ::CUDA_ERROR_INVALID_IMAGE.
Launches CUDA functions on multiple devices where thread blocks can cooperate and synchronize as they execute Invokes kernels as specified in the \ p launchParamsList array where each element of the array specifies all the parameters required to perform a single kernel launch . These kernels can cooperate and synchronize as they execute . The size of the array is specified by \ p numDevices . No two kernels can be launched on the same device . All the devices targeted by this multi - device launch must be identical . All devices must have a non - zero value for the device attribute :: CU_DEVICE_ATTRIBUTE_COOPERATIVE_MULTI_DEVICE_LAUNCH . All kernels launched must be identical with respect to the compiled code . Note that any __device__ __constant__ or __managed__ variables present in the module that owns the kernel launched on each device are independently instantiated on every device . It is the application ' s responsiblity to ensure these variables are initialized and used appropriately . The size of the grids as specified in blocks the size of the blocks themselves and the amount of shared memory used by each thread block must also match across all launched kernels . The streams used to launch these kernels must have been created via either :: cuStreamCreate or :: cuStreamCreateWithPriority . The NULL stream or :: CU_STREAM_LEGACY or :: CU_STREAM_PER_THREAD cannot be used . The kernels cannot make use of CUDA dynamic parallelism . By default the kernel won ' t begin execution on any GPU until all prior work in all the specified streams has completed . This behavior can be overridden by specifying the flag :: CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_PRE_LAUNCH_SYNC . When this flag is specified each kernel will only wait for prior work in the stream corresponding to that GPU to complete before it begins execution . Similarly by default any subsequent work pushed in any of the specified streams will not begin execution until the kernels on all GPUs have completed . This behavior can be overridden by specifying the flag :: CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_POST_LAUNCH_SYNC . When this flag is specified any subsequent work pushed in any of the specified streams will only wait for the kernel launched on the GPU corresponding to that stream to complete before it begins execution . Calling :: cuLaunchCooperativeKernelMultiDevice () sets persistent function state that is the same as function state set through :: cuLaunchKernel API when called individually for each element in \ p launchParamsList . Note that to use :: cuLaunchCooperativeKernelMultiDevice () the kernels must either have been compiled with toolchain version 3 . 2 or later so that it will contain kernel parameter information or have no kernel parameters . If either of these conditions is not met then :: cuLaunchCooperativeKernelMultiDevice () will return :: CUDA_ERROR_INVALID_IMAGE .
public static void RunCooperativeKernelMultiDevie(CudaLaunchParams[] launchParameterList, CudaCooperativeLaunchMultiDeviceFlags flags)
        {
            uint countDevices = (uint)launchParameterList.Length;
            CUResult res = DriverAPINativeMethods.Launch.cuLaunchCooperativeKernelMultiDevice(launchParameterList, countDevices, flags);
            Debug.WriteLine(String.Format("{0:G}, {1}: {2}", DateTime.Now, "cuLaunchCooperativeKernelMultiDevice", res));
            if (res != CUResult.Success) throw new CudaException(res);
        }
public static void RunCooperativeKernelMultiDevie ( CudaLaunchParams [ ] launchParameterList , CudaCooperativeLaunchMultiDeviceFlags flags ) { uint countDevices = ( uint ) launchParameterList . Length ; CUResult res = DriverAPINativeMethods . Launch . cuLaunchCooperativeKernelMultiDevice ( launchParameterList , countDevices , flags ) ; Debug . WriteLine ( String . Format ( " " , DateTime . Now , " " , res ) ) ; if ( res != CUResult . Success ) throw new CudaException ( res ) ; }
Launches CUDA functions on multiple devices where thread blocks can cooperate and synchronize as they execute
Launches CUDA functions on multiple devices where thread blocks can cooperate and synchronize as they execute
launchParameterList
CudaLaunchParams[]
List of launch parameters, one per device
List of launch parameters one per device
flags
CudaCooperativeLaunchMultiDeviceFlags
Flags to control launch behavior
Flags to control launch behavior