Kernel launch errors from cudaGetLastError are not handled

[aryan-programmer]

The enqueue_raw_kernel_launch_in_current_context does not check for an error with cudaGetLastError, thus kernel launch errors like cudaErrorInvalidConfiguration go uncaught, and the kernel launch silently fails.

Minimal example:

The following example demonstrates an example kernel, whose launch should fail since the maximum number of threads per block is 1024, and we are trying to launch it with 1500 threads.

#include <device_launch_parameters.h>
#include <stdio.h>
#include <iostream>
#include <cuda/api.hpp>

using namespace std;

__global__ void test_kernel() {
	int i = blockIdx.x * blockDim.x + threadIdx.x;
	if (i == 0) {
		printf("Hello CUDA\n");
	}
}

int main() {
	if (cuda::device::count() == 0) {
		std::cerr << "No CUDA devices on this system" << "\n";
		exit(EXIT_FAILURE);
	}

	cuda::device::current::set(cuda::device::get(0));
	auto device = cuda::device::current::get();

	try {
		cout << "Executing the kernel:" << endl;
		cuda::launch_configuration_t lc = cuda::launch_config_builder()
										 .overall_size(2048)
										 .block_dimensions(1500)
										 .build();
		cuda::launch(test_kernel, lc);
	} catch (std::exception ex) {
		cout << ex.what() << endl;
	}
	cuda::synchronize(device);

	return 0;
}

Current output:

Executing the kernel:

The kernel launch fails silently.

Expected output:

Executing the kernel:
Kernel launch failed: invalid configuration argument

The kernel launch error from cudaGetLastError is handled and converted into an exception that is caught here.

[eyalroz]

Please verify that this works for you now on the development branch.

[aryan-programmer]

It catches the error in the release build correctly.

However in the debug build while the launch parameter validation does catch the error properly, it also results in false alarms as follows:

Specifying the block_dimensions to be 32 results in the following error:

Executing the kernel:
specified block X-axis dimension 32 exceeds the maximum supported X dimension of 1024 for device 0

Specifying the block_dimensions to be 1024 results in the following error:

Executing the kernel:
specified block Y-axis dimension 1 exceeds the maximum supported Y dimension of 1024 for device 0

Both of these should have worked and printed "Hello World". Indeed, in release (with the block_dimensions being 32 or 1024) it outputs:

Executing the kernel:
Hello CUDA

[eyalroz]

Can you try again?

[aryan-programmer]

It works correctly now.