Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

after aligned d_stride :::get error ::::RuntimeError: CUDA error: an illegal memory access was encountered #87

Open
SWALLOWhajnal opened this issue Nov 7, 2024 · 3 comments

Comments

@SWALLOWhajnal
Copy link

I specifically aligned d_stride, but it still reported a memory problem

RuntimeError: CUDA error: an illegal memory access was encountered
CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect.
For debugging consider passing CUDA_LAUNCH_BLOCKING=1.
Compile with TORCH_USE_CUDA_DSA to enable device-side assertions

@SWALLOWhajnal
Copy link
Author

my batch_size=1 It's working.
my batch_size=2 It's RuntimeError: CUDA error: an illegal memory access was encountered
my batch_size=3 It's working.

@SWALLOWhajnal
Copy link
Author

Error from dcnv4_col2im_cuda.cuh


int shm_size = sizeof(opmath_t) * (G * block_multiplier * K) * 2;
if(!check_backward_warpp(d_stride, D)){
shm_size = sizeof(opmath_t) * ((G * block_multiplier * K) * 2 + G * block_multiplier * blockdimX * 3);
}
cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize,
shm_size);

kernel<<<num_blocks, num_threads, shm_size, stream>>>(
value, p_offset, grad_output, G, D, Q, kernel_h, kernel_w, stride_h,
stride_w, pad_h, pad_w, dilation_h, dilation_w, height_in, width_in,
height_out, width_out, offset_scale, remove_center, block_multiplier,
grad_im, grad_offset, padded_offset_dim);

cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("error in dcnv4_im2col_cuda: %s\n", cudaGetErrorString(err));
printf("launch arguments: gridDim=(%d, %d, %d), blockDim=(%d, %d, %d), "
"shm_size=%d\n\n",
num_blocks.x, num_blocks.y, num_blocks.z, num_threads.x,
num_threads.y, num_threads.z, shm_size);
AT_ASSERTM(false, "kernel launch error");
}
}

error in dcnv4_im2col_cuda: an illegal memory access was encountered
launch arguments: gridDim=(8192, 1, 1), blockDim=(12, 16, 1), shm_size=3456。
/DCNv4_op/src/cuda/dcnv4_col2im_cuda.cuh":470, please report a bug to PyTorch. kernel launch error

I can't fix this error.

@Q-M-E
Copy link

Q-M-E commented Dec 10, 2024

I have encountered the same issue.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

No branches or pull requests

2 participants