Set Max Dynamic Shared Memory

[bastianhagedorn]

Hi,

I need more than 48K dynamically allocated shared memory on a Titan RTX. The corresponding CUDA call would be

    cudaCheckError(cudaFuncSetAttribute(
        my_kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, 65536));

How can I set this function attribute using pycuda?

Thanks!

[inducer]

https://gitlab.tiker.net/inducer/pycuda/-/merge_requests/60 might do the trick.

[bastianhagedorn]

that looks good but I don't know yet how to use it. This is my python code so far:

  128     kernel = mod.get_function(kernel_name) 
  129     grid, block = (mygrid, myblock) 
~ 130     shared = int(max_dyn_smem_size)  
+ 131    # kernel.set_attribute(..) or something similar? 
  132     info(f"Launching {kernel_name}<<<{grid},{block},{shared}>>>") 
  133     kernel(
  134         *device_buffers,
  135         block=...
  136         grid=...
  137         shared=shared,
  138     )

What does the python API look like for setting this attribute?

Thanks

[bastianhagedorn]

got it kernel.set_attribute(cuda.function_attribute.MAX_DYNAMIC_SHARED_SIZE_BYTES, shared) thanks a lot for the quick fix!

[inducer]

I forgot to document those. Here's that: https://gitlab.tiker.net/inducer/pycuda/-/merge_requests/62