* removed gfx940 and gfx941

* removed gfx940 and gfx941

* Update "gfx94" to "gfx942" in init.cc

* Updated remaining "gfx94" updates to "gfx942"

* Update filenames and variables from gfx940 to gfx942

---------

Co-authored-by: akolliasAMD <akollias@amd.com>
Αυτή η υποβολή περιλαμβάνεται σε:
corey-derochie-amd
2025-03-20 09:34:53 -06:00
υποβλήθηκε από GitHub
γονέας 6403afff4a
υποβολή 6505639cf4
30 αρχεία άλλαξαν με 74 προσθήκες και 79 διαγραφές
@@ -34,7 +34,7 @@ THE SOFTWARE.
} while (0)
// Macro for collecting HW_REG_XCC_ID
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) || defined(__gfx950__)
#if defined(__gfx942__) || defined(__gfx950__)
#define GetXccId(val) \
asm volatile ("s_getreg_b32 %0, hwreg(HW_REG_XCC_ID)" : "=s" (val));
#else
@@ -147,7 +147,7 @@ int main(int argc, char** argv) {
HIPCHECK(hipStreamCreateWithFlags(&stream[0], hipStreamNonBlocking));
HIPCHECK(hipDeviceEnablePeerAccess(device_id[1], 0));
HIPCHECK(hipGetDeviceProperties(&prop[0], device_id[0]));
HIPCHECK(hipExtMallocWithFlags((void**)&flag[0], HIP_IPC_MEM_MIN_SIZE, strncmp(prop[0].gcnArchName, "gfx94", 5) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipExtMallocWithFlags((void**)&flag[0], HIP_IPC_MEM_MIN_SIZE, strncmp(prop[0].gcnArchName, "gfx942", 6) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipHostMalloc ((void**)&time_delta[0], sizeof(uint64_t), hipHostMallocDefault));
HIPCHECK(hipMalloc((void**)&abortFlag[0], sizeof(uint32_t)));
HIPCHECK(hipMemsetAsync(flag[0], 0, HIP_IPC_MEM_MIN_SIZE, stream[0]));
@@ -158,7 +158,7 @@ int main(int argc, char** argv) {
HIPCHECK(hipStreamCreateWithFlags(&stream[1], hipStreamNonBlocking));
HIPCHECK(hipDeviceEnablePeerAccess(device_id[0], 0));
HIPCHECK(hipGetDeviceProperties(&prop[1], device_id[1]));
HIPCHECK(hipExtMallocWithFlags((void**)&flag[1], HIP_IPC_MEM_MIN_SIZE, strncmp(prop[1].gcnArchName, "gfx94", 5) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipExtMallocWithFlags((void**)&flag[1], HIP_IPC_MEM_MIN_SIZE, strncmp(prop[1].gcnArchName, "gfx942", 6) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipHostMalloc((void**)&time_delta[1], sizeof(uint64_t), hipHostMallocDefault));
HIPCHECK(hipMalloc((void**)&abortFlag[1], sizeof(uint32_t)));
HIPCHECK(hipMemsetAsync(flag[1], 0, HIP_IPC_MEM_MIN_SIZE, stream[1]));
@@ -174,11 +174,11 @@ int main(int argc, char** argv) {
double vega_gpu_rtc_freq;
HIPCHECK(hipStreamSynchronize(stream[0]));
vega_gpu_rtc_freq = strncmp(prop[0].gcnArchName, "gfx94", 5) == 0 ? 1.0E8 : 2.5E7;
vega_gpu_rtc_freq = strncmp(prop[0].gcnArchName, "gfx942", 6) == 0 ? 1.0E8 : 2.5E7;
fprintf(stdout, "One-way latency in us: %g\n", double(*time_delta[0]) * 1e6 / NUM_LOOPS_RUN / vega_gpu_rtc_freq / 2);
HIPCHECK(hipStreamSynchronize(stream[1]));
vega_gpu_rtc_freq = strncmp(prop[1].gcnArchName, "gfx94", 5) == 0 ? 1.0E8 : 2.5E7;
vega_gpu_rtc_freq = strncmp(prop[1].gcnArchName, "gfx942", 6) == 0 ? 1.0E8 : 2.5E7;
fprintf(stdout, "One-way latency in us: %g\n", double(*time_delta[1]) * 1e6 / NUM_LOOPS_RUN / vega_gpu_rtc_freq / 2);
HIPCHECK(hipFree(flag[0]));
@@ -86,7 +86,7 @@ int main(int argc, char** argv) {
HIPCHECK(hipStreamCreateWithFlags(&stream[0], hipStreamNonBlocking));
HIPCHECK(hipDeviceEnablePeerAccess(device_id[1], 0));
HIPCHECK(hipGetDeviceProperties(&prop[0], device_id[0]));
HIPCHECK(hipExtMallocWithFlags((void**)&flag[0], HIP_IPC_MEM_MIN_SIZE, strncmp(prop[0].gcnArchName, "gfx94", 5) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipExtMallocWithFlags((void**)&flag[0], HIP_IPC_MEM_MIN_SIZE, strncmp(prop[0].gcnArchName, "gfx942", 6) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipMalloc((void**)&time_delta[0], HIP_IPC_MEM_MIN_SIZE));
HIPCHECK(hipMemsetAsync(flag[0], 0, HIP_IPC_MEM_MIN_SIZE, stream[0]));
HIPCHECK(hipStreamSynchronize(stream[0]));
@@ -95,7 +95,7 @@ int main(int argc, char** argv) {
HIPCHECK(hipStreamCreateWithFlags(&stream[1], hipStreamNonBlocking));
HIPCHECK(hipDeviceEnablePeerAccess(device_id[0], 0));
HIPCHECK(hipGetDeviceProperties(&prop[1], device_id[1]));
HIPCHECK(hipExtMallocWithFlags((void**)&flag[1], HIP_IPC_MEM_MIN_SIZE, strncmp(prop[1].gcnArchName, "gfx94", 5) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipExtMallocWithFlags((void**)&flag[1], HIP_IPC_MEM_MIN_SIZE, strncmp(prop[1].gcnArchName, "gfx942", 6) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipMalloc((void**)&time_delta[1], HIP_IPC_MEM_MIN_SIZE));
HIPCHECK(hipMemsetAsync(flag[1], 0, HIP_IPC_MEM_MIN_SIZE, stream[1]));
HIPCHECK(hipStreamSynchronize(stream[1]));
@@ -109,11 +109,11 @@ int main(int argc, char** argv) {
double vega_gpu_rtc_freq;
HIPCHECK(hipStreamSynchronize(stream[0]));
vega_gpu_rtc_freq = strncmp(prop[0].gcnArchName, "gfx94", 5) == 0 ? 1.0E8 : 2.5E7;
vega_gpu_rtc_freq = strncmp(prop[0].gcnArchName, "gfx942", 6) == 0 ? 1.0E8 : 2.5E7;
fprintf(stdout, "One-way latency in us: %g\n", double(*time_delta[0]) * 1e6 / NUM_LOOPS_RUN / vega_gpu_rtc_freq / 2);
HIPCHECK(hipStreamSynchronize(stream[1]));
vega_gpu_rtc_freq = strncmp(prop[1].gcnArchName, "gfx94", 5) == 0 ? 1.0E8 : 2.5E7;
vega_gpu_rtc_freq = strncmp(prop[1].gcnArchName, "gfx942", 6) == 0 ? 1.0E8 : 2.5E7;
fprintf(stdout, "One-way latency in us: %g\n", double(*time_delta[1]) * 1e6 / NUM_LOOPS_RUN / vega_gpu_rtc_freq / 2);
HIPCHECK(hipFree(flag[0]));
@@ -426,7 +426,7 @@ int main(int argc,char* argv[])
static const char *ring_4p3l = "0 1 2 3|0 1 3 2|0 2 1 3|0 2 3 1|0 3 1 2|0 3 2 1";
static const char *ring_8p1h = "0 1 3 2 4 5 7 6|6 7 5 4 2 3 1 0|0 1 5 4 6 7 3 2|2 3 7 6 4 5 1 0";
static const char *ring_16p1h = "0 1 3 2 6 7 15 14 10 11 9 8 12 13 5 4|0 1 2 3 7 6 13 12 8 9 10 11 15 14 5 4|0 2 3 7 6 14 15 11 10 8 9 13 12 4 5 1|4 5 13 12 8 9 11 10 14 15 7 6 2 3 1 0|4 5 14 15 11 10 9 8 12 13 6 7 3 2 1 0|1 5 4 12 13 9 8 10 11 15 14 6 7 3 2 0";
static const char *ring_gfx940_8p = "0 1 2 3 4 5 6 7|0 1 2 3 4 5 7 6|0 2 4 1 3 6 5 7|0 2 4 6 1 7 3 5|0 3 1 5 2 7 4 6|0 3 5 1 6 2 7 4|0 4 1 7 3 6 2 5|7 6 5 4 3 2 1 0|6 7 5 4 3 2 1 0|7 5 6 3 1 4 2 0|5 3 7 1 6 4 2 0|6 4 7 2 5 1 3 0|4 7 2 6 1 5 3 0|5 2 6 3 7 1 4 0";
static const char *ring_gfx942_8p = "0 1 2 3 4 5 6 7|0 1 2 3 4 5 7 6|0 2 4 1 3 6 5 7|0 2 4 6 1 7 3 5|0 3 1 5 2 7 4 6|0 3 5 1 6 2 7 4|0 4 1 7 3 6 2 5|7 6 5 4 3 2 1 0|6 7 5 4 3 2 1 0|7 5 6 3 1 4 2 0|5 3 7 1 6 4 2 0|6 4 7 2 5 1 3 0|4 7 2 6 1 5 3 0|5 2 6 3 7 1 4 0";
setupPeers(connection_info, &is_xgmi);
if (!r) {
@@ -436,8 +436,8 @@ int main(int argc,char* argv[])
if (nGpu == 8 && !cr8g) {
hipDeviceProp_t prop;
HIPCHECK(hipGetDeviceProperties(&prop, 0));
if (strncmp(prop.gcnArchName, "gfx94", 5) == 0) {
r = (char *)ring_gfx940_8p;
if (strncmp(prop.gcnArchName, "gfx942", 6) == 0) {
r = (char *)ring_gfx942_8p;
if(!workgroups) workgroups = 28;
} else {
r = (char *)ring_8p1h;
@@ -521,11 +521,11 @@ int main(int argc,char* argv[])
profiling_data[i] = (struct profiling_data_t *)malloc(sizeof(struct profiling_data_t)*iters);
HIPCHECK(hipMalloc((void**) &d_profiling_data[i], sizeof(struct profiling_data_t)*iters));
HIPCHECK(hipExtMallocWithFlags((void**) &transfer_data[i], sizeof(struct transfer_data_t), strncmp(prop.gcnArchName, "gfx94", 5) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipExtMallocWithFlags((void**) &transfer_data[i], sizeof(struct transfer_data_t), strncmp(prop.gcnArchName, "gfx942", 6) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
for (int j = 0; j < workgroups; j++) {
HIPCHECK(hipExtMallocWithFlags((void**) &buff[i*MAX_WORKGROUPS+j], 2*N*sizeof(float), strncmp(prop.gcnArchName, "gfx94", 5) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipExtMallocWithFlags((void**) &buff[i*MAX_WORKGROUPS+j], 2*N*sizeof(float), strncmp(prop.gcnArchName, "gfx942", 6) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
// additional fine grained buffer for local doublecopy, only need 1 buffer (not used by remote)
HIPCHECK(hipExtMallocWithFlags((void**) &buff_fine[i*MAX_WORKGROUPS+j], N*sizeof(float), strncmp(prop.gcnArchName, "gfx94", 5) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipExtMallocWithFlags((void**) &buff_fine[i*MAX_WORKGROUPS+j], N*sizeof(float), strncmp(prop.gcnArchName, "gfx942", 6) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipMalloc((void**) &buff_coarse[i*MAX_WORKGROUPS+j], 2*N*sizeof(float)));
//randomize test data
hipLaunchKernelGGL(initTestDataKernel,
@@ -670,7 +670,7 @@ int main(int argc,char* argv[])
hipDeviceProp_t prop;
HIPCHECK(hipGetDeviceProperties(&prop, i));
double vega_gpu_rtc_freq, bw_std_dev = 0, mean_write_cycle = 0;
if (strncmp(prop.gcnArchName, "gfx94", 5) == 0)
if (strncmp(prop.gcnArchName, "gfx942", 6) == 0)
vega_gpu_rtc_freq = 1.0E8;
else
vega_gpu_rtc_freq = 2.5E7;
@@ -1,7 +1,7 @@
<system version="2">
<cpu numaid="2" affinity="000000ff,ffff0000,00000000,000000ff,ffff0000,00000000" arch="x86_64" vendor="AuthenticAMD" familyid="175" modelid="144">
<pci busid="0000:82:00.0" class="0x120000" vendor="0x1002" device="0x74a0" subsystem_vendor="0x1002" subsystem_device="0x74a0" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="0" sm="228" gcn="940" arch="38911" rank="0" gdr="1">
<gpu dev="0" sm="228" gcn="942" arch="38911" rank="0" gdr="1">
<xgmi target="0000:c2:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:02:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:42:00.0" count="1" tclass="0x120000"/>
@@ -15,7 +15,7 @@
</cpu>
<cpu numaid="3" affinity="ffffff00,00000000,00000000,ffffff00,00000000,00000000" arch="x86_64" vendor="AuthenticAMD" familyid="175" modelid="144">
<pci busid="0000:c2:00.0" class="0x120000" vendor="0x1002" device="0x74a0" subsystem_vendor="0x1002" subsystem_device="0x74a0" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="1" sm="228" gcn="940" arch="38911" rank="1" gdr="1">
<gpu dev="1" sm="228" gcn="942" arch="38911" rank="1" gdr="1">
<xgmi target="0000:82:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:02:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:42:00.0" count="1" tclass="0x120000"/>
@@ -29,7 +29,7 @@
</cpu>
<cpu numaid="0" affinity="00000000,00000000,00ffffff,00000000,00000000,00ffffff" arch="x86_64" vendor="AuthenticAMD" familyid="175" modelid="144">
<pci busid="0000:02:00.0" class="0x120000" vendor="0x1002" device="0x74a0" subsystem_vendor="0x1002" subsystem_device="0x74a0" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="2" sm="228" gcn="940" arch="38911" rank="2" gdr="1">
<gpu dev="2" sm="228" gcn="942" arch="38911" rank="2" gdr="1">
<xgmi target="0000:82:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:c2:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:42:00.0" count="1" tclass="0x120000"/>
@@ -43,7 +43,7 @@
</cpu>
<cpu numaid="1" affinity="00000000,0000ffff,ff000000,00000000,0000ffff,ff000000" arch="x86_64" vendor="AuthenticAMD" familyid="175" modelid="144">
<pci busid="0000:42:00.0" class="0x120000" vendor="0x1002" device="0x74a0" subsystem_vendor="0x1002" subsystem_device="0x74a0" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="3" sm="228" gcn="940" arch="38911" rank="3" gdr="1">
<gpu dev="3" sm="228" gcn="942" arch="38911" rank="3" gdr="1">
<xgmi target="0000:82:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:c2:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:02:00.0" count="1" tclass="0x120000"/>
@@ -129,11 +129,10 @@ NodeModelDesc model_descs[] = {
{"topo_8p1h_5.xml", " 8gfx910 2H3XGMI 8NIC 2AMD B"},
{"topo_16p1h.xml", "16gfx910 2H3XGMI 8NIC 4AMD A"},
{"topo_16p1h_vm.xml", "16gfx910 2H3XGMI 8NIC 4AMD B"},
// GFX 940
{"topo_4p_940.xml", " 4gfx940 1H3XGMI 4NIC 4AMD2 A"},
// GFX 942
{"topo_8p_940.xml", " 8gfx942 1H7XGMI 8NIC 2Intel A"},
{"topo_8p_940vm.xml", " 8gfx942 1H7XGMI 8NIC 2Intel B"},
{"topo_4p_942.xml", " 4gfx942 1H3XGMI 4NIC 4AMD2 A"},
{"topo_8p_942.xml", " 8gfx942 1H7XGMI 8NIC 2Intel A"},
{"topo_8p_942vm.xml", " 8gfx942 1H7XGMI 8NIC 2Intel B"},
{"topo_16p_gio-1s-1rp-cascade.xml", "16gfx942 2H7XGMI 1NIC 2AMD A"},
{"topo_16p_gio-3s-1rp-split-flat.xml", "16gfx942 2H7XGMI 1NIC 2AMD B"},
};