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

	// Workaround for the fact that the assembler reports an error when attempting to 
	// make vector loads of doubles.
	template <> struct ModifiedLoad<double2, CG> {																												
		__device__ __forceinline__ static void Ld(double2 &dest, double2* d_ptr, int offset) {													
			asm("ld.global.cg.f64 %0, [%1];" : "=d"(dest.x) : _B40C_ASM_PTR_(d_ptr + offset));																	
			asm("ld.global.cg.f64 %0, [%1];" : "=d"(dest.y) : _B40C_ASM_PTR_(d_ptr + offset + 1));																	
		}																																							
	};																																								
	template <> struct ModifiedLoad<double4, CG> {																												
		__device__ __forceinline__ static void Ld(double4 &dest, double4* d_ptr, int offset) {													
			asm("ld.global.cg.f64 %0, [%1];" : "=d"(dest.x) : _B40C_ASM_PTR_(d_ptr + offset));																	
			asm("ld.global.cg.f64 %0, [%1];" : "=d"(dest.y) : _B40C_ASM_PTR_(d_ptr + offset + 1));																	
			asm("ld.global.cg.f64 %0, [%1];" : "=d"(dest.z) : _B40C_ASM_PTR_(d_ptr + offset + 2));																	
			asm("ld.global.cg.f64 %0, [%1];" : "=d"(dest.w) : _B40C_ASM_PTR_(d_ptr + offset + 3));																	
		}																																							
	};																																								
	
	// NONE-modified load 
	template <typename T> struct ModifiedLoad<T, NONE>
	{
		__device__ __forceinline__ static void Ld(T &dest, T* d_ptr, int offset) {
			dest = d_ptr[offset]; 
		}
	};
	
	// NONE-modified load 
	template <> struct ModifiedLoad<char, NONE>
	{
		__device__ __forceinline__ static void Ld(signed char &dest, char* d_ptr, int offset) {
			dest = d_ptr[offset]; 
		}
	};
	
#else 

	// Nothing is cached in these architectures: load normally
	template <typename T, CacheModifier CACHE_MODIFIER> struct ModifiedLoad
	{
		__device__ __forceinline__ static void Ld(T &dest, T* d_ptr, int offset) {
			dest = d_ptr[offset]; 
		}
	};
	
	// Accomodate bizarre introduction of "signed" for char loads
	template <CacheModifier CACHE_MODIFIER> struct ModifiedLoad<char, CACHE_MODIFIER>
	{
		__device__ __forceinline__ static void Ld(signed char &dest, char* d_ptr, int offset) {
			dest = d_ptr[offset]; 
		}
	};

#endif






	 */
	#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));														\
			}																																						\
		};																																								

