"SimpleKernel" example throws exception

[AlexBatsev]

I've tried to run SimpleKernel example and has got
ILGPU.Runtime.Cuda.CudaException: 'a PTX JIT compilation failed'

Inside CudaKernel constructor, the function CudaAPI.Current.LoadModule(...) returned
CudaError cudaStatus = CUDA_ERROR_INVALID_PTX;

string errorLog = "
_ptxas application ptx input, line 5; fatal   : Unsupported .version 6.0; current version is '5.0'
ptxas fatal   : Ptx assembly aborted due to errors_
"

Indeed, on line 5 one can see ".version 6.0":

kernel.PTXAssembly =

//
// Generated by ILGPU v0.5.1
//

.version 6.0
.target sm_50
.address_size 64


.visible .entry ILGPUKernel(
.param .align 4 .b8 _index_1318[4],
.param .align 8 .b8 _dataView_1319[16],
.param .b32 _constant_1320
)
{
	.reg .pred	%p<2>;
	.reg .b32	%r<8>;
	.reg .b64	%rd<4>;

	ld.param.b32	%r2, [_index_1318];
	ld.param.b64	%rd1, [_dataView_1319];
	ld.param.b32	%r3, [_dataView_1319+8];
	ld.param.b32	%r4, [_constant_1320];
	mov.b32	%r5, %ctaid.x;
	mov.b32	%r6, %ntid.x;
	mov.b32	%r7, %tid.x;
	mad.lo.s32	%r1, %r5, %r6, %r7;
	setp.ge.s32	%p1, %r1, %r2;
	@%p1 ret;

	L\_13170:
	mul.wide.u32	%rd3, %r1, 4;
	add.u64	%rd2, %rd1, %rd3;
	add.s32	%r7, %r1, %r4;
	st.b32	[%rd2], %r7;
	ret;
}

If it is important, my graphic card is Quadro M2000M.

Is this a bug or do I misunderstand something?
Thank you.

[MoFtZ]

hi @AlexBatsev, have you tried updating your CUDA drivers? This appears to be a similar problem to #29.

[AlexBatsev]

After I've updated drivers the problem's gone.
After updating ILGPU to 0.6.0, the .version changed to "6.4".
Question then: maybe it makes sense to check if driver is capable of certain version and generate shader accordingly?
Otherwise each time ILGPU package is updated, drivers should be updated as well. Right?
What if I have old graphic card, for which driver updates are not available any longer?
BTW, while debugging I found out that the ILGPU code looks nice and is easy to perceive and understand. Very good work, IMHO.

[m4rs-mt]

It makes sense to add additional driver version checks to check compatibility. We should add them as soon as possible to improve usability :)

Otherwise each time ILGPU package is updated, drivers should be updated as well. Right?

In general: no. We had to update the PTX version to support new GPU features. However, the upcoming release of the ILGPU compiler has a completely new intrinsic processing pipeline. This allows us to define version-dependent implementations and specialized code generation for different GPU targets. This should help us to work around the problem of general GPU support for older hardware platforms that you mentioned.

[AlexBatsev]

@m4rs-mt thank you for answers.