Library/ManagedCuda/CudaKernel.cs
/// <summary>
/// Launches a CUDA function where thread blocks can cooperate and synchronize as they execute
/// <para/>
/// Invokes the kernel \p f on a \p gridDimX x \p gridDimY x \p gridDimZ
/// grid of blocks.Each block contains \p blockDimX x \p blockDimY x
/// \p blockDimZ threads.
/// <para/>
/// \p sharedMemBytes sets the amount of dynamic shared memory that will be
/// available to each thread block.
/// <para/>
/// The device on which this kernel is invoked must have a non-zero value for
/// the device attribute::CU_DEVICE_ATTRIBUTE_COOPERATIVE_LAUNCH.
/// <para/>
/// The total number of blocks launched 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.
/// <para/>
/// The kernel cannot make use of CUDA dynamic parallelism.
/// <para/>
/// Kernel parameters must be specified via \p kernelParams.  If \p f
/// has N parameters, then \p kernelParams needs to be an array of N
/// pointers.  Each of \p kernelParams [0]
/// through \p kernelParams [N-1]
/// must point to a region of memory from which the actual kernel
/// parameter will be copied.  The number of kernel parameters and their
/// offsets and sizes do not need to be specified as that information is
/// retrieved directly from the kernel's image.
/// <para/>
/// Calling ::cuLaunchCooperativeKernel() sets persistent function state that is
/// the same as function state set through ::cuLaunchKernel API
/// <para/>
/// When the kernel \p f is launched via ::cuLaunchCooperativeKernel(), the previous
/// block shape, shared size and parameter info associated with \p f
/// is overwritten.
/// <para/>
/// Note that to use ::cuLaunchCooperativeKernel(), the kernel \p f 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 ::cuLaunchCooperativeKernel() will
/// return ::CUDA_ERROR_INVALID_IMAGE.
/// </summary>
Launches a CUDA function where thread blocks can cooperate and synchronize as they execute



\p sharedMemBytes sets the amount of dynamic shared memory that will be
available to each thread block.

The device on which this kernel is invoked must have a non-zero value for
the device attribute::CU_DEVICE_ATTRIBUTE_COOPERATIVE_LAUNCH.



The kernel cannot make use of CUDA dynamic parallelism.

Kernel parameters must be specified via \p kernelParams.  If \p f
has N parameters, then \p kernelParams needs to be an array of N
pointers.  Each of \p kernelParams [0]
through \p kernelParams [N-1]
must point to a region of memory from which the actual kernel
parameter will be copied.  The number of kernel parameters and their
offsets and sizes do not need to be specified as that information is
retrieved directly from the kernel's image.

Calling ::cuLaunchCooperativeKernel() sets persistent function state that is
the same as function state set through ::cuLaunchKernel API

When the kernel \p f is launched via ::cuLaunchCooperativeKernel(), the previous
block shape, shared size and parameter info associated with \p f
is overwritten.

Note that to use ::cuLaunchCooperativeKernel(), the kernel \p f 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 ::cuLaunchCooperativeKernel() will
return ::CUDA_ERROR_INVALID_IMAGE.
Launches a CUDA function where thread blocks can cooperate and synchronize as they execute \ p sharedMemBytes sets the amount of dynamic shared memory that will be available to each thread block . The device on which this kernel is invoked must have a non - zero value for the device attribute :: CU_DEVICE_ATTRIBUTE_COOPERATIVE_LAUNCH . The kernel cannot make use of CUDA dynamic parallelism . Kernel parameters must be specified via \ p kernelParams . If \ p f has N parameters then \ p kernelParams needs to be an array of N pointers . Each of \ p kernelParams [ 0 ] through \ p kernelParams [ N - 1 ] must point to a region of memory from which the actual kernel parameter will be copied . The number of kernel parameters and their offsets and sizes do not need to be specified as that information is retrieved directly from the kernel ' s image . Calling :: cuLaunchCooperativeKernel () sets persistent function state that is the same as function state set through :: cuLaunchKernel API When the kernel \ p f is launched via :: cuLaunchCooperativeKernel () the previous block shape shared size and parameter info associated with \ p f is overwritten . Note that to use :: cuLaunchCooperativeKernel () the kernel \ p f 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 :: cuLaunchCooperativeKernel () will return :: CUDA_ERROR_INVALID_IMAGE .
public virtual void RunCooperative(CUstream stream, params object[] parameters)
        {
            int paramCount = parameters.Length;
            IntPtr[] paramsList = new IntPtr[paramCount];
            GCHandle[] GCHandleList = new GCHandle[paramCount];
            for (int i = 0; i < paramCount; i++)
            {
                GCHandleList[i] = GCHandle.Alloc(parameters[i], GCHandleType.Pinned);
                paramsList[i] = GCHandleList[i].AddrOfPinnedObject();
            }
            res = DriverAPINativeMethods.Launch.cuLaunchCooperativeKernel(_function, _gridDim.x, _gridDim.y, _gridDim.z, _blockDim.x, _blockDim.y, _blockDim.z, _sharedMemSize, stream, paramsList);
            Debug.WriteLine(String.Format("{0:G}, {1}: {2}, Kernel: {3}", DateTime.Now, "cuLaunchCooperativeKernel", res, _kernelName));
            for (int i = 0; i < paramCount; i++)
            {
                GCHandleList[i].Free();
            }
            if (res != CUResult.Success) throw new CudaException(res);
        }
public virtual void RunCooperative ( CUstream stream , params object [ ] parameters ) { int paramCount = parameters . Length ; IntPtr [ ] paramsList = new IntPtr [ paramCount ] ; GCHandle [ ] GCHandleList = new GCHandle [ paramCount ] ; for ( int i = 0 ; i < paramCount ; i ++ ) { GCHandleList [ i ] = GCHandle . Alloc ( parameters [ i ] , GCHandleType . Pinned ) ; paramsList [ i ] = GCHandleList [ i ] . AddrOfPinnedObject ( ) ; } res = DriverAPINativeMethods . Launch . cuLaunchCooperativeKernel ( _function , _gridDim . x , _gridDim . y , _gridDim . z , _blockDim . x , _blockDim . y , _blockDim . z , _sharedMemSize , stream , paramsList ) ; Debug . WriteLine ( String . Format ( " " , DateTime . Now , " " , res , _kernelName ) ) ; for ( int i = 0 ; i < paramCount ; i ++ ) { GCHandleList [ i ] . Free ( ) ; } if ( res != CUResult . Success ) throw new CudaException ( res ) ; }
Launches a CUDA function where thread blocks can cooperate and synchronize as they execute
Launches a CUDA function where thread blocks can cooperate and synchronize as they execute
//Get pointers to kernel parameters //Launch the kernel //Free pinned managed parameters //First free resources then throw potential exception!
stream
CUstream