@@ -102,9 +102,12 @@ __device__ __forceinline__ ReadOnlyDeviceMacroProperty<T, I, J, K, W> ReadOnlyDe
102
102
if (!d_ptr) {
103
103
return ReadOnlyDeviceMacroProperty<T, I, J, K, W>(nullptr , nullptr );
104
104
}
105
- #endif
106
105
return ReadOnlyDeviceMacroProperty<T, I, J, K, W>(reinterpret_cast <T*>(d_ptr),
107
106
reinterpret_cast <unsigned int *>(d_ptr + (I * J * K * W * sizeof (T)))); // Read-write flag resides in 8 bits at the end of the buffer
107
+ #else
108
+
109
+ return ReadOnlyDeviceMacroProperty<T, I, J, K, W>(reinterpret_cast <T*>(d_ptr));
110
+ #endif
108
111
}
109
112
template <typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W, unsigned int N>
110
113
__device__ __forceinline__ DeviceMacroProperty<T, I, J, K, W> DeviceEnvironment::getMacroProperty (const char (&name)[N]) const {
@@ -113,9 +116,11 @@ __device__ __forceinline__ DeviceMacroProperty<T, I, J, K, W> DeviceEnvironment:
113
116
if (!d_ptr) {
114
117
return DeviceMacroProperty<T, I, J, K, W>(nullptr , nullptr );
115
118
}
116
- #endif
117
119
return DeviceMacroProperty<T, I, J, K, W>(reinterpret_cast <T*>(d_ptr),
118
120
reinterpret_cast <unsigned int *>(d_ptr + (I * J * K * W * sizeof (T)))); // Read-write flag resides in 8 bits at the end of the buffer
121
+ #else
122
+ return DeviceMacroProperty<T, I, J, K, W>(reinterpret_cast <T*>(d_ptr));
123
+ #endif
119
124
}
120
125
#endif // __CUDACC_RTC__
121
126
0 commit comments