warp icon indicating copy to clipboard operation
warp copied to clipboard

Declare array or vector from pointer inside kernel

Open BernardoCovas opened this issue 1 year ago • 3 comments

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

BernardoCovas avatar Apr 07 '24 17:04 BernardoCovas

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 avatar Apr 10 '24 01:04 mmacklin

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?

mmacklin avatar Apr 10 '24 04:04 mmacklin

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

BernardoCovas avatar Apr 10 '24 09:04 BernardoCovas

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.

christophercrouzet avatar Jul 26 '24 02:07 christophercrouzet

Hello @christophercrouzet , @mmacklin Thank you very much!

BernardoCovas avatar Jul 29 '24 16:07 BernardoCovas