Clear Filters
Clear Filters

Using a "CUDAKernel" type object within a parfor loop

3 views (last 30 days)
Thanks in advance for your help.
I'm trying to make use of a CUDAKernel object from within a parfor loop, but when I do so I am met with the message:
Warning: Cannot load an object of class 'CUDAKernel':
No matching constructor signature found.
> In parallel.internal.pool.optionallyDeserialize (line 7)
In parallel.internal.parfor.cppRemoteParallelFunction (line 25)
As a workaround for this issue, I have tried passing in a string which contains the name of the kernel as follows:
cudaKernel = parallel.gpu.CUDAKernel(cudaKernelName+".ptx", cudaKernelName+".cu");
However, directly constructing the kernel in each iteration of the for loop greatly reduces performance.
I am wondering, is there any efficient way that I can pass an existing CUDAKernel object into a parfor loop?

Accepted Answer

Joss Knight
Joss Knight on 22 Jun 2023
A CUDAKernel object cannot be serialized, as you've found, so you will need to construct it separately on each worker. However, you can do this efficiently using parallel.pool.Constant:
kernelConst = parallel.pool.Constant(@()parallel.gpu.CUDAKernel(cudaKernelName+".ptx", cudaKernelName+".cu"));
Then inside the parfor you access the object using kernelConst.Value:
parfor idx = 1:numIterations
% ...
feval(kernelConst.Value, args, etc);
This efficiently constructs the object on each worker just once.
Note that if you don't have a multi-GPU setup, the execution of your kernel object will be serialized anyway, which means the kernels on each worker will not be running at the same time.
Joseph DeCunha
Joseph DeCunha on 22 Jun 2023
Edited: Joseph DeCunha on 22 Jun 2023
@Joss Knight: Implementing this approach on my end I'm encountering the error "Attempting to access the property or method of an invalid object". Do you see this behaviour as well? and if so, do you have any advice on how to proceed?
Edit: Nevermind. I was trying to construct the CUDAKernel constant from an existing kernel through:
kernelConst = parallel.pool.Constant(existingKernel);
When I do as you suggested and create the kernel constant by passing the CUDAKernel constructor I see the expected behaviour.

Sign in to comment.

More Answers (1)

Aditya Singh
Aditya Singh on 22 Jun 2023
Edited: Aditya Singh on 22 Jun 2023
Hello Joseph,
As per my understanding, you are facing an issue in using CUDAKernel in parfor loop, so you tried a workaround. But in this approach, you are making a kernel instance each time, which you want to avoid.
The workaround for not having to create a kernel every time is to utilize a pre-compiled kernel. This approach assumes that the kernel code undergoes minimal modifications with each iteration.
Suppose that you have a CUDA kernel called `vectorAddKernel` that adds two vectors of size `n` and writes the result to a third vector of the same size. The kernel code is presented below:
% Define the kernel source code with placeholders for input parameters
kernelSource = ['#include <cuda.h>\n' ...
'template <int n>\n' ...
'extern "C" __global__ void vectorAddKernel(float* A, float* B, float* C)\n'
'{\n' ...
' int i = threadIdx.x + blockDim.x * blockIdx.x;\n' ...
' if (i < n) {\n' ...
' C[i] = A[i] + B[i];\n' ...
' }\n' ...
% Define the kernel parameters
n = 1024;
numIterations = 10;
% Compile the templated kernel code to PTX code using nvcc and then load
% the file
kernelPTX = fileread('vectorAddKernel.ptx');
compiledKernel = parallel.gpu.CUDAKernel(kernelPTX, 'vectorAddKernel');
% Inside the parfor loop, call the compiled kernel with the necessary arguments
parfor idx = 1:numIterations
% Call kernel with arguments
blockSize = [256,1,1]; % block size
gridSize = [ceil(n/256),1,1]; % grid size
output{idx} = zeros(n, 1, 'single', 'gpuArray');
feval(compiledKernel, blockSize, gridSize, {gpuArray(single(rand(n, 1))), gpuArray(single(rand(n, 1))), output{idx}});
For reference, kindly see:
Hope this helps!
Joss Knight
Joss Knight on 22 Jun 2023

You can directly compile to PTX from MATLAB using mexcuda with the -ptx option. You should probably add that (and check that the code runs).

Joseph DeCunha
Joseph DeCunha on 22 Jun 2023
Aditya, if you run this code you will find it experiences the same issue as my original post. Your CUDAKernel object is inaccessible from within the parfor loop.

Sign in to comment.


Find more on Parallel Computing Fundamentals in Help Center and File Exchange




Community Treasure Hunt

Find the treasures in MATLAB Central and discover how the community can help you!

Start Hunting!