Getting error cudaErrorInvalidConfiguration.


I am getting the following error when I run my code with cuda-memcheck:
Program hit cudaErrorInvalidConfiguration (error 9) due to “invalid configuration argument” on CUDA API call to cudaLaunchKernel.

When I googled this error, I found out that this error is given out for improper kernel configuration. However, when I printed out my kernel configuration I get the following config:
gridDim.x, gridDim.y, gridDim.z, blockDim.x, blockDim.y, blockDim.z --> 1, 3, 2, 1024, 1, 1 .

All the values are within kernel specification. To cross check I am pasting deviceQuery result here:

/home/gyana/NVIDIA_CUDA-10.0_Samples/1_Utilities/deviceQuery/deviceQuery Starting…

CUDA Device Query (Runtime API) version (CUDART static linking)

Detected 1 CUDA Capable device(s)

Device 0: “GeForce GTX 1050”
CUDA Driver Version / Runtime Version 10.1 / 10.0
CUDA Capability Major/Minor version number: 6.1
Total amount of global memory: 4042 MBytes (4238737408 bytes)
( 5) Multiprocessors, (128) CUDA Cores/MP: 640 CUDA Cores
GPU Max Clock rate: 1493 MHz (1.49 GHz)
Memory Clock rate: 3504 Mhz
Memory Bus Width: 128-bit
L2 Cache Size: 524288 bytes
Maximum Texture Dimension Size (x,y,z) 1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384)
Maximum Layered 1D Texture Size, (num) layers 1D=(32768), 2048 layers
Maximum Layered 2D Texture Size, (num) layers 2D=(32768, 32768), 2048 layers
Total amount of constant memory: 65536 bytes
Total amount of shared memory per block: 49152 bytes
Total number of registers available per block: 65536
Warp size: 32
Maximum number of threads per multiprocessor: 2048
Maximum number of threads per block: 1024
Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535)
Maximum memory pitch: 2147483647 bytes
Texture alignment: 512 bytes
Concurrent copy and kernel execution: Yes with 2 copy engine(s)
Run time limit on kernels: Yes
Integrated GPU sharing Host Memory: No
Support host page-locked memory mapping: Yes
Alignment requirement for Surfaces: Yes
Device has ECC support: Disabled
Device supports Unified Addressing (UVA): Yes
Device supports Compute Preemption: Yes
Supports Cooperative Kernel Launch: Yes
Supports MultiDevice Co-op Kernel Launch: Yes
Device PCI Domain ID / Bus ID / location ID: 0 / 1 / 0
Compute Mode:
< Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 10.1, CUDA Runtime Version = 10.0, NumDevs = 1
Result = PASS

Can someone tell me what am I doing wrong?

Can you post a code snippet showing how exactly you launch the kernel?

A possible third argument given to the <<< >>> kernel launch syntax would request shared memory resources per block. If this exceeds 48kb, the same error would be thrown.

A possible fourth argument given would specify a CUDA stream. If it is an invalid stream handle, you would also see an error.