<div dir="ltr">
<div dir="ltr"><div>Hi all,<br></div><div><br></div><div>redux sync intrinsics are not working as expected.</div><div><br></div><div>clang/test/<a href="http://redux-builtins.cu" target="_blank">redux-builtins.cu</a> has usage as</div><div><span style="font-family:monospace"><br></span></div><div><span style="font-family:monospace">out = __nvvm_redux_sync_add(val, 0xFF);</span></div><div><br></div><div>out is the write location for the warp,</div><div>val is the thread's contributed value,</div><div>and 0xFF is the mask for a fully active warp.</div><div><br></div><div>So far all usage of this builtin has resulted in an Illegal instruction.</div><div>This is an nvcc application using the nvcc builtin to reduce across a warp:</div><div><span style="font-family:monospace"><br></span></div><div><span style="font-family:monospace">#include "cuda_runtime.h"<br>#include "device_launch_parameters.h"<br>#include <cuda_profiler_api.h><br>#include <stdio.h><br>#include <stdlib.h><br>#include <cuda_runtime.h><br><br>__inline__ __device__<br>unsigned warpReduceRedux(unsigned val) {<br>    return __reduce_add_sync(0xFF, val);<br>}<br><br>__global__ void reduceKernel(unsigned *in, unsigned* out, int N) {<br>    unsigned sum = in[threadIdx.x];<br>    sum = warpReduceRedux(sum);<br>    if (threadIdx.x == 0)<br>        out[0] = sum;<br>}<br><br>int main()<br>{<br>    const int SIZE = 32;<br>    const int ARRAY_BYTES = SIZE * sizeof(unsigned);<br><br>    // generate the input array on the host<br>    unsigned h_in[SIZE];<br>    unsigned sum = 0.0f;<br>    for (int i = 0; i < SIZE; i++) {<br>        h_in[i] = i;<br>        sum += h_in[i];<br>    }<br><br>    // declare GPU memory pointers<br>    unsigned * d_in, *d_out;<br><br>    // allocate GPU memory<br>    cudaMalloc((void **)&d_in, ARRAY_BYTES);<br>    cudaMalloc((void **)&d_out, sizeof(unsigned));<br><br>    // transfer the input array to the GPU<br>    cudaMemcpy(d_in, h_in, ARRAY_BYTES, cudaMemcpyHostToDevice);<br><br>    // offload to device<br>    reduceKernel<<<1, SIZE>>>(d_in, d_out, SIZE);<br><br>    // copy back the sum from GPU<br>    unsigned h_out;<br>    cudaMemcpy(&h_out, d_out, sizeof(unsigned), cudaMemcpyDeviceToHost);<br>    printf("%u\n", h_out);<br>}</span></div><div><br></div><div>cuda-memcheck is clear and has verifiable output.</div><div>The same application, substituting the nvcc builtin for the clang one then building with clang:</div><div><span style="font-family:monospace"><br></span></div><div><span style="font-family:monospace">__inline__ __device__<br>unsigned warpReduceRedux(unsigned val) {<br>    return __nvvm_redux_sync_add(val, 0xFF);<br>}</span><br></div><div><br></div><div> compiles but does not pass cuda-memcheck and does not provide the correct output:</div><div><br></div><div><span style="font-family:monospace">========= CUDA-MEMCHECK<br>========= Illegal Instruction<br>========= at 0x00000cf0 in reduceKernel(unsigned int*, unsigned int*, int)<br>========= by thread (0,0,0) in block (0,0,0)</span><br><br></div><div>What is the usage for these? I've also attached the PTX emitted by these apps in case there's a backend issue to be found.<br><div class="gmail-adL"><br></div></div></div>

</div>