in the opencl node there is the built-in iterations.
How can I reference that in the opencl code though? such as for modifying a seed value with the iteration?
Thanks for any help!
How to reference the iteration number in OpenCL?
2045 5 0-
- Barrett Meeker
- Member
- 54 posts
- Joined: March 2015
- Offline
-
- elovikov
- Member
- 122 posts
- Joined: June 2019
- Offline
I'm not an expert in opencl so not sure if it's even possible to get this natively on device. It's probably in order command queue and as far as I know you can't get the kernel index in queue on the device.
The only way I can think of is incrementing detail attribute in write back kernel.
Consider this simple example (hip attached). It just modifies point position on current iteration index. So the more iterations the more points modified.
(geo has to have a detail integer attribute "iteration")
With that you can use iteration number without race conditions on your main kernel.
Again I'm not sure if this is the right approach, may be there is a smarter way. But if you already have a write back kernel you can utilize it without much overhead. If not, well you have to use the dummy one.
The only way I can think of is incrementing detail attribute in write back kernel.
Consider this simple example (hip attached). It just modifies point position on current iteration index. So the more iterations the more points modified.
(geo has to have a detail integer attribute "iteration")
kernel void kernelName( int P_length, global float * P , int iteration_length, global int * iteration ) { int idx = get_global_id(0); if (idx >= P_length) return; float3 pos = vload3(idx, P); int iter = iteration[0]; pos.y = (idx == iter) ? pos.y + 0.2f : pos.y; vstore3(pos, idx, P); } kernel void writeBack( int P_length, global float * P , int iteration_length, global int * iteration ) { int idx = get_global_id(0); if (idx == P_length - 1) iteration[0]++; }
With that you can use iteration number without race conditions on your main kernel.
Again I'm not sure if this is the right approach, may be there is a smarter way. But if you already have a write back kernel you can utilize it without much overhead. If not, well you have to use the dummy one.
-
- elovikov
- Member
- 122 posts
- Joined: June 2019
- Offline
Actually after a little bit of thinking I find a better way.
I think in this case it's probably better to organize kernels in compile blocks instead of using sop iterations feature. With this you don't have to use the same work-item and basically just increment iteration on the separate kernel.
Performance should be the same.
I think in this case it's probably better to organize kernels in compile blocks instead of using sop iterations feature. With this you don't have to use the same work-item and basically just increment iteration on the separate kernel.
Performance should be the same.
-
- Barrett Meeker
- Member
- 54 posts
- Joined: March 2015
- Offline
Thanks for those examples!
I was under the impression though it would be good to use the built in opencl iteration vs a compile block because using the iterations on the opencl node means the data stays on the graphics card rather than the card having to get all the data each iteration if using a solver or compile block. That was just my understanding though, not sure if that is correct.
I thought of one other way that I haven't tried but I think might work - use the writeback to check if the work id is 0, if it is then iterate a global int *
I think that will work? I might try it. I have some opencl nodes inside a solver and I'd kind of like to try sub-stepping them internally if that makes sense...
I was under the impression though it would be good to use the built in opencl iteration vs a compile block because using the iterations on the opencl node means the data stays on the graphics card rather than the card having to get all the data each iteration if using a solver or compile block. That was just my understanding though, not sure if that is correct.
I thought of one other way that I haven't tried but I think might work - use the writeback to check if the work id is 0, if it is then iterate a global int *
I think that will work? I might try it. I have some opencl nodes inside a solver and I'd kind of like to try sub-stepping them internally if that makes sense...
-
- elovikov
- Member
- 122 posts
- Joined: June 2019
- Offline
I actually also thought that builtin iterations is the only way to keep data on the device.
But here: https://vimeo.com/241568199 [vimeo.com] (on 29 minutes mark) Jeff Lait said that iterations are just shortcut to compile block with for loop that contains only one opencl sop. If you need more stages in your solver you can use explicit compile blocks.
I think the "node compiler" is smart enough and treats for loops with only opencl sops in that way so it puts loop body into one queue (so no copy to the host involved).
Obviously I'm just guessing
I don't know technical details about node compilation but I'd say it safe to use compile blocks with chained opencl sops.
But here: https://vimeo.com/241568199 [vimeo.com] (on 29 minutes mark) Jeff Lait said that iterations are just shortcut to compile block with for loop that contains only one opencl sop. If you need more stages in your solver you can use explicit compile blocks.
I think the "node compiler" is smart enough and treats for loops with only opencl sops in that way so it puts loop body into one queue (so no copy to the host involved).
Obviously I'm just guessing
I don't know technical details about node compilation but I'd say it safe to use compile blocks with chained opencl sops.
-
- Barrett Meeker
- Member
- 54 posts
- Joined: March 2015
- Offline
-
- Quick Links