Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Declare array or vector from pointer inside kernel #206

Closed
BernardoCovas opened this issue Apr 7, 2024 · 5 comments
Closed

Declare array or vector from pointer inside kernel #206

BernardoCovas opened this issue Apr 7, 2024 · 5 comments
Assignees
Labels
feature request Request for something to be added

Comments

@BernardoCovas
Copy link

Hi! Is there a way that I can have a list of integers inside a kernel representing pointers to array data of known sizes?
Curretly I am working with big image arrays of different sizes that barely fit in memory, and thus I can not copy them into some sort of contiguous array due to copy delays and memory space. Moreover, I will be loading and unloading a portion of these images in a dynamic way as the computation progresses. Which images will be loaded or unloaded will depend on other factors, but some will remain while others will be unloaded and new ones will be loaded. Since images can have vastly different sizes, I don't think I can use a sort of rolling buffer. I would like to use warp for processing, but currently I have not found a way to declare inside a warp kernel that an integer is an array pointer. Which image I am going to access will depend on computation, and will vary from iteration to iteration inside the same kernel. What I wanted to achive is something like the following:

image0 = wp.array(...) # 10000 x 10000 image
image1 = wp.array(...) # 10000 x 10000 image
image2 = wp.array(...) # 10000 x 10000 image

image_source = wp.array(...) # 10000 x 10000 image

@wp.kernel
def kn(image_src, image_trgts, image_trgts_shape):
    x, y = wp.tid()
    target_image_index = (...) # target image index will depend on some computation for the current image pixel
    target_x = (...) # target image coordinates will depend on some other computation
    target_y = (...) # target image coordinates will depend on some other computation
 
    # get the sape of target image from the shapes array
    tgt_h = image_trgts_shape[target_image, 0]
    tgt_w = image_trgts_shape[target_image, 1]

    # declare the array for the target image that I want to access
    target = wp.array(pointer=image_trgts[target_image], shape=(tgt_h, tgt_w, 3))
    pixel_rgb = wp.vec3f(
        target[target_y, target_x, 0],
        target[target_y, target_x, 1],
        target[target_y, target_x, 2])
    # finally work with target image pixel
    (...)

The real computation has plenty more steps, however this example outlines the issue.

Thank you in advance

@mmacklin
Copy link
Collaborator

Hi @BernardoCovas, this looks like an interesting use case, and I can see that having the ability to construct an array at kernel time from ptr would be useful indeed.

Let me discuss with the team and get back to you on how we can support this.

Thanks,
Miles

@mmacklin
Copy link
Collaborator

Just want to ask a clarifying question - in your example, would image_trgts be an array of e.g.: uint64 addresses representing the array memory for an image?

@BernardoCovas
Copy link
Author

Hi @mmacklin , thanks for the reply
Yes! I would have to create it somehow beforehand

@mmacklin mmacklin assigned mmacklin and nvlukasz and unassigned mmacklin Apr 11, 2024
@shi-eric shi-eric added the feature request Request for something to be added label Apr 11, 2024
@christophercrouzet
Copy link
Member

Hi @BernardoCovas, thanks for opening this issue!

We've just added a new wp.array(ptr=...) built-in with the commit 261f2a9, so it will be part of the next release of Warp.

@BernardoCovas
Copy link
Author

Hello @christophercrouzet , @mmacklin
Thank you very much!

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
feature request Request for something to be added
Projects
None yet
Development

No branches or pull requests

5 participants