光合开发者社区-论坛

光合基金

程序性能对比问题

2024-08-04 19:42:07

144

测试程序的核函数运行时间发现DCU与V100差距较大，DCU的运行速度只有V100的50%多一点，请问是什么原因？

V100显卡是从云平台租的，DCU由实验室自行购置，安装了24.04版本的驱动和DTK，分别使用cudaEventRecord,hipEventRecord的方式进行和核函数运行时间的测量，两个版本的代码除api外，其他部分没有不同。

回复 @网友xQ7Rwo

__global__ void

CalcSimGPU(float **d_Abd, float *d_sim_matrix, int *order_m, int *order_n, int iter, int file_count,

int OrderN, float *Dist_1, float *Dist_2, int *Order_1, int *Order_2,

int *Order_d) {

int tid = threadIdx.x + blockIdx.x * blockDim.x;

if (tid < iter) {

__shared__ float s_dist_1[SHARED_SIZE];

__shared__ float s_dist_2[SHARED_SIZE];

for (int count = 0; count < SHARED_SIZE; count++) {

s_dist_1[count] = Dist_1[count];

}

for (int count = 0; count < SHARED_SIZE; count++) {

s_dist_2[count] = Dist_2[count];

}

__syncthreads();

int m = order_m[tid];

int n = order_n[tid];

//int p = m * file_count + n - (m + 2) * (m + 1) * 0.5;

float *Abd_1 = d_Abd[m];

float *Abd_2 = d_Abd[n];

float Reg_1[REG_SIZE];

float Reg_2[REG_SIZE];

float total = 0;

for (int i = 0; i < OrderN; i++) {

int order_1 = Order_1[i];

int order_2 = Order_2[i];

int order_d = Order_d[i] + REG_SIZE;

float dist_1;

float dist_2;

if (i < SHARED_SIZE) {

dist_1 = 1 - s_dist_1[i];

dist_2 = 1 - s_dist_2[i];

} else {

dist_1 = 1 - Dist_1[i];

dist_2 = 1 - Dist_2[i];

}

float c1_1;

float c1_2;

float c2_1;

float c2_2;

if (order_1 >= 0) {

c1_1 = Abd_1[order_1];

c1_2 = Abd_2[order_1];

} else {

c1_1 = Reg_1[order_1 + REG_SIZE];

c1_2 = Reg_2[order_1 + REG_SIZE];

}

if (order_2 >= 0) {

c2_1 = Abd_1[order_2];

c2_2 = Abd_2[order_2];

} else {

c2_1 = Reg_1[order_2 + REG_SIZE];

c2_2 = Reg_2[order_2 + REG_SIZE];

}

float min_1 = fminf(c1_1, c1_2);

float min_2 = fminf(c2_1, c2_2);

total += (min_1 + min_2);

Reg_1[order_d] = (c1_1 - min_1) * dist_1 + (c2_1 - min_2) * dist_2;

Reg_2[order_d] = (c1_2 - min_1) * dist_1 + (c2_2 - min_2) * dist_2;

}

total *= 0.01f;

total = fminf(1.0f, fmaxf(0.0f, total));

d_sim_matrix[tid] = total;

} else { return; }

}

CUDA和DCU的sharedSize设置的是4096，blocksize设置的是16，DCU是实验室课题组购买的，通过rocminfo查询的信息如图所示，DCU的驱动和DTK都使用的是24.04版本。