 	NONE,
 	CG,
 	CS, 
-	CA
+	CA,
+	WB
 };
 
 

	 */
	#define B40C_DEFINE_GLOBAL_LOAD(base_type, dest_type, short_type, ptx_type, reg_mod)																												\
		template <> struct ModifiedLoad<base_type, CG> {																												\
			__device__ __forceinline__ static void Ld(dest_type &dest, base_type* d_ptr, int offset) {																\
				asm("ld.global.cg."#ptx_type" %0, [%1];" : "="#reg_mod(dest) : _B40C_ASM_PTR_(d_ptr + offset));														\
			}																																						\
		};																																							\
		template <> struct ModifiedLoad<base_type, CS> {																												\
			__device__ __forceinline__ static void Ld(dest_type &dest, base_type* d_ptr, int offset) {																\
				asm("ld.global.cs."#ptx_type" %0, [%1];" : "="#reg_mod(dest) : _B40C_ASM_PTR_(d_ptr + offset));														\
			}																																						\
		};																																							\
		template <> struct ModifiedLoad<base_type, CA> {																												\
			__device__ __forceinline__ static void Ld(dest_type &dest, base_type* d_ptr, int offset) {																\
				asm("ld.global.ca."#ptx_type" %0, [%1];" : "="#reg_mod(dest) : _B40C_ASM_PTR_(d_ptr + offset));														\
			}																																						\
		};																																							\
		template <> struct ModifiedLoad<short_type##1, CG> {																											\
			__device__ __forceinline__ static void Ld(short_type##1 &dest, short_type##1* d_ptr, int offset) {														\
				asm("ld.global.cg."#ptx_type" %0, [%1];" : "="#reg_mod(dest.x) : _B40C_ASM_PTR_(d_ptr + offset));														\
			}																																						\
		};																																							\
		template <> struct ModifiedLoad<short_type##1, CS> {																											\
			__device__ __forceinline__ static void Ld(short_type##1 &dest, short_type##1* d_ptr, int offset) {														\
				asm("ld.global.cs."#ptx_type" %0, [%1];" : "="#reg_mod(dest.x) : _B40C_ASM_PTR_(d_ptr + offset));														\
			}																																						\
		};																																							\
		template <> struct ModifiedLoad<short_type##1, CA> {																											\
			__device__ __forceinline__ static void Ld(short_type##1 &dest, short_type##1* d_ptr, int offset) {														\
				asm("ld.global.ca."#ptx_type" %0, [%1];" : "="#reg_mod(dest.x) : _B40C_ASM_PTR_(d_ptr + offset));														\
			}																																						\
		};																																							\
		template <> struct ModifiedLoad<short_type##2, CG> {																											\
			__device__ __forceinline__ static void Ld(short_type##2 &dest, short_type##2* d_ptr, int offset) {														\
				asm("ld.global.cg.v2."#ptx_type" {%0, %1}, [%2];" : "="#reg_mod(dest.x), "="#reg_mod(dest.y) : _B40C_ASM_PTR_(d_ptr + offset));						\
			}																																						\
		};																																							\
		template <> struct ModifiedLoad<short_type##2, CS> {																											\
			__device__ __forceinline__ static void Ld(short_type##2 &dest, short_type##2* d_ptr, int offset) {														\
				asm("ld.global.cs.v2."#ptx_type" {%0, %1}, [%2];" : "="#reg_mod(dest.x), "="#reg_mod(dest.y) : _B40C_ASM_PTR_(d_ptr + offset));						\
			}																																						\
		};																																							\
		template <> struct ModifiedLoad<short_type##2, CA> {																											\
			__device__ __forceinline__ static void Ld(short_type##2 &dest, short_type##2* d_ptr, int offset) {														\
				asm("ld.global.ca.v2."#ptx_type" {%0, %1}, [%2];" : "="#reg_mod(dest.x), "="#reg_mod(dest.y) : _B40C_ASM_PTR_(d_ptr + offset));						\
			}																																						\
		};																																							\
		template <> struct ModifiedLoad<short_type##4, CG> {																											\
			__device__ __forceinline__ static void Ld(short_type##4 &dest, short_type##4* d_ptr, int offset) {														\
				asm("ld.global.cg.v4."#ptx_type" {%0, %1, %2, %3}, [%4];" : "="#reg_mod(dest.x), "="#reg_mod(dest.y), "="#reg_mod(dest.z), "="#reg_mod(dest.w) : _B40C_ASM_PTR_(d_ptr + offset));	\
			}																																						\
		};																																							\
		template <> struct ModifiedLoad<short_type##4, CS> {																											\
			__device__ __forceinline__ static void Ld(short_type##4 &dest, short_type##4* d_ptr, int offset) {														\
				asm("ld.global.cs.v4."#ptx_type" {%0, %1, %2, %3}, [%4];" : "="#reg_mod(dest.x), "="#reg_mod(dest.y), "="#reg_mod(dest.z), "="#reg_mod(dest.w) : _B40C_ASM_PTR_(d_ptr + offset));	\
			}																																						\
		};																																							\
		template <> struct ModifiedLoad<short_type##4, CA> {																											\
			__device__ __forceinline__ static void Ld(short_type##4 &dest, short_type##4* d_ptr, int offset) {														\
				asm("ld.global.ca.v4."#ptx_type" {%0, %1, %2, %3}, [%4];" : "="#reg_mod(dest.x), "="#reg_mod(dest.y), "="#reg_mod(dest.z), "="#reg_mod(dest.w) : _B40C_ASM_PTR_(d_ptr + offset));	\
			}																																						\
		};


	 */
	#define B40C_DEFINE_BASE_GLOBAL_LOAD(base_type, ptx_type, reg_mod)																								\
		template <> struct ModifiedLoad<base_type, CG> {																												\
			__device__ __forceinline__ static void Ld(base_type &val, base_type* d_ptr, int offset) {																\
				asm("ld.global.cg."#ptx_type" %0, [%1];" : "="#reg_mod(val) : _B40C_ASM_PTR_(d_ptr + offset));														\
			}																																						\
		};																																							\
		template <> struct ModifiedLoad<base_type, CS> {																												\
			__device__ __forceinline__ static void Ld(base_type &val, base_type* d_ptr, int offset) {																\
				asm("ld.global.cs."#ptx_type" %0, [%1];" : "="#reg_mod(val) : _B40C_ASM_PTR_(d_ptr + offset));														\
			}																																						\
		};																																							\
		template <> struct ModifiedLoad<base_type, CA> {																												\
			__device__ __forceinline__ static void Ld(base_type &val, base_type* d_ptr, int offset) {																\
				asm("ld.global.ca."#ptx_type" %0, [%1];" : "="#reg_mod(val) : _B40C_ASM_PTR_(d_ptr + offset));														\
			}																																						\
		};																																								

