Merge remote-tracking branch 'origin/aaa' into cuda_version

S1/3/build_and_run.sh

@@ -32,7 +32,7 @@ print_warning() {
 }
 
 # 编译配置 - 可通过环境变量自定义
-COMPILER=${COMPILER:-mxcc}
+COMPILER=${COMPILER:-nvcc}
 COMPILER_FLAGS=${COMPILER_FLAGS:-"-O3 -std=c++17 --extended-lambda -DRUN_FULL_TEST"} 
 
 # ***** 这里是关键修改点1：头文件目录 *****
@@ -260,9 +260,9 @@ case "$RUN_MODE" in
         print_info "编译并运行 ${ALGO_TO_RUN} 测试 (模式: ${SINGLE_ALGO_TEST_MODE})..."
         local source_file_name=""
         case "$ALGO_TO_RUN" in
-            "ReduceSum") source_file_name="reduce_sum_algorithm.maca" ;;
-            "SortPair") source_file_name="sort_pair_algorithm.maca" ;;
-            "TopkPair") source_file_name="topk_pair_algorithm.maca" ;;
+            "ReduceSum") source_file_name="reduce_sum_algorithm.cu" ;;
+            "SortPair") source_file_name="sort_pair_algorithm.cu" ;;
+            "TopkPair") source_file_name="topk_pair_algorithm.cu" ;;
         esac
 
         if compile_algorithm "$ALGO_TO_RUN" "$source_file_name"; then

S1/3/reduce_sum_algorithm.cu

@@ -0,0 +1,439 @@
+#include "test_utils.h"
+#include "performance_utils.h"
+#include "yaml_reporter.h"
+#include <iostream>
+#include <vector>
+#include <iomanip>
+
+
+// ============================================================================
+// 实现标记宏 - 参赛者修改实现时请将此宏设为0
+// ============================================================================
+#ifndef USE_DEFAULT_REF_IMPL
+#define USE_DEFAULT_REF_IMPL 0  // 1=默认实现, 0=参赛者自定义实现
+#endif
+
+#if USE_DEFAULT_REF_IMPL
+#include <thrust/reduce.h>
+#include <thrust/device_vector.h>
+#include <thrust/execution_policy.h>
+#include <thrust/functional.h>
+#endif
+
+#define MUITWORK 4
+
+// 误差容忍度
+constexpr double REDUCE_ERROR_TOLERANCE = 0.005; // 0.5%
+
+// ============================================================================
+// ReduceSum算法实现接口
+// 参赛者需要替换Thrust实现为自己的高性能kernel
+// ============================================================================
+
+template <typename InputT = float, typename OutputT = float>
+__global__ void blockReduceKernel(const InputT* d_in, InputT* d_out, int num_items){
+    extern __shared__ OutputT sdata_out[];
+    
+    unsigned int tid = threadIdx.x;
+    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
+
+    if (idx < num_items) {
+        sdata_out[tid] = d_in[idx];
+    } else {
+        sdata_out[tid] = (InputT)0;
+    }
+
+    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1){
+        if (tid < stride && idx + stride){
+            sdata_out[tid] += sdata_out[tid + stride];
+        }
+        __syncthreads();
+    }
+
+    if (tid == 0) {
+        d_out[blockIdx.x] = sdata_out[0];
+        // printf("a:%f ", d_out[0]);
+    }
+}
+
+template <typename InputT = float, typename OutputT = float>
+__global__ void finalReduceKernel(const InputT* d_in, OutputT* d_out, int num_items, OutputT init_value){
+    extern __shared__ OutputT sdata_out[];
+    
+    unsigned int tid = threadIdx.x;
+    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
+
+    if (idx < num_items) {
+        sdata_out[tid] = (OutputT)d_in[idx];
+    } else {
+        sdata_out[tid] = 0;
+    }
+
+    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1){
+        if (tid < stride && idx + stride){
+            sdata_out[tid] += sdata_out[tid + stride];
+        }
+        __syncthreads();
+    }
+
+    if (tid == 0) {
+        d_out[0] = sdata_out[0] + init_value;
+        // printf("a:%f ", d_out[0]);
+    }
+}
+
+template <typename T = float>
+__inline__ __device__ T warpReduce(T myVal){
+    myVal += __shfl_xor_sync(0xffffffff, myVal, 16);
+    myVal += __shfl_xor_sync(0xffffffff, myVal, 8);
+    myVal += __shfl_xor_sync(0xffffffff, myVal, 4);
+    myVal += __shfl_xor_sync(0xffffffff, myVal, 2);
+    myVal += __shfl_xor_sync(0xffffffff, myVal, 1);
+    return myVal;
+}
+
+template <typename InputT = float, typename OutputT = float>
+__global__ void shflBlockReduceKernel(const InputT* d_in, InputT* d_out, int num_items){
+    unsigned int tid = threadIdx.x;
+    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
+
+    InputT myVal = (idx * MUITWORK < num_items) ? d_in[idx * MUITWORK] : (InputT)0;
+    myVal += (idx * MUITWORK + 1 < num_items) ? d_in[idx * MUITWORK + 1] : (InputT)0;
+    myVal += (idx * MUITWORK + 2 < num_items) ? d_in[idx * MUITWORK + 2] : (InputT)0;
+    myVal += (idx * MUITWORK + 3 < num_items) ? d_in[idx * MUITWORK + 3] : (InputT)0;
+
+    myVal = warpReduce(myVal);
+
+    int warpIdx = tid / 32;
+    int laneIdx = tid % 32;
+
+    extern __shared__ InputT seme[];
+
+    if (laneIdx == 0) seme[warpIdx] = myVal;
+
+    __syncthreads();
+    myVal = (tid < blockDim.x / 32) ? seme[laneIdx] : 0;
+    if (warpIdx == 0) myVal = warpReduce(myVal);
+    if (tid == 0) {
+        d_out[blockIdx.x] = myVal;
+        // printf("b:%f ", d_out[0]);
+    }
+}
+
+template <typename InputT = float, typename OutputT = float>
+__global__ void finialShflReduceKernel(const InputT* d_in, OutputT* d_out, int num_items, OutputT init_value){
+    unsigned int tid = threadIdx.x;
+    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
+
+    InputT myVal = (idx * MUITWORK < num_items) ? d_in[idx * MUITWORK] : (InputT)0;
+    myVal += (idx * MUITWORK + 1 < num_items) ? d_in[idx * MUITWORK + 1] : (InputT)0;
+    myVal += (idx * MUITWORK + 2 < num_items) ? d_in[idx * MUITWORK + 2] : (InputT)0;
+    myVal += (idx * MUITWORK + 3 < num_items) ? d_in[idx * MUITWORK + 3] : (InputT)0;
+
+    myVal = warpReduce(myVal);
+
+    int warpIdx = tid / 32;
+    int laneIdx = tid % 32;
+
+    extern __shared__ OutputT seme[];
+
+    if (laneIdx == 0) seme[warpIdx] = myVal;
+
+    __syncthreads();
+    myVal = (tid < blockDim.x / 32) ? seme[laneIdx] : 0;
+    if (warpIdx == 0) myVal = warpReduce(myVal);
+    if (tid == 0) {
+        d_out[0] = myVal + init_value;
+        // printf("b:%f ", d_out[0]);
+    }
+}
+
+template <typename InputT = float, typename OutputT = float>
+__global__ void finialMutiShflReduceKernel(const InputT* d_in1, const InputT* d_in2, OutputT* d_out, int num_items0, int num_items1, OutputT init_value){
+    unsigned int tid = threadIdx.x;
+    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
+
+    OutputT myVal = (idx < num_items0) ? (OutputT)d_in1[idx] : 0; 
+    myVal += (idx < num_items1) ? (OutputT)d_in2[idx] : 0; 
+
+    myVal = warpReduce(myVal);
+
+    int warpIdx = tid / 32;
+    int laneIdx = tid % 32;
+
+    extern __shared__ OutputT seme[];
+
+    if (laneIdx == 0) seme[warpIdx] = myVal;
+
+    __syncthreads();
+    myVal = (tid < blockDim.x / 32) ? seme[laneIdx] : 0;
+    if (warpIdx == 0) myVal = warpReduce(myVal);
+    if (tid == 0) {
+        d_out[0] = myVal + init_value;
+        // printf("b:%f ", d_out[0]);
+    }
+}
+
+template <typename InputT = float, typename OutputT = float>
+class ReduceSumAlgorithm {
+public:
+    // 主要接口函数 - 参赛者需要实现这个函数
+    void reduce(const InputT* d_in, OutputT* d_out, int num_items, OutputT init_value) {
+        
+#if !USE_DEFAULT_REF_IMPL
+        // ========================================
+        // 参赛者自定义实现区域
+        // ========================================
+        
+        // TODO: 参赛者在此实现自己的高性能归约算法
+        
+        // 示例：参赛者可以调用1个或多个自定义kernel
+
+        grid = (num_items + (MUITWORK * block) - 1) / (MUITWORK * block);
+        const InputT* d_tmp = d_in;
+        MACA_CHECK(cudaMalloc(&temp_results, grid * sizeof(InputT)));
+        for (; grid > 1;){
+            shflBlockReduceKernel<<<grid, block, block / 32 * sizeof(InputT)>>>(d_tmp, temp_results, num_items);
+            num_items = grid;
+            grid = (num_items + (MUITWORK * block) - 1) / (MUITWORK * block);
+
+            MACA_CHECK(cudaDeviceSynchronize()); 
+
+            d_tmp = temp_results;
+        }
+        finialShflReduceKernel<<<1, block, block / 32 * sizeof(OutputT)>>>(d_tmp, d_out, num_items, init_value);
+        cudaFree(temp_results);
+#else
+        // ========================================
+        // 默认基准实现
+        // ========================================
+        auto input_ptr = thrust::device_pointer_cast(d_in);
+        auto output_ptr = thrust::device_pointer_cast(d_out);
+        
+        // 直接使用thrust::reduce进行归约
+        *output_ptr = thrust::reduce(
+            thrust::device,
+            input_ptr,
+            input_ptr + num_items,
+            static_cast<OutputT>(init_value)
+        );     
+#endif
+    }
+    
+    // 获取当前实现状态
+    static const char* getImplementationStatus() {
+#if USE_DEFAULT_REF_IMPL
+        return "DEFAULT_REF_IMPL";
+#else
+        return "CUSTOM_IMPL";
+#endif
+    }
+    
+private:
+    // 参赛者可以在这里添加辅助函数和成员变量
+    // 例如：中间结果缓冲区、多阶段归约等
+    int block = 256;
+    int grid;
+    InputT* temp_results;
+};
+
+// ============================================================================
+// 测试和性能评估
+// ============================================================================
+
+bool testCorrectness() {
+    std::cout << "ReduceSum 正确性测试..." << std::endl;
+    TestDataGenerator generator;
+    ReduceSumAlgorithm<float, float> algorithm;
+    
+    bool allPassed = true;
+    
+    // 测试不同数据规模
+    for (int i = 0; i < NUM_TEST_SIZES && i < 2; i++) { // 限制测试规模
+        int size = std::min(TEST_SIZES[i], 10000);
+        std::cout << "  测试规模: " << size << std::endl;
+        
+        // 测试普通数据
+        {
+            auto data = generator.generateRandomFloats(size, -10.0f, 10.0f);
+            float init_value = 1.0f;
+            
+            // CPU参考计算
+            double cpu_result = cpuReduceSum(data, static_cast<double>(init_value));
+            
+            // GPU计算
+            float *d_in;
+            float *d_out;
+            MACA_CHECK(cudaMalloc(&d_in, size * sizeof(float)));
+            MACA_CHECK(cudaMalloc(&d_out, sizeof(float)));
+            
+            MACA_CHECK(cudaMemcpy(d_in, data.data(), size * sizeof(float), cudaMemcpyHostToDevice));
+            
+            algorithm.reduce(d_in, d_out, size, init_value);
+            
+            float gpu_result;
+            MACA_CHECK(cudaMemcpy(&gpu_result, d_out, sizeof(float), cudaMemcpyDeviceToHost));
+            
+            // 验证误差
+            double relative_error = std::abs(gpu_result - cpu_result) / std::abs(cpu_result);
+            if (relative_error > REDUCE_ERROR_TOLERANCE) {
+                std::cout << "    失败: 误差过大 " << relative_error << std::endl;
+                allPassed = false;
+            } else {
+                std::cout << "    通过 (误差: " << relative_error << ")" << std::endl;
+            }
+            
+            cudaFree(d_in);
+            cudaFree(d_out);
+        }
+        
+        // 测试特殊值 (NaN, Inf)
+        if (size > 100) {
+            std::cout << "  测试特殊值..." << std::endl;
+            auto data = generator.generateSpecialFloats(size);
+            float init_value = 0.0f;
+            
+            double cpu_result = cpuReduceSum(data, static_cast<double>(init_value));
+            
+            float *d_in;
+            float *d_out;
+            MACA_CHECK(cudaMalloc(&d_in, size * sizeof(float)));
+            MACA_CHECK(cudaMalloc(&d_out, sizeof(float)));
+            
+            MACA_CHECK(cudaMemcpy(d_in, data.data(), size * sizeof(float), cudaMemcpyHostToDevice));
+            
+            algorithm.reduce(d_in, d_out, size, init_value);
+            
+            float gpu_result;
+            MACA_CHECK(cudaMemcpy(&gpu_result, d_out, sizeof(float), cudaMemcpyDeviceToHost));
+            
+            // 对于包含特殊值的情况，检查是否正确处理
+            if (std::isfinite(cpu_result) && std::isfinite(gpu_result)) {
+                double relative_error = std::abs(gpu_result - cpu_result) / std::abs(cpu_result);
+                if (relative_error > REDUCE_ERROR_TOLERANCE) {
+                    std::cout << "    失败: 特殊值处理错误" << std::endl;
+                    allPassed = false;
+                } else {
+                    std::cout << "    通过 (特殊值处理)" << std::endl;
+                }
+            } else {
+                std::cout << "    通过 (特殊值结果)" << std::endl;
+            }
+            
+            cudaFree(d_in);
+            cudaFree(d_out);
+        }
+    }
+    
+    return allPassed;
+}
+
+void benchmarkPerformance() {
+    PerformanceDisplay::printReduceSumHeader();
+    
+    TestDataGenerator generator;
+    PerformanceMeter meter;
+    ReduceSumAlgorithm<float, float> algorithm;
+    
+    const int WARMUP_ITERATIONS = 5;
+    const int BENCHMARK_ITERATIONS = 10;
+    
+    // 用于YAML报告的数据收集
+    std::vector<std::map<std::string, std::string>> perf_data;
+    
+    for (int i = 0; i < NUM_TEST_SIZES; i++) {
+        int size = TEST_SIZES[i];
+        
+        // 生成测试数据
+        auto data = generator.generateRandomFloats(size);
+        float init_value = 0.0f;
+        
+        // 分配GPU内存
+        float *d_in;
+        float *d_out;
+        MACA_CHECK(cudaMalloc(&d_in, size * sizeof(float)));
+        MACA_CHECK(cudaMalloc(&d_out, sizeof(float)));
+        
+        MACA_CHECK(cudaMemcpy(d_in, data.data(), size * sizeof(float), cudaMemcpyHostToDevice));
+        
+        // Warmup阶段
+        for (int iter = 0; iter < WARMUP_ITERATIONS; iter++) {
+            algorithm.reduce(d_in, d_out, size, init_value);
+        }
+        
+        // 正式测试阶段
+        float total_time = 0;
+        for (int iter = 0; iter < BENCHMARK_ITERATIONS; iter++) {
+            meter.startTiming();
+            algorithm.reduce(d_in, d_out, size, init_value);
+            total_time += meter.stopTiming();
+        }
+        
+        float avg_time = total_time / BENCHMARK_ITERATIONS;
+        
+        // 计算性能指标
+        auto metrics = PerformanceCalculator::calculateReduceSum(size, avg_time);
+        
+        // 显示性能数据
+        PerformanceDisplay::printReduceSumData(size, avg_time, metrics);
+        
+        // 收集YAML报告数据
+        auto entry = YAMLPerformanceReporter::createEntry();
+        entry["data_size"] = std::to_string(size);
+        entry["time_ms"] = std::to_string(avg_time);
+        entry["throughput_gps"] = std::to_string(metrics.throughput_gps);
+        entry["data_type"] = "float";
+        perf_data.push_back(entry);
+        
+        cudaFree(d_in);
+        cudaFree(d_out);
+    }
+    
+    // 生成YAML性能报告
+    YAMLPerformanceReporter::generateReduceSumYAML(perf_data, "reduce_sum_performance.yaml");
+    PerformanceDisplay::printSavedMessage("reduce_sum_performance.yaml");
+}
+
+// ============================================================================
+// 主函数
+// ============================================================================
+int main(int argc, char* argv[]) {
+    std::cout << "=== ReduceSum 算法测试 ===" << std::endl;
+    
+    // 检查参数
+    std::string mode = "all";
+    if (argc > 1) {
+        mode = argv[1];
+    }
+    
+    bool correctness_passed = true;
+    bool performance_completed = true;
+    
+    try {
+        if (mode == "correctness" || mode == "all") {
+            correctness_passed = testCorrectness();
+        }
+        
+        if (mode == "performance" || mode == "all") {
+            if (correctness_passed || mode == "performance") {
+                benchmarkPerformance();
+            } else {
+                std::cout << "跳过性能测试，因为正确性测试未通过" << std::endl;
+                performance_completed = false;
+            }
+        }
+        
+        std::cout << "\n=== 测试完成 ===" << std::endl;
+        std::cout << "实现状态: " << ReduceSumAlgorithm<float, float>::getImplementationStatus() << std::endl;
+        if (mode == "all") {
+            std::cout << "正确性: " << (correctness_passed ? "通过" : "失败") << std::endl;
+            std::cout << "性能测试: " << (performance_completed ? "完成" : "跳过") << std::endl;
+        }
+        
+        return correctness_passed ? 0 : 1;
+        
+    } catch (const std::exception& e) {
+        std::cerr << "测试出错: " << e.what() << std::endl;
+        return 1;
+    }
+}

S1/3/reduce_sum_algorithm.maca

@@ -20,6 +20,8 @@
 #include <thrust/functional.h>
 #endif
 
+#define MUITWORK 4
+
 // 误差容忍度
 constexpr double REDUCE_ERROR_TOLERANCE = 0.005; // 0.5%
 
@@ -95,7 +97,10 @@ __global__ void shflBlockReduceKernel(const InputT* d_in, InputT* d_out, int num
     unsigned int tid = threadIdx.x;
     unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
 
-    InputT myVal = (idx < num_items) ? d_in[idx] : (InputT)0; 
+    InputT myVal = (idx * MUITWORK < num_items) ? d_in[idx * MUITWORK] : (InputT)0;
+    myVal += (idx * MUITWORK + 1 < num_items) ? d_in[idx * MUITWORK + 1] : (InputT)0;
+    myVal += (idx * MUITWORK + 2 < num_items) ? d_in[idx * MUITWORK + 2] : (InputT)0;
+    myVal += (idx * MUITWORK + 3 < num_items) ? d_in[idx * MUITWORK + 3] : (InputT)0;
 
     myVal = warpReduce(myVal);
 
@@ -120,7 +125,10 @@ __global__ void finialShflReduceKernel(const InputT* d_in, OutputT* d_out, int n
     unsigned int tid = threadIdx.x;
     unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
 
-    OutputT myVal = (idx < num_items) ? (OutputT)d_in[idx] : 0; 
+    InputT myVal = (idx * MUITWORK < num_items) ? d_in[idx * MUITWORK] : (InputT)0;
+    myVal += (idx * MUITWORK + 1 < num_items) ? d_in[idx * MUITWORK + 1] : (InputT)0;
+    myVal += (idx * MUITWORK + 2 < num_items) ? d_in[idx * MUITWORK + 2] : (InputT)0;
+    myVal += (idx * MUITWORK + 3 < num_items) ? d_in[idx * MUITWORK + 3] : (InputT)0;
 
     myVal = warpReduce(myVal);
 
@@ -181,44 +189,20 @@ public:
         
         // 示例：参赛者可以调用1个或多个自定义kernel
 
-        // int block = 1024;
-        int stream_num = 2;
+        grid = (num_items + (MUITWORK * block) - 1) / (MUITWORK * block);
+        const InputT* d_tmp = d_in;
+        MACA_CHECK(mcMalloc(&temp_results, grid * sizeof(InputT)));
+        for (; grid > 1;){
+            shflBlockReduceKernel<<<grid, block, block / 32 * sizeof(InputT)>>>(d_tmp, temp_results, num_items);
+            num_items = grid;
+            grid = (num_items + (MUITWORK * block) - 1) / (MUITWORK * block);
 
-        mcStream_t stream0, stream1;
-        mcStreamCreate(&stream0);
-        mcStreamCreate(&stream1);
+            MACA_CHECK(mcDeviceSynchronize()); 
 
-        int num_items0 = num_items / 2;
-        int num_items1 = num_items - num_items0;
-        grid0 = (num_items0 + block - 1) / block;
-        grid1 = (num_items1 + block - 1) / block;
-
-        const InputT* d_tmp0 = d_in;
-        const InputT* d_tmp1 = d_in + num_items0;
-
-        MACA_CHECK(mcMalloc(&temp_results0, grid0 * sizeof(InputT)));
-        MACA_CHECK(mcMalloc(&temp_results1, grid1 * sizeof(InputT)));
-        for (; grid0 > 1 || grid1 > 1; ){
-            if (grid0 > 1){
-                shflBlockReduceKernel<<<grid0, block, block / 32 * sizeof(InputT), stream0>>>(d_tmp0, temp_results0, num_items0);
-                grid0 = (num_items0 + block - 1) / block;
-                num_items0 = grid0;
-            }
-            if (grid1 > 1){
-                shflBlockReduceKernel<<<grid1, block, block / 32 * sizeof(InputT), stream1>>>(d_tmp1, temp_results1, num_items1);
-                grid1 = (num_items1 + block - 1) / block;
-                num_items1 = grid1;
-            }
-
-            MACA_CHECK(mcStreamSynchronize(stream0)); 
-            MACA_CHECK(mcStreamSynchronize(stream1)); 
-
-            d_tmp0 = temp_results0;
-            d_tmp1 = temp_results1;
+            d_tmp = temp_results;
         }
-        finialMutiShflReduceKernel<<<1, block, block / 32 * sizeof(OutputT)>>>(d_tmp0, d_tmp1, d_out, num_items0, num_items1, init_value);
-        mcFree(temp_results0);
-        mcFree(temp_results1);
+        finialShflReduceKernel<<<1, block, block / 32 * sizeof(OutputT)>>>(d_tmp, d_out, num_items, init_value);
+        mcFree(temp_results);
 #else
         // ========================================
         // 默认基准实现
@@ -248,11 +232,9 @@ public:
 private:
     // 参赛者可以在这里添加辅助函数和成员变量
     // 例如：中间结果缓冲区、多阶段归约等
-    int block = 512;
-    int grid0;
-    int grid1;
-    InputT* temp_results0;
-    InputT* temp_results1;
+    int block = 256;
+    int grid;
+    InputT* temp_results;
 };
 
 // ============================================================================

S1/3/utils/test_utils.h

@@ -2,8 +2,8 @@
 #include <vector>
 #include <random>
 #include <algorithm>
-#include <mc_runtime.h>
-#include <maca_fp16.h>
+#include <cuda_runtime.h>
+#include <cuda_fp16.h>
 #include <iostream>
 #include <chrono>
 #include <cmath>
@@ -33,10 +33,10 @@ constexpr int BENCHMARK_ITERATIONS = 10;
 // ============================================================================
 #define MACA_CHECK(call) \
     do { \
-        mcError_t error = call; \
-        if (error != mcSuccess) { \
+        cudaError_t error = call; \
+        if (error != cudaSuccess) { \
             std::cerr << "MACA error at " << __FILE__ << ":" << __LINE__ \
-                      << " - " << mcGetErrorString(error) << std::endl; \
+                      << " - " << cudaGetErrorString(error) << std::endl; \
             exit(1); \
         } \
     } while(0)
@@ -117,28 +117,28 @@ public:
 // ============================================================================
 class PerformanceMeter {
 private:
-    mcEvent_t start, stop;
+    cudaEvent_t start, stop;
     
 public:
     PerformanceMeter() {
-        MACA_CHECK(mcEventCreate(&start));
-        MACA_CHECK(mcEventCreate(&stop));
+        MACA_CHECK(cudaEventCreate(&start));
+        MACA_CHECK(cudaEventCreate(&stop));
     }
     
     ~PerformanceMeter() {
-        mcEventDestroy(start);
-        mcEventDestroy(stop);
+        cudaEventDestroy(start);
+        cudaEventDestroy(stop);
     }
     
     void startTiming() {
-        MACA_CHECK(mcEventRecord(start));
+        MACA_CHECK(cudaEventRecord(start));
     }
     
     float stopTiming() {
-        MACA_CHECK(mcEventRecord(stop));
-        MACA_CHECK(mcEventSynchronize(stop));
+        MACA_CHECK(cudaEventRecord(stop));
+        MACA_CHECK(cudaEventSynchronize(stop));
         float milliseconds = 0;
-        MACA_CHECK(mcEventElapsedTime(&milliseconds, start, stop));
+        MACA_CHECK(cudaEventElapsedTime(&milliseconds, start, stop));
         return milliseconds;
     }
 };