ptrendx commented on a change in pull request #7789: Fix Moderngpu usages in
MXNet for CUDA 9
URL: https://github.com/apache/incubator-mxnet/pull/7789#discussion_r137602633
##########
File path:
src/operator/contrib/ctc_include/contrib/moderngpu/include/device/intrinsics.cuh
##########
@@ -104,32 +104,40 @@ MGPU_DEVICE uint prmt_ptx(uint a, uint b, uint index) {
#endif // __CUDA_ARCH__ >= 200
-
+#if CUDA_VERSION >= 9000
////////////////////////////////////////////////////////////////////////////////
-// shfl_up
-
-__device__ __forceinline__ float shfl_up(float var,
- unsigned int delta, int width = 32) {
+// shfl_add
+MGPU_DEVICE int shfl_add(int x, int offset, int width = WARP_SIZE, unsigned
int threadmask = 0xFFFFFFFF) {
+ int result = 0;
#if __CUDA_ARCH__ >= 300
- var = __shfl_up(var, delta, width);
+ int mask = (WARP_SIZE - width)<< 8;
+ asm(
+ "{.reg .s32 r0;"
Review comment:
It works right now (this diff view makes it difficult to see but it is
actually very minor change to the code that is just after that).
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
[email protected]
With regards,
Apache Git Services
