google/clspv

Q: how to trigger the use of uniform buffers.

stolk opened this issue · 2 comments

stolk commented

So, this is more of a support question than an issue, sorry about that! If inappropriate, I will delete it.

I would like clspv to generate SPIR-V with the use of uniform buffers for arguments, instead of storage buffers.

Is there some OpenCL syntax that will induce the use of uniform buffers? Or a clspv option?

The result I am trying to achieve, is having the kernel use constant memory which in CUDA speak would be __constant__ prefixed.

So far I tried --pod-ubo which did not help me.

So in short: How to get the src data of the kernel below in constant memory?

#define uint32_t	uint

__kernel
void foo
(
	__global const uint32_t* __restrict__ src,
	__global uint32_t* __restrict__ dst
)
{
	const uint32_t pindex = get_global_id(0);
	uint32_t s = src[pindex];
	dst[pindex] = s ^ 0xffffffff;
}

Thanks!

You can try --constant-args-ubo.
Note that you might also want to have a look at: --max-ubo-size, relaxed-ubo-layout and std430-ubo-layout.

I think it is completely fine to open issue to ask this question. But know that we have a discord server as well (linked in the README), it's up to you to use the one you prefer.

Feel free to reopen if one needs more help on that topic.