Joyee
August 16, 2011, 6:08am
1
Hi all,
I have a error when I am trying to compile Merrill’s radix sort under win-XP + VS2005.
error: asm operand type size(1) does not match type/size implied by constraint ‘r’
it occurs in the following code
#define B40C_DEFINE_GLOBAL_LOAD(base_type, dest_type, short_type, ptx_type, reg_mod)\
asm("ld.global.cg."#ptx_type" %0, [%1];" : "="#reg_mod(dest) : _B40C_ASM_PTR_(d_ptr + offset));\
...
B40C_DEFINE_GLOBAL_LOAD(char, signed char, char, s8, r)
Check out the bottom of page 8 in the “Using Inline PTX Assembly In Cuda”
error: an asm operand must have scalar type
The type and size implied by a PTX asm constraint must match that of the associated operand. Example where size does not match:
For ‘char’ type variable “ciâ€,
asm(“add.s32 %0,%1,%2;”:“=r”(ci):“r”(j),“r”(k));
error: asm operand type size(1) does not match type/size implied by constraint ‘r’
In order to use ‘char’ type variables “ciâ€, “cjâ€, and “ck†in the above asm statement, code
segment similar to the following may be used,
int temp = ci;
asm(“add.s32 %0,%1,%2;”:“=r”(temp):“r”((int)cj),“r”((int)ck));
ci = temp;
Another example where type does not match:
For ‘float’ type variable “fiâ€,
asm(“add.s32 %0,%1,%2;”:“=r”(fi):“r”(j),“r”(k));
error: asm operand type size(4) does not match type/size implied by constraint ‘r’
SeanB
September 16, 2011, 5:21am
3
You can also use my radix sort:http://www.moderngpu.com/sort/mgpusort.html
The library should compile under VC2005, although the tests require VC2008’s tr1 for mt19937. The .sln I included is for vs9. But for this you only need to compile cucpp.cpp, mgpusort.cpp, and sorttables.cpp. If you want it as a DLL (not required), use the mgpusort.def manifest.
sean