Using a "CUDAKernel" type object within a parfor loop

Question

Joseph DeCunha 2023-6-21

0
链接

此问题的直接链接

https://ww2.mathworks.cn/matlabcentral/answers/1986459-using-a-cudakernel-type-object-within-a-parfor-loop

评论： Joss Knight 2023-6-23

采纳的回答： Joss Knight

在 MATLAB Online 中打开

Hello,

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?

Joseph

0 个评论
显示 -2更早的评论隐藏 -2更早的评论

请先登录，再进行评论。

请先登录，再回答此问题。

Answer 1

Joss Knight 2023-6-22

2
链接

此回答的直接链接

https://ww2.mathworks.cn/matlabcentral/answers/1986459-using-a-cudakernel-type-object-within-a-parfor-loop#answer_1260484

在 MATLAB Online 中打开

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);
end

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.

3 个评论
显示 1更早的评论隐藏 1更早的评论

Joseph DeCunha 2023-6-22

编辑：Joseph DeCunha 2023-6-22

在 MATLAB Online 中打开

@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.

Joss Knight 2023-6-23

That's good to know, thanks!

请先登录，再进行评论。

Answer 2

Aditya Singh 2023-6-22

0
链接

此回答的直接链接

https://ww2.mathworks.cn/matlabcentral/answers/1986459-using-a-cudakernel-type-object-within-a-parfor-loop#answer_1260439

编辑：Aditya Singh 2023-6-22

在 MATLAB Online 中打开

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}}); 
end 

For reference, kindly see:

Hope this helps!

4 个评论
显示 2更早的评论隐藏 2更早的评论

Joss Knight 2023-6-22

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 2023-6-22

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.

请先登录，再进行评论。