SunnySun

#include <cooperative_groups.h>
#include <maca_async.h>
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <cstdio>
#include <cmath>

namespace cg = cooperative_groups;

constexpr int BLOCK_SIZE = 256;
constexpr int TILE_SIZE = 6144;      // 每个tile的元素数量
constexpr int NUM_TILES = 2048;         // 处理的tile数量
constexpr int COMPUTE_ITERATIONS = 1; // 每个tile的计算工作量

// compute func
__device__ float compute_kernel(float data, int iterations) {
    float result = data;
    for (int i = 0; i < iterations; ++i) {
        result = result * 0.99f + 0.01f;
    }
    return result;
}

__global__ void async_copy_pipeline(float* global_data, float* global_output) {
    auto cta = cg::this_thread_block();
    __shared__ float smem[2][TILE_SIZE];
    // calc offset
    const int cta_offset = blockIdx.x * NUM_TILES * TILE_SIZE;

    float *src = global_data + cta_offset;

    int read_buf  = 0;   // compute
    int write_buf = 1;   // copy

    // the first tile 异步拷贝
    cg::memcpy_async(cta, smem[read_buf], src, sizeof(float) * TILE_SIZE);
    cg::wait(cta);
    cta.sync();

    for (int tile_idx = 0; tile_idx < NUM_TILES; ++tile_idx) {
        if (tile_idx < NUM_TILES - 1) {
            // 发起next tile 的异步拷贝
            float* next_src = src + (tile_idx + 1) * TILE_SIZE;
            cg::memcpy_async(cta, smem[write_buf], next_src, sizeof(float) * TILE_SIZE);
        }
        // current tile compute
        for (int i = threadIdx.x; i < TILE_SIZE; i += blockDim.x) {
            smem[read_buf][i] = compute_kernel(smem[read_buf][i], COMPUTE_ITERATIONS);
        }

        if (tile_idx < NUM_TILES - 1) {
            cg::wait(cta);
        }
        cta.sync();

        float *dst = global_output + cta_offset + tile_idx * TILE_SIZE;
        for (int i = threadIdx.x; i < TILE_SIZE; i += blockDim.x) {
            dst[i] = smem[read_buf][i];
        }
        if (tile_idx < NUM_TILES - 1) {
            int temp = read_buf;
            read_buf = write_buf;
            write_buf = temp;
        }
    }
}

// 同步版本用于对比
__global__ void sync_copy_sequential(float* global_data, float* global_output) {
    __shared__ float smem[TILE_SIZE];

    float* src = global_data + blockIdx.x * NUM_TILES * TILE_SIZE;
    float* dst = global_output + blockIdx.x * NUM_TILES * TILE_SIZE;

    for (int tile_idx = 0; tile_idx < NUM_TILES; ++tile_idx) {
        float* current_src = src + tile_idx * TILE_SIZE;
        float* current_dst = dst + tile_idx * TILE_SIZE;

        //copy to smem
        for (int i = threadIdx.x; i < TILE_SIZE; i += blockDim.x) {
            smem[i] = current_src[i];
        }
        __syncthreads();

        for (int i = threadIdx.x; i < TILE_SIZE; i += blockDim.x) {
            smem[i] = compute_kernel(smem[i], COMPUTE_ITERATIONS);
        }
        __syncthreads();

        for (int i = threadIdx.x; i < TILE_SIZE; i += blockDim.x) {
            current_dst[i] = smem[i];
        }
        __syncthreads();
    }
}

int main() {
    cudaError_t cudaStatus = cudaSetDevice(0);
    const int num_blocks = 4;
    const int total_elements = num_blocks * NUM_TILES * TILE_SIZE;

    printf("Testing async copy pipeline with %d blocks, %d tiles per block, %d elements per tile\n",
           num_blocks, NUM_TILES, TILE_SIZE);

    thrust::host_vector<float> h_input(total_elements);
    thrust::host_vector<float> h_output_async(total_elements);
    thrust::host_vector<float> h_output_sync(total_elements);
    for (int j = 0; j < total_elements; ++j) h_output_async[j] = static_cast<float>(-1);
    for (int j = 0; j < total_elements; ++j) h_output_sync[j] = static_cast<float>(-1);

    for (int i = 0; i < total_elements; ++i) {
        h_input[i] = static_cast<float>(i) * 0.1f;
    }

    thrust::device_vector<float> d_input = h_input;
    thrust::device_vector<float> d_output_async = h_output_async;
    thrust::device_vector<float> d_output_sync = h_output_sync;

    // test async
    cudaEvent_t start_async, stop_async;
    cudaEventCreate(&start_async);
    cudaEventCreate(&stop_async);

    cudaEventRecord(start_async);
    async_copy_pipeline<<<num_blocks, BLOCK_SIZE>>>(
        d_input.data().get(),d_output_async.data().get()
    );
    cudaEventRecord(stop_async);
    cudaDeviceSynchronize();

    float async_ms = 0;
    cudaEventElapsedTime(&async_ms, start_async, stop_async);

    // test sync
    cudaEvent_t start_sync, stop_sync;
    cudaEventCreate(&start_sync);
    cudaEventCreate(&stop_sync);

    cudaEventRecord(start_sync);
    sync_copy_sequential<<<num_blocks, BLOCK_SIZE>>>(
        d_input.data().get(),d_output_sync.data().get()
    );
    cudaEventRecord(stop_sync);
    cudaDeviceSynchronize();

    float sync_ms = 0;
    cudaEventElapsedTime(&sync_ms, start_sync, stop_sync);

    // 验证结果一致性
    thrust::host_vector<float> h_result_async = d_output_async;
    thrust::host_vector<float> h_result_sync = d_output_sync;

    bool results_match = true;
    for (int i = 0; i < total_elements; ++i) {
        if (fabs(h_result_async[i] - h_result_sync[i]) > 1e-6) {
            results_match = false;
            break;
        }
    }

    printf("Results: %s\n", results_match ? "MATCH" : "MISMATCH");
    printf("Async pipeline time: %.3f ms\n", async_ms);
    printf("Sync sequential time: %.3f ms\n", sync_ms);
    printf("Speedup: %.2fx\n", sync_ms / async_ms);

    // clear
    cudaEventDestroy(start_async);
    cudaEventDestroy(stop_async);
    cudaEventDestroy(start_sync);
    cudaEventDestroy(stop_sync);

    return 0;
}

上面的代码是用于 测试异步拷贝加计算的运行时间 以及 同步拷贝加计算的运行时间，预期结果是使用异步拷贝性能优于同步拷贝，但是在MetaX C500上测试，结果如下（结果显示异步拷贝耗时更长）：

Testing async copy pipeline with 4 blocks, 2048 tiles per block, 6144 elements per tile
Results: MATCH
Async pipeline time: 30.342 ms
Sync sequential time: 9.460 ms
Speedup: 0.31x

同样的代码，在L20上测试，结果如下（L20上异步拷贝的耗时比同步短）：

Testing async copy pipeline with 4 blocks, 2048 tiles per block, 6144 elements per tile
Results: MATCH
Async pipeline time: 4.221 ms
Sync sequential time: 17.118 ms
Speedup: 4.06x

代码中用到了异步拷贝，主要参考developer.metax-tech.com/api/client/document/preview/559/C500_MXMACAC%2B%2BProgrammingGuide_CN.html#ehg6356t2io61 中说明的使用方法

问题：为什么在C500上异步拷贝的性能比同步的还要差呢？请问是异步拷贝的接口使用有问题还是什么原因呢？请指教，谢谢