inline __device__ void atomic_add()

in candle-flash-attn-v1/kernels/fmha/gmem_tile.h [259:277]

• 18 lines of code
• 7 McCabe index (conditional complexity)

    inline __device__ void atomic_add(const uint4 (&src)[STGS_PER_LOOP], int mi) {
        static_assert(BYTES_PER_ELEMENT == 4);  // Only do atomic add on floats
        int row_ = tidx_ / THREADS_PER_ROW;
        #pragma unroll
        for( int ii = 0; ii < STGS_PER_LOOP; ++ii ) {
            int jj = mi * STGS_PER_LOOP + ii;
            if ((!col_predicate) || (row_ + jj * ROWS_PER_STG >= this->actual_seqlen_q)) {
                break;
            }

            if( !HAS_INCOMPLETE_STG || (jj < STGS - 1 || this->is_active_for_last_stg_) ) {
                float *ptr_ = reinterpret_cast<float *>(this->ptr_ + jj * ROWS_PER_STG * this->row_stride_in_bytes);
                #pragma unroll
                for (int jj = 0; jj < 4; ++jj) {
                    atomicAdd(ptr_ + jj, reinterpret_cast<const float(&)[4]>(src[ii])[jj]);
                }
            }
        }
    }