When I define a large (>1024 byte) fixed-size array within a kernel, the compiler appears to create a global memory allocation which persists past the completion of the kernel and, as far as I have been able to tell, there is no way to free this allocation for the remainder of the program. I have at…

[image] CUDA thread local array memory is not released after finished kernel execution CUDA Programming and Performance I have encountered a problem of which the cuda thread local array memory is not released after kernel execution. I am using CUDA 9.1 (comput…

[image] john310: When I define a large (>1024 byte) fixed-size array within a kernel Allocating large-ish arrays in functions has never been a great idea, IMHO, but stylistic preferences differ. Alternatively, you could dynamically allocate memory for the workspace needed by the kernel (with …

For my application the performance impact of the large array is negligible compared to other bottlenecks, and makes for some more convenient code. That being said, I am more concerned about the discovery that there is any scenario in which a kernel is creating a global memory allocation which persis…

I am not sure where the notion of “spilled memory” comes from here. What I am suggesting is that all large data objects should be allocated outside of functions, and pointers or references to these data objects then passed to the functions, completely independent of the use of CUDA. For CUDA specifi…

Thanks, appreciate the terminology clarification as yes I was mingling the concepts of truly spilled registers and local arrays potentially being optimized into register memory. In my application, the REG_ARR_LN is actually generated by a preprocessor and can be either small enough to potentially be…

Interesting to see that the issue was discussed five years ago. Was a bug filed at the time? If so, what is the status of that bug? My assumption is that always freeing the memory has negative performance implications, which is why the equivalent action of the suggested workaround is is not applied …

Given that was 5 years ago, any information on whether this is something that is still being worked on? Edit: Also, when I try inserting the cudaDeviceSetLimit(cudaLimitStackSize, 0); after my kernel call, I get an illegal memory access was encountered on the subsequent cudaMemGetInfo call, so it d…

[image] john310: Given that was 5 years ago, any information on whether this is something that is still being worked on? Nothing is still being worked on. I did file a bug (2126657) ~6 years ago related to this, however it tracked separate observations from what was reported by the original …

#include <iostream> constexpr int N = 1 << 20; constexpr int TPB = 128; constexpr int BPG = (N + TPB - 1) / TPB; constexpr int REG_ARR_LN = 32 * 64; __global__ void kernel(int loop_num) { uint32_t spilled[REG_ARR_LN]; for (int i = 0; i < loop_num; ++i) { spilled[i] = spilled[(i - 1) % REG…

You may wish to learn more about code verification, debugging, and error reporting in CUDA. Unit 12 of this online course covers relevant topics and may be of interest. You have illegal code . This has nothing to do with your usage of the cudaDeviceSetLimit() api. If you run your code under compu…

Yep noticed that, that’s what I get for trying to respond quickly; missed the cudaDeviceSynchronize after kernel call to successfully register the error that happened there. So nevermind, the cudaDeviceSetLimit api indeed solves the problem for me. Thank you!

Global memory "leak" after register spill?

Accelerated Computing CUDA CUDA Programming and Performance

Robert_Crovella August 20, 2024, 9:53pm 5

Topic		Replies	Views
global memory lost CUDA Programming and Performance	9	6663	November 4, 2010
Global memory occupied until cudaDeviceReset() or app exits CUDA Programming and Performance	0	2550	June 25, 2014
Unaccounted memory consumption while running the kernel CUDA Programming and Performance	2	332	August 29, 2023
How to release register CUDA Programming and Performance cuda	6	1202	April 28, 2022
using cudaMalloc and cudaFree within a loop unspecified launch failure! CUDA Programming and Performance	21	37985	April 23, 2009
Global memory's lifespan : Application. How do you do that? How could you allocate a variable to CUDA Programming and Performance	17	2852	October 3, 2011
Can we do malloc inside a __global__ function CUDA Programming and Performance	26	10113	February 21, 2010
CUDA thread local array memory is not released after finished kernel execution CUDA Programming and Performance	2	1428	December 1, 2019
malloc memory in kernel linked via in/out variable CUDA Programming and Performance	10	2111	October 17, 2015
Registers, spilling and local memory - Unexpected behaviour? CUDA Programming and Performance	3	2027	January 20, 2020

Global memory "leak" after register spill?

Related topics