非临时缓存以及访存对计算性能的思考

作者: zenny_chen | 来源:发表于2018-07-19 17:10 被阅读123次

以下是基于OpenCL的测试代码：

#include <stdio.h>
#include <stdbool.h>
#include <string.h>
#include <stdlib.h>
#include <sys/time.h>

#ifdef __APPLE__
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif

#ifndef var
#define var     __auto_type
#endif

#define TEST_LOOP_COUNT     1000

static const char *sKernelCode =
"kernel void test(global int4 *pMemBuf)  \n"
"{  \n"
"    const uint itemID = get_global_id(0);  \n"
"    pMemBuf[itemID] += 1;  \n"
"}\n";

/// 以一般访存模式处理数据
/// @param pUserData 指向主机端用户数据的缓存
/// @param pDeviceData 指向设备数据的缓存
/// @param count 指定缓存中int数据元素的个数
/// @return 返回计算后的设备数据的某个整数值
extern int DoDataProcessNormal(int *pUserData, const int *pDeviceData, size_t count);

/// 以非临时不影响Cache的访存模式处理数据
/// @param pUserData 指向主机端用户数据的缓存
/// @param pDeviceData 指向设备数据的缓存
/// @param count 指定缓存中int数据元素的个数
/// @return 返回计算后的设备数据的某个整数值
extern int DoDataProcessNonTmp(int *pUserData, const int *pDeviceData, size_t count);

enum CalculateType
{
    CalculateType_NORMAL_SINGLE_CORE,
    CalculateType_NORMAL_DUO_CORE,
    CalculateType_NONTMP_SINGLE_CORE,
    CalculateType_NONTMP_DUO_CORE
};

static int NormalSingleCoreHandler(int *pUserData, const int *pDeviceData, size_t count)
{
    return DoDataProcessNormal(pUserData, pDeviceData, count);
}

static int NormalDuoCoreHandler(int *pUserData, const int *pDeviceData, size_t count)
{
    __block int tmp = 0;
    dispatch_async(dispatch_get_global_queue(QOS_CLASS_USER_INITIATED, 0), ^{
        tmp = DoDataProcessNormal(pUserData, pDeviceData, count / 2);
    });
    
    int sum = DoDataProcessNormal(&pUserData[count / 2], &pDeviceData[count / 2], count / 2);
    
    while(tmp == 0)
        asm("pause");
    sum += tmp;
    
    return sum;
}

static int NonTmpSingleCoreHandler(int *pUserData, const int *pDeviceData, size_t count)
{
    return DoDataProcessNonTmp(pUserData, pDeviceData, count);
}

static int NonTmpDuoCoreHandler(int *pUserData, const int *pDeviceData, size_t count)
{
    __block int tmp = 0;
    dispatch_async(dispatch_get_global_queue(QOS_CLASS_USER_INITIATED, 0), ^{
        tmp = DoDataProcessNonTmp(pUserData, pDeviceData, count / 2);
    });
    
    int sum = DoDataProcessNonTmp(&pUserData[count / 2], &pDeviceData[count / 2], count / 2);
    
    while(tmp == 0)
        asm("pause");
    sum += tmp;
    
    return sum;
}

static int (* const sHandlerList[])(int*, const int*, size_t) = {
    NormalSingleCoreHandler, NormalDuoCoreHandler,
    NonTmpSingleCoreHandler, NonTmpDuoCoreHandler
};

int main(void)
{
    cl_int ret;
    
    printf("The kernel code is: \n%s\n", sKernelCode);
    
    cl_platform_id platform_id = NULL;
    cl_device_id device_id = NULL;
    cl_context context = NULL;
    cl_command_queue command_queue = NULL;
    cl_mem srcDstMemObj = NULL;
    cl_program program = NULL;
    cl_kernel kernel = NULL;
    
    int *pHostBuffer = NULL;
    int *pUserData = NULL;
    
    // 获得OpenCL平台
    clGetPlatformIDs(1, &platform_id, NULL);
    if(platform_id == NULL)
    {
        puts("Get OpenCL platform failed!");
        goto FINISH;
    }
    
    // 获得OpenCL计算设备，这里使用GPU类型的计算设备
    clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL);
    if(device_id == NULL)
    {
        puts("No GPU available as a compute device!");
        goto FINISH;
    }
    
    // 根据设备ID来创建上下文
    context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);
    if(context == NULL)
    {
        puts("Context not established!");
        goto FINISH;
    }
    
    // 根据上下文与设备ID来创建命令队列
    command_queue = clCreateCommandQueue(context, device_id, 0, &ret);
    if(command_queue == NULL)
    {
        puts("Command queue cannot be created!");
        goto FINISH;
    }
    
    // 我们分配4MB的存储空间作为GPU计算的缓存长度
    const size_t contentLength = sizeof(int) * 1024 * 1024;
    
    // 在主机端分配并初始化4MB的缓存数据
    pHostBuffer = malloc(contentLength);
    pUserData = malloc(contentLength);
    for(int i = 0; i < contentLength / sizeof(int); i++)
    {
        pHostBuffer[i] = i + 1;
        pUserData[i] = i;
    }
    
    srcDstMemObj = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR, contentLength, pHostBuffer, &ret);
    if(srcDstMemObj == NULL)
    {
        puts("Memory object failed to create!");
        goto FINISH;
    }
    
    int *pDeviceBuffer = clEnqueueMapBuffer(command_queue, srcDstMemObj, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, contentLength, 0, NULL, NULL, &ret);
    
    if(pDeviceBuffer == pHostBuffer)
        puts("OK! Good!");
    else
    {
        // 若从GPU端映射得到的存储器地址与原先主机端的不同，则将数据从主机端传递到GPU端
        ret = clEnqueueWriteBuffer(command_queue, srcDstMemObj, CL_TRUE, 0, contentLength, pHostBuffer, 0, NULL, NULL);
        if(ret != CL_SUCCESS)
        {
            puts("Data transfer failed");
            goto FINISH;
        }
    }
    
    const var length = strlen(sKernelCode);
    program = clCreateProgramWithSource(context, 1, (const char**)&sKernelCode, (const size_t[]){length}, &ret);

    ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
    if (ret != CL_SUCCESS)
    {
        size_t len;
        char buffer[8 * 1024];
        
        printf("Error: Failed to build program executable!\n");
        clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
        printf("%s\n", buffer);
        goto FINISH;
    }
    
    // 创建内核函数
    kernel = clCreateKernel(program, "test", &ret);
    if(kernel == NULL)
    {
        puts("Kernel failed to create!");
        goto FINISH;
    }
    
    size_t maxWorkGroupSize = 0;
    ret = clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(maxWorkGroupSize), &maxWorkGroupSize, NULL);
    if(ret != CL_SUCCESS)
    {
        puts("Query max workgroup size failed!");
        goto FINISH;
    }
    printf("Current work-group size: %zu\n", maxWorkGroupSize);
    
    ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&srcDstMemObj);
    
    if(ret != CL_SUCCESS)
    {
        puts("Set arguments error!");
        goto FINISH;
    }
    
    struct timeval tBegin[TEST_LOOP_COUNT], tEnd[TEST_LOOP_COUNT];
    int64_t sum = 0;
    
    for(int i = 0; i < TEST_LOOP_COUNT; i++)
    {
        gettimeofday(&tBegin[i], NULL);

        // 将内核执行命令排入命令队列
        ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL,
                                     (const size_t[]){contentLength / sizeof(int) / 4},
                                     (const size_t[]){maxWorkGroupSize}, 0,
                                     NULL, NULL);
        if(ret != CL_SUCCESS)
        {
            puts("kernel1 execution failed");
            goto FINISH;
        }
        
        // 这里用clFinish做命令执行同步
        clFinish(command_queue);

        sum += sHandlerList[CalculateType_NONTMP_DUO_CORE](pUserData, pDeviceBuffer, contentLength / sizeof(int));
        
        gettimeofday(&tEnd[i], NULL);
    }
    
    var timeSpentForMicroseconds = 1000000LL * (tEnd[9].tv_sec - tBegin[9].tv_sec) + (tEnd[9].tv_usec - tBegin[9].tv_usec);
    for(int i = 10; i < TEST_LOOP_COUNT; i++)
    {
        const var spent = 1000000LL * (tEnd[i].tv_sec - tBegin[i].tv_sec) + (tEnd[i].tv_usec - tBegin[i].tv_usec);
        if(timeSpentForMicroseconds > spent)
            timeSpentForMicroseconds = spent;
    }

    printf("Time spent: %.2fms\n", (double)timeSpentForMicroseconds / 1000.0);
    printf("sum = %lld\n", sum);
    
    // 做数据校验
    int index = 0;
    for(int i = 0; i < 10; i++)
    {
        for(int j = 0; j < 10; j++, index++)
            printf("[%d] = %d, ", index, pDeviceBuffer[index]);
        puts("");
    }

FINISH:
    
    if(pHostBuffer != NULL)
        free(pHostBuffer);
    if(pUserData != NULL)
        free(pUserData);

    if(srcDstMemObj != NULL)
        clReleaseMemObject(srcDstMemObj);
    
    if(kernel != NULL)
        clReleaseKernel(kernel);
    
    if(program != NULL)
        clReleaseProgram(program);
    
    if(command_queue != NULL)
        clReleaseCommandQueue(command_queue);
    
    if(context != NULL)
        clReleaseContext(context);
    
    puts("Program complete");
    
    return 0;
}

下面是附属的汇编代码：

//
//  asm.s
//  OpenCLTest
//
//  Created by Zenny Chen on 2018/7/18.
//  Copyright © 2018年 Zenny Chen. All rights reserved.
//

.text
.align 4
.intel_syntax

/// int DoDataProcessNormal(int *pUserData, const int *pDeviceData, size_t count)
.globl _DoDataProcessNormal, _DoDataProcessNonTmp

_DoDataProcessNormal:

    // rdi: pUserData
    // rsi: pDeviceData
    // rdx: count

    mov         eax, 1
    vmovd       xmm15, eax
    vinserti128 ymm15, ymm15, xmm15, 1
    vpshufd     ymm15, ymm15, 0
    vpxor       ymm14, ymm14, ymm14

    // count /= 4 * 8
    shr         edx, 5

DoDataProcessNormal_LOOP:

    vmovdqa     ymm0, [rdi + 0]
    vmovdqa     ymm8, [rsi + 0]

    vmovdqa     ymm1, [rdi + 32]
    vmovdqa     ymm9, [rsi + 32]

    vmovdqa     ymm2, [rdi + 64]
    vmovdqa     ymm10, [rsi + 64]

    vmovdqa     ymm3, [rdi + 96]
    vmovdqa     ymm11, [rsi + 96]

    vpaddd      ymm0, ymm0, ymm15
    vpaddd      ymm8, ymm8, ymm9
    vpaddd      ymm1, ymm1, ymm15
    vpaddd      ymm10, ymm10, ymm11
    vpaddd      ymm2, ymm2, ymm15
    vpaddd      ymm3, ymm3, ymm15
    vpaddd      ymm8, ymm8, ymm10

    vmovdqa     [rdi + 0], ymm0
    vmovdqa     [rdi + 32], ymm1
    vmovdqa     [rdi + 64], ymm2
    vmovdqa     [rdi + 96], ymm3

    vpaddd      ymm14, ymm14, ymm8

    add         rdi, 32 * 4
    add         rsi, 32 * 4

    sub         edx, 1
    jnz         DoDataProcessNormal_LOOP

    vextracti128    xmm0, ymm14, 1
    vpaddd      xmm0, xmm0, xmm14
    // 4 int -> 2 int
    vphaddd     xmm0, xmm0, xmm15
    // 2 int -> 1 int
    vphaddd     xmm0, xmm0, xmm15
    vmovd       eax, xmm0

    ret


/// int DoDataProcessNonTmp(int *pUserData, const int *pDeviceData, size_t count)
_DoDataProcessNonTmp:

    // rdi: pUserData
    // rsi: pDeviceData
    // rdx: count

    mov         eax, 1
    vmovd       xmm15, eax
    vinserti128 ymm15, ymm15, xmm15, 1
    vpshufd     ymm15, ymm15, 0
    vpxor       ymm14, ymm14, ymm14

    // count /= 4 * 8
    shr         edx, 5

DoDataProcessNonTmp_LOOP:

    vmovdqa     ymm0, [rdi + 0]
    vmovntdqa   ymm8, [rsi + 0]

    vmovdqa     ymm1, [rdi + 32]
    vmovntdqa   ymm9, [rsi + 32]

    vmovdqa     ymm2, [rdi + 64]
    vmovntdqa   ymm10, [rsi + 64]

    vmovdqa     ymm3, [rdi + 96]
    vmovntdqa   ymm11, [rsi + 96]

    vpaddd      ymm0, ymm0, ymm15
    vpaddd      ymm8, ymm8, ymm9
    vpaddd      ymm1, ymm1, ymm15
    vpaddd      ymm10, ymm10, ymm11
    vpaddd      ymm2, ymm2, ymm15
    vpaddd      ymm3, ymm3, ymm15
    vpaddd      ymm8, ymm8, ymm10

    vmovdqa     [rdi + 0], ymm0
    vmovdqa     [rdi + 32], ymm1
    vmovdqa     [rdi + 64], ymm2
    vmovdqa     [rdi + 96], ymm3

    vpaddd      ymm14, ymm14, ymm8

    add         rdi, 32 * 4
    add         rsi, 32 * 4

    sub         edx, 1
    jnz         DoDataProcessNormal_LOOP

    vextracti128    xmm0, ymm14, 1
    vpaddd      xmm0, xmm0, xmm14
    // 4 int -> 2 int
    vphaddd     xmm0, xmm0, xmm15
    // 2 int -> 1 int
    vphaddd     xmm0, xmm0, xmm15
    vmovd       eax, xmm0

    ret

OpenCL的核心GPU实现一般不会分配出真正意义上Write-Combined类型的存储区域，因此我们下面上一个对应的Metal API版本：

//
//  main.m
//  NonTemporalTest
//
//  Created by Zenny Chen on 2018/7/19.
//  Copyright © 2018年 CodeLearning. All rights reserved.
//

@import Foundation;
@import Metal;

#ifndef var
#define var     __auto_type
#endif

#define TEST_LOOP_COUNT     1000

/// 以一般访存模式处理数据
/// @param pUserData 指向主机端用户数据的缓存
/// @param pDeviceData 指向设备数据的缓存
/// @param count 指定缓存中int数据元素的个数
/// @return 返回计算后的设备数据的某个整数值
extern int DoDataProcessNormal(int *pUserData, const int *pDeviceData, size_t count);

/// 以非临时不影响Cache的访存模式处理数据
/// @param pUserData 指向主机端用户数据的缓存
/// @param pDeviceData 指向设备数据的缓存
/// @param count 指定缓存中int数据元素的个数
/// @return 返回计算后的设备数据的某个整数值
extern int DoDataProcessNonTmp(int *pUserData, const int *pDeviceData, size_t count);

/// 以非临时访存模式读取所指定的设备端数据
/// @param pDeviceData 指向设备数据的缓存
/// @param count 指定缓存中int数据元素的个数
/// @return 返回计算后的设备数据的某个整数值
extern int NonTmpReadOnce(const int *pDeviceData, size_t count);

/// 以普通的访存模式读取所指定的设备端数据
/// @param pDeviceData 指向设备数据的缓存
/// @param count 指定缓存中int数据元素的个数
/// @return 返回计算后的设备数据的某个整数值
extern int NormalReadOnce(const int *pDeviceData, size_t count);

enum CalculateType
{
    CalculateType_NORMAL_SINGLE_CORE,
    CalculateType_NORMAL_DUO_CORE,
    CalculateType_NONTMP_SINGLE_CORE,
    CalculateType_NONTMP_DUO_CORE,
    CalculateType_NONTMP_READ_ONCE,
    CalculateType_NORMAL_READ_ONCE
};

static int NormalSingleCoreHandler(int *pUserData, const int *pDeviceData, size_t count)
{
    return DoDataProcessNormal(pUserData, pDeviceData, count);
}

static int NormalDuoCoreHandler(int *pUserData, const int *pDeviceData, size_t count)
{
    __block int tmp = 0;
    dispatch_async(dispatch_get_global_queue(QOS_CLASS_USER_INITIATED, 0), ^{
        tmp = DoDataProcessNormal(pUserData, pDeviceData, count / 2);
    });
    
    int sum = DoDataProcessNormal(&pUserData[count / 2], &pDeviceData[count / 2], count / 2);
    
    while(tmp == 0)
    asm("pause");
    sum += tmp;
    
    return sum;
}

static int NonTmpSingleCoreHandler(int *pUserData, const int *pDeviceData, size_t count)
{
    return DoDataProcessNonTmp(pUserData, pDeviceData, count);
}

static int NonTmpDuoCoreHandler(int *pUserData, const int *pDeviceData, size_t count)
{
    __block int tmp = 0;
    dispatch_async(dispatch_get_global_queue(QOS_CLASS_USER_INITIATED, 0), ^{
        tmp = DoDataProcessNonTmp(pUserData, pDeviceData, count / 2);
    });
    
    int sum = DoDataProcessNonTmp(&pUserData[count / 2], &pDeviceData[count / 2], count / 2);
    
    while(tmp == 0)
    asm("pause");
    sum += tmp;
    
    return sum;
}

static int NonTmpReadOnceHandler(int *pUserData, const int *pDeviceData, size_t count)
{
    return NonTmpReadOnce(pDeviceData, count);
}

static int NormalReadOnceHandler(int *pUserData, const int *pDeviceData, size_t count)
{
    return NormalReadOnce(pDeviceData, count);
}

static int (* const sHandlerList[])(int*, const int*, size_t) = {
    NormalSingleCoreHandler, NormalDuoCoreHandler,
    NonTmpSingleCoreHandler, NonTmpDuoCoreHandler,
    NonTmpReadOnceHandler, NormalReadOnceHandler
};

int main(int argc, const char * argv[])
{
    @autoreleasepool
    {
        // 创建默认计算设备
        var device = MTLCreateSystemDefaultDevice();
        
        // 创建库
        var library = device.newDefaultLibrary;
        
        // 创建计算函数
        var function = [library newFunctionWithName:@"test"];
        [library release];
        
        // 创建计算流水线
        var pipelineState = [device newComputePipelineStateWithFunction:function error:NULL];
        [function release];
        
        // 获得当前上下文中一个线程组中最多可以容纳多少个线程
        const var threadgroupSize = pipelineState.maxTotalThreadsPerThreadgroup;
        NSLog(@"Current threadgroup size: %tu", threadgroupSize);
        
        // 创建命令队列
        var commandQueue = device.newCommandQueue;
        
        const var elemCount = 1024 * 1024;
        
        // 初始化数据
        int *hostBuffer = malloc(elemCount * sizeof(int));
        int *userData = malloc(elemCount * sizeof(int));
        
        for(int i = 0; i < elemCount; i++)
        {
            hostBuffer[i] = i + 1;
            userData[i] = i;
        }
        
        // 创建缓存对象
        var memBuffer = [device newBufferWithBytes:hostBuffer length:elemCount * sizeof(int) options:MTLResourceCPUCacheModeWriteCombined];
        
        const MTLSize threadsPerGroup = {threadgroupSize, 1, 1};
        const MTLSize nThreadgroups = {elemCount / 4 / threadgroupSize, 1, 1};
        
        NSTimeInterval beginTime[TEST_LOOP_COUNT], endTime[TEST_LOOP_COUNT];
        int64_t sum = 0;
        
        for(int i = 0; i < TEST_LOOP_COUNT; i++)
        {
            // 获取命令缓存
            var commandBuffer = commandQueue.commandBuffer;
            
            // 获取命令编码器并设置其流水线状态
            var commandEncoder = commandBuffer.computeCommandEncoder;
            [commandEncoder setComputePipelineState:pipelineState];
            
            // 对命令编码器设置参数，
            // 我们在Metal Shading文件中所看到的参数次序就是根据这个次序安排的
            [commandEncoder setBuffer:memBuffer offset:0 atIndex:0];
            
            // 分派计算线程
            [commandEncoder dispatchThreadgroups:nThreadgroups threadsPerThreadgroup:threadsPerGroup];
            [commandEncoder endEncoding];
            
            beginTime[i] = NSProcessInfo.processInfo.systemUptime;
            
            // 提交
            [commandBuffer commit];
            
            // 这里挂起当前线程，等待命令完全执行完毕后再继续执行后续指令
            [commandBuffer waitUntilCompleted];
            
            sum += sHandlerList[CalculateType_NONTMP_READ_ONCE](userData, memBuffer.contents, elemCount);
            
            endTime[i] = NSProcessInfo.processInfo.systemUptime;
        }
        
        var minTimeSpent = endTime[9] - beginTime[9];
        for(int i = 10; i < TEST_LOOP_COUNT; i++)
        {
            const var timeSpent = endTime[i] - beginTime[i];
            if(minTimeSpent > timeSpent)
            minTimeSpent = timeSpent;
        }
        NSLog(@"Time spent: %.2fms\n", minTimeSpent * 1000.0);
        NSLog(@"sum = %lld\n", sum);
        
        // 释放资源
        [memBuffer release];
        [pipelineState release];
        [commandQueue release];
        [device release];
        
        free(hostBuffer);
        free(userData);
    }
    
    return 0;
}

以下是附属的Metal源文件内容：

//
//  compute.metal
//  NonTemporalTest
//
//  Created by Zenny Chen on 2018/7/19.
//  Copyright © 2018年 CodeLearning. All rights reserved.
//

#include <metal_stdlib>
using namespace metal;

kernel void test(device int4 *memBuffer [[ buffer(0) ]],
                 uint gid [[ thread_position_in_grid ]])
{
    memBuffer[gid] += 1;
}

为了能更好地验证非临时访存与普通访存对WC存储区域读数据的性能，我们在汇编文件中新增了两个函数进行对比测试，我们下面贴出完整的汇编代码：

//
//  asm.s
//  NonTemporalTest
//
//  Created by Zenny Chen on 2018/7/19.
//  Copyright © 2018年 CodeLearning. All rights reserved.
//

.text
.align 4
.intel_syntax

/// int DoDataProcessNormal(int *pUserData, const int *pDeviceData, size_t count)
.globl _DoDataProcessNormal, _DoDataProcessNonTmp

/// int NonTmpReadOnce(const int *pDeviceData, size_t count)
.globl _NonTmpReadOnce, _NormalReadOnce

_DoDataProcessNormal:

    // rdi: pUserData
    // rsi: pDeviceData
    // rdx: count

    mov         eax, 1
    vmovd       xmm15, eax
    vinserti128 ymm15, ymm15, xmm15, 1
    vpshufd     ymm15, ymm15, 0
    vpxor       ymm14, ymm14, ymm14

    // count /= 4 * 8
    shr         edx, 5

DoDataProcessNormal_LOOP:

    vmovdqa     ymm0, [rdi + 0]
    vmovdqa     ymm8, [rsi + 0]

    vmovdqa     ymm1, [rdi + 32]
    vmovdqa     ymm9, [rsi + 32]

    vmovdqa     ymm2, [rdi + 64]
    vmovdqa     ymm10, [rsi + 64]

    vmovdqa     ymm3, [rdi + 96]
    vmovdqa     ymm11, [rsi + 96]

    vpaddd      ymm0, ymm0, ymm15
    vpaddd      ymm8, ymm8, ymm9
    vpaddd      ymm1, ymm1, ymm15
    vpaddd      ymm10, ymm10, ymm11
    vpaddd      ymm2, ymm2, ymm15
    vpaddd      ymm3, ymm3, ymm15
    vpaddd      ymm8, ymm8, ymm10

    vmovdqa     [rdi + 0], ymm0
    vmovdqa     [rdi + 32], ymm1
    vmovdqa     [rdi + 64], ymm2
    vmovdqa     [rdi + 96], ymm3

    vpaddd      ymm14, ymm14, ymm8

    add         rdi, 32 * 4
    add         rsi, 32 * 4

    sub         edx, 1
    jnz         DoDataProcessNormal_LOOP

    vextracti128    xmm0, ymm14, 1
    vpaddd      xmm0, xmm0, xmm14
    // 4 int -> 2 int
    vphaddd     xmm0, xmm0, xmm15
    // 2 int -> 1 int
    vphaddd     xmm0, xmm0, xmm15
    vmovd       eax, xmm0

    ret


/// int DoDataProcessNonTmp(int *pUserData, const int *pDeviceData, size_t count)
_DoDataProcessNonTmp:

    // rdi: pUserData
    // rsi: pDeviceData
    // rdx: count

    mov         eax, 1
    vmovd       xmm15, eax
    vinserti128 ymm15, ymm15, xmm15, 1
    vpshufd     ymm15, ymm15, 0
    vpxor       ymm14, ymm14, ymm14

    // count /= 4 * 8
    shr         edx, 5

DoDataProcessNonTmp_LOOP:

    vmovdqa     ymm0, [rdi + 0]
    vmovntdqa   ymm8, [rsi + 0]

    vmovdqa     ymm1, [rdi + 32]
    vmovntdqa   ymm9, [rsi + 32]

    vmovdqa     ymm2, [rdi + 64]
    vmovntdqa   ymm10, [rsi + 64]

    vmovdqa     ymm3, [rdi + 96]
    vmovntdqa   ymm11, [rsi + 96]

    vpaddd      ymm0, ymm0, ymm15
    vpaddd      ymm8, ymm8, ymm9
    vpaddd      ymm1, ymm1, ymm15
    vpaddd      ymm10, ymm10, ymm11
    vpaddd      ymm2, ymm2, ymm15
    vpaddd      ymm3, ymm3, ymm15
    vpaddd      ymm8, ymm8, ymm10

    vmovdqa     [rdi + 0], ymm0
    vmovdqa     [rdi + 32], ymm1
    vmovdqa     [rdi + 64], ymm2
    vmovdqa     [rdi + 96], ymm3

    vpaddd      ymm14, ymm14, ymm8

    add         rdi, 32 * 4
    add         rsi, 32 * 4

    sub         edx, 1
    jnz         DoDataProcessNormal_LOOP

    vextracti128    xmm0, ymm14, 1
    vpaddd      xmm0, xmm0, xmm14
    // 4 int -> 2 int
    vphaddd     xmm0, xmm0, xmm15
    // 2 int -> 1 int
    vphaddd     xmm0, xmm0, xmm15
    vmovd       eax, xmm0

    ret


// int NonTmpReadOnce(const int *pDeviceData, size_t count)
_NonTmpReadOnce:

    // rdi: pUserData
    // rsi: count
    vpxor       ymm15, ymm15, ymm15
    vpxor       ymm14, ymm14, ymm14

    // count /= 8 * 8
    shr         esi, 6

NonTmpReadOnce_LOOP:

    vmovntdqa   ymm0, [rdi + 0 * 32]
    vmovntdqa   ymm1, [rdi + 1 * 32]
    vmovntdqa   ymm2, [rdi + 2 * 32]
    vmovntdqa   ymm3, [rdi + 3 * 32]
    vmovntdqa   ymm4, [rdi + 4 * 32]
    vmovntdqa   ymm5, [rdi + 5 * 32]
    vmovntdqa   ymm6, [rdi + 6 * 32]
    vmovntdqa   ymm7, [rdi + 7 * 32]

    vpaddd      ymm0, ymm0, ymm1
    vpaddd      ymm2, ymm2, ymm3
    vpaddd      ymm4, ymm4, ymm5
    vpaddd      ymm6, ymm6, ymm7
    vpaddd      ymm0, ymm0, ymm2
    vpaddd      ymm4, ymm4, ymm6

    vpaddd      ymm15, ymm0, ymm4

    sub         esi, 1
    jne         NonTmpReadOnce_LOOP

    vextracti128    xmm0, ymm15, 1
    vpaddd      xmm0, xmm0, xmm15
    // 4 int -> 2 int
    vphaddd     xmm0, xmm0, xmm14
    // 2 int -> 1 int
    vphaddd     xmm0, xmm0, xmm14
    vmovd       eax, xmm0

    ret


// int NormalReadOnce(const int *pDeviceData, size_t count)
_NormalReadOnce:

    // rdi: pUserData
    // rsi: count
    vpxor       ymm15, ymm15, ymm15
    vpxor       ymm14, ymm14, ymm14

    // count /= 8 * 8
    shr         esi, 6

NormalReadOnce_LOOP:

    vmovdqa     ymm0, [rdi + 0 * 32]
    vmovdqa     ymm1, [rdi + 1 * 32]
    vmovdqa     ymm2, [rdi + 2 * 32]
    vmovdqa     ymm3, [rdi + 3 * 32]
    vmovdqa     ymm4, [rdi + 4 * 32]
    vmovdqa     ymm5, [rdi + 5 * 32]
    vmovdqa     ymm6, [rdi + 6 * 32]
    vmovdqa     ymm7, [rdi + 7 * 32]

    vpaddd      ymm0, ymm0, ymm1
    vpaddd      ymm2, ymm2, ymm3
    vpaddd      ymm4, ymm4, ymm5
    vpaddd      ymm6, ymm6, ymm7
    vpaddd      ymm0, ymm0, ymm2
    vpaddd      ymm4, ymm4, ymm6

    vpaddd      ymm15, ymm0, ymm4

    sub         esi, 1
    jne         NormalReadOnce_LOOP

    vextracti128    xmm0, ymm15, 1
    vpaddd      xmm0, xmm0, xmm15
    // 4 int -> 2 int
    vphaddd     xmm0, xmm0, xmm14
    // 2 int -> 1 int
    vphaddd     xmm0, xmm0, xmm14
    vmovd       eax, xmm0

    ret

我们可以看到，WC（Write Combining）存储器类型由于不被映射到Cache，因此访问速度非常慢。但它仍然可以通过大规模并行访问来减少访存消耗的时间。通过上述代码测试例子，我们可以发现，对于WC的存储区域，使用非临时读与一般正常的读对性能在对此区域的每个数据元素只读一次的情况下几乎没啥区别。但如果对其中某个数据元素访问多次的话，那么使用非临时访问模式就要比普通模式强大很多了！在x86微处理器架构中，像MOVNTDQA这种指令属于非临时暗示性指令，如果它所操作的存储区域是WC存储区域，那么该暗示性指令的实现可以是将此数据元素加载到一个临时内部缓存，该缓存等价于一条对齐的Cache行，但不会将此数据填充到Cache中去。在Cache中任一有存储器类型混叠的行将会被窥探并冲刷出去。后继对WC的Cache行的未读部分进行MOVNTDQA的读将会从临时内部缓存得到数据，如果数据可用的话。临时内部缓存可以被处理器以任何理由、在任何时刻冲刷出去。
当读取数据的时候，非临时暗示是通过使用一个写绑定（WC）存储器类型协议来实现的。使用此协议，处理器并不把数据读到Cache层级，它也不会从存储器将对应的Cache行取到Cache层级中去。如果正在读的存储器区域并不是一个WC存储区域，那么此区域的存储器类型将会覆盖掉非临时暗示。也就是说，对于一个非WC存储区域而言，使用MOVNTDQA的读就跟使用MOVDQA的读一样了。
因为WC协议使用了一个弱次序的存储器一致性模型，因此，如果多个处理器（或多个处理器核心）对于所引用的存储器位置使用了不同的存储器类型，或是一条处理器的读与系统中其他代理的写进行同步，那么由一条MFENCE指令所实现的一个栅栏操作应该要与MOVNTDQA指令联合使用。流式加载暗示的一个处理器实现并不覆盖有效的存储器类型，但流式加载暗示的实现却是依赖于处理器的。比如，一个处理器实现可以选择忽略此暗示，并将此指令对于任一存储器类型都作为一条普通的MOVDQA进行处理。另外，其他一些实现也可以优化由MOVNTDQA在WC存储器类型上所生成的Cache读，以减少Cache逐出。
从上述测试结果来看，对64个连续的int数据类型进行重复读操作，在Intel Core i7 4650U上使用非临时缓存进行操作，仅需2.8ms，而使用普通的读操作则需要高达10.5ms。因此，总结下来，WC存储器类型配合非临时读，对于需要多次访问相同存储位置元素的任务而言还是非常有帮助的，尤其再加上多核心的支持，效果更佳。而如果此存储区域的数据在CPU端仅访问一次并且核心数又较少，那么我们就别把它设置为WC类型，而是使用普通类型加普通读操作才是最好的。

非临时缓存以及访存对计算性能的思考

相关文章

网友评论

延伸阅读

深度阅读

栏目导航

热点阅读

高性能计算相关