oscarbg

//goes well until ptxas which says something about invalid arguments so at least f16x2 modified atom instruction is recognized?
u32
atomf16x2(u32 a, u32 b)
{
	u32 d;

	asm("atom.global.add.f16x2 	%0, [%1], %2;" : "=r"(d) : "r"(a), "r"(b));
	//atom.global.add.u32 	%r5, [%rd2], 10;
	//asm("mul.wide.s16 %0, %1, %2;" : "=r"(d) : "h"(a), "h"(b));

oscarbg

// -*- compile-command: "nvcc -m 32 -arch sm_50 -Xptxas=-v,-abi=no -cubin int_mul.cu" ; -*-

#include <stdint.h>

//
//
//

#define KERNEL_QUALIFIERS                   __global__
#define KERNEL_QUALIFIERS_EXTERN            extern     KERNEL_QUALIFIERS