全局内存
合并访问
struc_add.h
#pragma once
#include "Global.h"
#include <stdio.h>
#include <stdlib.h>
#include <conio.h>
#include <assert.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#define NUM_ELEMENTS 4096
typedef struct
{
u32 a;
u32 b;
u32 c;
u32 d;
} INTERLEAVED_T;
typedef INTERLEAVED_T INTERLEAVED_ARRAY_T[NUM_ELEMENTS];
typedef u32 ARRAY_MEMBER_T[NUM_ELEMENTS];
typedef struct
{
ARRAY_MEMBER_T a;
ARRAY_MEMBER_T b;
ARRAY_MEMBER_T c;
ARRAY_MEMBER_T d;
} NON_INTERLEAVED_T;
//************************************
// Method: add_test_non_interleaved_cpu
// Brief: struct 内容非交叉 CPU
// Access: public
// Returns: __host__ float
// Qualifier:
// Param(I/O): NON_INTERLEAVED_T * const host_dest_ptr
// Param(I/O): const NON_INTERLEAVED_T * const host_src_ptr
// Param(I/O): const u32 iter
// Param(I/O): const u32 num_elements
//************************************
__host__ float add_test_non_interleaved_cpu(
NON_INTERLEAVED_T * const host_dest_ptr,
const NON_INTERLEAVED_T * const host_src_ptr,
const u32 iter,
const u32 num_elements);
//************************************
// Method: add_test_interleaved_cpu
// Brief: struct 内容交叉 CPU
// Access: public
// Returns: __host__ float
// Qualifier:
// Param(I/O): INTERLEAVED_T * const host_dest_ptr
// Param(I/O): const INTERLEAVED_T * const host_src_ptr
// Param(I/O): const u32 iter
// Param(I/O): const u32 num_elements
//************************************
__host__ float add_test_interleaved_cpu(INTERLEAVED_T * const host_dest_ptr, const INTERLEAVED_T * const host_src_ptr, const u32 iter, const u32 num_elements);
//************************************
// Method: add_kernel_interleaved
// Brief: struct 内容交叉 GPU
// Access: public
// Returns: __global__ void
// Qualifier:
// Param(I/O): INTERLEAVED_T * const dest_ptr
// Param(I/O): const INTERLEAVED_T * const src_ptr
// Param(I/O): const u32 iter
// Param(I/O): const u32 num_elements
//************************************
__global__ void add_kernel_interleaved(INTERLEAVED_T* const dest_ptr, const INTERLEAVED_T * const src_ptr, const u32 iter, const u32 num_elements);
//************************************
// Method: add_kernel_non_interleaved
// Brief: struct 内容非交叉 GPU
// Access: public
// Returns: __global__ void
// Qualifier:
// Param(I/O): NON_INTERLEAVED_T * const dest_ptr
// Param(I/O): const NON_INTERLEAVED_T * const src_ptr
// Param(I/O): const u32 iter
// Param(I/O): const u32 num_elements
//************************************
__global__ void add_kernel_non_interleaved(NON_INTERLEAVED_T * const dest_ptr, const NON_INTERLEAVED_T * const src_ptr, const u32 iter, const u32 num_elements);
//************************************
// Method: add_test_interleaved
// Brief: struct 内容交叉 GPU host
// Access: public
// Returns: __host__ float
// Qualifier:
// Param(I/O): INTERLEAVED_T * const host_dest_ptr
// Param(I/O): const INTERLEAVED_T * const host_src_ptr
// Param(I/O): const u32 iter
// Param(I/O): const u32 num_elements
//************************************
__host__ float add_test_interleaved(INTERLEAVED_T * const host_dest_ptr, const INTERLEAVED_T * const host_src_ptr, const u32 iter, const u32 num_elements);
//************************************
// Method: add_test_non_interleaved
// Brief: struct 内容非交叉 GPU host
// Access: public
// Returns: __host__ float
// Qualifier:
// Param(I/O): NON_INTERLEAVED_T * const dest_ptr
// Param(I/O): const NON_INTERLEAVED_T * const src_ptr
// Param(I/O): const u32 iter
// Param(I/O): const u32 num_elements
//************************************
__host__ float add_test_non_interleaved(NON_INTERLEAVED_T * const host_dest_ptr, const NON_INTERLEAVED_T * const host_src_ptr, const u32 iter, const u32 num_elements);
struct_add.cu
#include "struct_add.h"
#include <timer.h>
__host__ float add_test_non_interleaved_cpu(NON_INTERLEAVED_T * const host_dest_ptr, const NON_INTERLEAVED_T * const host_src_ptr, const u32 iter, const u32 num_elements)
{
StartTimer();
for (u32 tid = 0; tid < num_elements;tid++)
{
for (u32 i = 0; i < iter;i++)
{
host_dest_ptr->a[tid] += host_src_ptr->a[tid];
host_dest_ptr->b[tid] += host_src_ptr->b[tid];
host_dest_ptr->c[tid] += host_src_ptr->c[tid];
host_dest_ptr->d[tid] += host_src_ptr->d[tid];
}
}
const float delta = GetTimer();
return delta;
}
__host__ float add_test_interleaved_cpu(INTERLEAVED_T * const host_dest_ptr, const INTERLEAVED_T * const host_src_ptr, const u32 iter, const u32 num_elements)
{
StartTimer();
for (u32 tid = 0; tid < num_elements;tid++)
{
for (u32 i = 0; i < iter;i++)
{
host_dest_ptr[tid].a += host_src_ptr[tid].a;
host_dest_ptr[tid].b += host_src_ptr[tid].b;
host_dest_ptr[tid].c += host_src_ptr[tid].c;
host_dest_ptr[tid].d += host_src_ptr[tid].d;
}
}
const float delta = GetTimer();
return delta;
}
__global__ void add_kernel_interleaved(INTERLEAVED_T* const dest_ptr, const INTERLEAVED_T * const src_ptr, const u32 iter, const u32 num_elements)
{
const u32 tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < num_elements)
{
for (u32 i = 0; i < iter; i++)
{
dest_ptr[tid].a += src_ptr[tid].a;
dest_ptr[tid].b += src_ptr[tid].b;
dest_ptr[tid].c += src_ptr[tid].c;
dest_ptr[tid].d += src_ptr[tid].d;
}
}
}
__global__ void add_kernel_non_interleaved(NON_INTERLEAVED_T * const dest_ptr, const NON_INTERLEAVED_T * const src_ptr, const u32 iter, const u32 num_elements)
{
const u32 tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < num_elements)
{
for (u32 i = 0; i < iter;i++)
{
dest_ptr->a[tid] += src_ptr->a[tid];
dest_ptr->b[tid] += src_ptr->b[tid];
dest_ptr->c[tid] += src_ptr->c[tid];
dest_ptr->d[tid] += src_ptr->d[tid];
}
}
}
__host__ float add_test_interleaved(INTERLEAVED_T * const host_dest_ptr, const INTERLEAVED_T * const host_src_ptr, const u32 iter, const u32 num_elements)
{
const u32 num_threads = 256;
const u32 num_blocks = (num_elements + num_threads - 1) / num_threads;
const size_t num_bytes = sizeof(INTERLEAVED_T)*num_elements;
INTERLEAVED_T * device_dest_ptr;
INTERLEAVED_T * device_src_ptr;
CUDA_CALL(cudaMalloc((void**)&device_src_ptr, num_bytes));
CUDA_CALL(cudaMalloc((void**)&device_dest_ptr, num_bytes));
cudaEvent_t kernel_start, kernel_stop;
cudaEventCreate(&kernel_start, 0);
cudaEventCreate(&kernel_stop, 0);
cudaStream_t test_stream;
CUDA_CALL(cudaStreamCreate(&test_stream));
CUDA_CALL(cudaMemcpy(device_src_ptr, host_src_ptr, num_bytes, cudaMemcpyHostToDevice));
CUDA_CALL(cudaEventRecord(kernel_start, 0));
add_kernel_interleaved << <num_blocks, num_threads >> >(device_dest_ptr, device_src_ptr, iter, num_elements);
CUDA_CALL(cudaEventRecord(kernel_stop, 0));
CUDA_CALL(cudaEventSynchronize(kernel_stop));
float delta = 0.0f;
CUDA_CALL(cudaEventElapsedTime(&delta, kernel_start, kernel_stop));
CUDA_CALL(cudaFree(device_src_ptr));
CUDA_CALL(cudaFree(device_dest_ptr));
CUDA_CALL(cudaEventDestroy(kernel_start));
CUDA_CALL(cudaEventDestroy(kernel_stop));
CUDA_CALL(cudaStreamDestroy(test_stream));
return delta;
}
__host__ float add_test_non_interleaved(NON_INTERLEAVED_T * const host_dest_ptr, const NON_INTERLEAVED_T * const host_src_ptr, const u32 iter, const u32 num_elements)
{
const u32 num_threads = 256;
const u32 num_blocks = (num_elements + num_threads - 1) / num_threads;
// 应该是 NON_INTERLEAVED_T
const size_t num_bytes = sizeof(NON_INTERLEAVED_T);// *num_elements;
NON_INTERLEAVED_T * device_dest_ptr;
NON_INTERLEAVED_T * device_src_ptr;
CUDA_CALL(cudaMalloc((void**)&device_src_ptr, num_bytes));
CUDA_CALL(cudaMalloc((void**)&device_dest_ptr, num_bytes));
cudaEvent_t kernel_start, kernel_stop;
cudaEventCreate(&kernel_start, 0);
cudaEventCreate(&kernel_stop, 0);
cudaStream_t test_stream;
CUDA_CALL(cudaStreamCreate(&test_stream));
CUDA_CALL(cudaMemcpy(device_src_ptr, host_src_ptr, num_bytes, cudaMemcpyHostToDevice));
CUDA_CALL(cudaEventRecord(kernel_start, 0));
add_kernel_non_interleaved << <num_blocks, num_threads >> >(device_dest_ptr, device_src_ptr, iter, num_elements);
CUDA_CALL(cudaEventRecord(kernel_stop, 0));
CUDA_CALL(cudaEventSynchronize(kernel_stop));
float delta = 0.0f;
CUDA_CALL(cudaEventElapsedTime(&delta, kernel_start, kernel_stop));
CUDA_CALL(cudaFree(device_src_ptr));
CUDA_CALL(cudaFree(device_dest_ptr));
CUDA_CALL(cudaEventDestroy(kernel_start));
CUDA_CALL(cudaEventDestroy(kernel_stop));
CUDA_CALL(cudaStreamDestroy(test_stream));
return delta;
}
struct_add_main.cpp
/* CPU和GPU全局内存, 分别用交叉和非交叉方式来计算struct元素之和.*/
#include "Global.h"
#include "struct_add.h"
int main()
{
getCpuInfo();
getGpuInfo();
// 数据初始化
INTERLEAVED_T * host_src_interleaved_ptr = NULL;
INTERLEAVED_T * host_dest_interleaved_ptr = NULL;
host_src_interleaved_ptr = (INTERLEAVED_T *)malloc(sizeof(INTERLEAVED_T)*NUM_ELEMENTS);
host_dest_interleaved_ptr = (INTERLEAVED_T *)malloc(sizeof(INTERLEAVED_T)*NUM_ELEMENTS);
NON_INTERLEAVED_T host_src_non_interleaved_ptr;
NON_INTERLEAVED_T host_dest_non_interleaved_ptr;
for (int i = 0; i < NUM_ELEMENTS;i++)
{
host_src_interleaved_ptr[i].a = rand() % 4;
host_src_interleaved_ptr[i].b = rand() % 4;
host_src_interleaved_ptr[i].c = rand() % 4;
host_src_interleaved_ptr[i].d = rand() % 4;
host_src_non_interleaved_ptr.a[i] = rand() % 4;
host_src_non_interleaved_ptr.b[i] = rand() % 4;
host_src_non_interleaved_ptr.c[i] = rand() % 4;
host_src_non_interleaved_ptr.d[i] = rand() % 4;
}
const u32 iter = 10000;
// CPU 非交叉
float delta_non_interleaved_cpu = add_test_non_interleaved_cpu(&host_dest_non_interleaved_ptr, &host_src_non_interleaved_ptr, iter, NUM_ELEMENTS);
// CPU 交叉
float delta_interleaved_cpu = add_test_interleaved_cpu(host_dest_interleaved_ptr, host_src_interleaved_ptr, iter, NUM_ELEMENTS);
// GPU 非交叉
float delta_non_interleaved_gpu = add_test_non_interleaved(&host_dest_non_interleaved_ptr, &host_src_non_interleaved_ptr, iter, NUM_ELEMENTS);
delta_non_interleaved_gpu = add_test_non_interleaved(&host_dest_non_interleaved_ptr, &host_src_non_interleaved_ptr, iter, NUM_ELEMENTS);
// GPU 交叉
float delta_interleaved_gpu = add_test_interleaved(host_dest_interleaved_ptr, host_src_interleaved_ptr, iter, NUM_ELEMENTS);
delta_interleaved_gpu = add_test_interleaved(host_dest_interleaved_ptr, host_src_interleaved_ptr, iter, NUM_ELEMENTS);
// 输出
printf("耗时统计:\n");
printf("CPU\t非交叉: \t%.3f ms\n", delta_non_interleaved_cpu);
printf("CPU\t交叉: \t\t%.3f ms\n", delta_interleaved_cpu);
printf("GPU\t非交叉: \t%.3f ms\n", delta_non_interleaved_gpu);
printf("GPU\t交叉: \t\t%.3f ms\n", delta_interleaved_gpu);
return 0;
}
测试结果
相对来说, CPU的交叉形式比CPU的非交叉形式快.
GPU相反, 非交叉比交叉快.
但是, 测试结果的差异没有书上的大.