Skip to content

Commit

Permalink
temp push
Browse files Browse the repository at this point in the history
  • Loading branch information
@NguyenNhuDi
NguyenNhuDi committed Dec 31, 2024
1 parent fb16e25 commit eaca955
Show file tree
Hide file tree
Showing 13 changed files with 39 additions and 1 deletion.
2 changes: 2 additions & 0 deletions library/src/rng/config/lfsr113_config.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -46,6 +46,7 @@ struct generator_config_selector<ROCRAND_RNG_PSEUDO_LFSR113, T>
case target_arch::gfx90a: return 64;
case target_arch::gfx908: return 256;
case target_arch::gfx906: return 256;
case target_arch::gfx1201: return 128;
default:
return generator_config_defaults<ROCRAND_RNG_PSEUDO_LFSR113, T>::threads;
}
Expand All @@ -63,6 +64,7 @@ struct generator_config_selector<ROCRAND_RNG_PSEUDO_LFSR113, T>
case target_arch::gfx90a: return 2048;
case target_arch::gfx908: return 1024;
case target_arch::gfx906: return 2048;
case target_arch::gfx1201: return 512;
default:
return generator_config_defaults<ROCRAND_RNG_PSEUDO_LFSR113, T>::blocks;
}
Expand Down
2 changes: 2 additions & 0 deletions library/src/rng/config/mrg31k3p_config.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -46,6 +46,7 @@ struct generator_config_selector<ROCRAND_RNG_PSEUDO_MRG31K3P, T>
case target_arch::gfx90a: return 512;
case target_arch::gfx908: return 1024;
case target_arch::gfx906: return 256;
case target_arch::gfx1201: return 1024;
default:
return generator_config_defaults<ROCRAND_RNG_PSEUDO_MRG31K3P, T>::threads;
}
Expand All @@ -63,6 +64,7 @@ struct generator_config_selector<ROCRAND_RNG_PSEUDO_MRG31K3P, T>
case target_arch::gfx90a: return 2048;
case target_arch::gfx908: return 600;
case target_arch::gfx906: return 1792;
case target_arch::gfx1201: return 512;
default:
return generator_config_defaults<ROCRAND_RNG_PSEUDO_MRG31K3P, T>::blocks;
}
Expand Down
2 changes: 2 additions & 0 deletions library/src/rng/config/mrg32k3a_config.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -46,6 +46,7 @@ struct generator_config_selector<ROCRAND_RNG_PSEUDO_MRG32K3A, T>
case target_arch::gfx1030: return 256;
case target_arch::gfx908: return 1024;
case target_arch::gfx906: return 256;
case target_arch::gfx1201: return 1024;
default:
return generator_config_defaults<ROCRAND_RNG_PSEUDO_MRG32K3A, T>::threads;
}
Expand All @@ -63,6 +64,7 @@ struct generator_config_selector<ROCRAND_RNG_PSEUDO_MRG32K3A, T>
case target_arch::gfx90a: return 2048;
case target_arch::gfx908: return 600;
case target_arch::gfx906: return 2048;
case target_arch::gfx1201: return 256;
default:
return generator_config_defaults<ROCRAND_RNG_PSEUDO_MRG32K3A, T>::blocks;
}
Expand Down
2 changes: 2 additions & 0 deletions library/src/rng/config/mt19937_config.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -44,6 +44,7 @@ struct generator_config_selector<ROCRAND_RNG_PSEUDO_MT19937, T>
case target_arch::gfx942: return 256;
case target_arch::gfx90a: return 1024;
case target_arch::gfx908: return 512;
case target_arch::gfx1201: return 64;
default:
return generator_config_defaults<ROCRAND_RNG_PSEUDO_MT19937, T>::threads;
}
Expand All @@ -59,6 +60,7 @@ struct generator_config_selector<ROCRAND_RNG_PSEUDO_MT19937, T>
case target_arch::gfx942: return 1024;
case target_arch::gfx90a: return 64;
case target_arch::gfx908: return 64;
case target_arch::gfx1201: return 512;
default:
return generator_config_defaults<ROCRAND_RNG_PSEUDO_MT19937, T>::blocks;
}
Expand Down
2 changes: 2 additions & 0 deletions library/src/rng/config/mtgp32_config.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -46,6 +46,7 @@ struct generator_config_selector<ROCRAND_RNG_PSEUDO_MTGP32, T>
case target_arch::gfx90a: return 256;
case target_arch::gfx908: return 256;
case target_arch::gfx906: return 256;
case target_arch::gfx1201: return 256;
default:
return generator_config_defaults<ROCRAND_RNG_PSEUDO_MTGP32, T>::threads;
}
Expand All @@ -63,6 +64,7 @@ struct generator_config_selector<ROCRAND_RNG_PSEUDO_MTGP32, T>
case target_arch::gfx90a: return 512;
case target_arch::gfx908: return 480;
case target_arch::gfx906: return 448;
case target_arch::gfx1201: return 512;
default:
return generator_config_defaults<ROCRAND_RNG_PSEUDO_MTGP32, T>::blocks;
}
Expand Down
2 changes: 2 additions & 0 deletions library/src/rng/config/philox4_32_10_config.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -46,6 +46,7 @@ struct generator_config_selector<ROCRAND_RNG_PSEUDO_PHILOX4_32_10, T>
case target_arch::gfx90a: return 512;
case target_arch::gfx908: return 512;
case target_arch::gfx906: return 64;
case target_arch::gfx1201: return 1024;
default:
return generator_config_defaults<ROCRAND_RNG_PSEUDO_PHILOX4_32_10, T>::threads;
}
Expand All @@ -63,6 +64,7 @@ struct generator_config_selector<ROCRAND_RNG_PSEUDO_PHILOX4_32_10, T>
case target_arch::gfx90a: return 2048;
case target_arch::gfx908: return 3840;
case target_arch::gfx906: return 896;
case target_arch::gfx1201: return 1024;
default:
return generator_config_defaults<ROCRAND_RNG_PSEUDO_PHILOX4_32_10, T>::blocks;
}
Expand Down
2 changes: 2 additions & 0 deletions library/src/rng/config/threefry2_32_20_config.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -46,6 +46,7 @@ struct generator_config_selector<ROCRAND_RNG_PSEUDO_THREEFRY2_32_20, T>
case target_arch::gfx90a: return 512;
case target_arch::gfx908: return 512;
case target_arch::gfx906: return 256;
case target_arch::gfx1201: return 1024;
default:
return generator_config_defaults<ROCRAND_RNG_PSEUDO_THREEFRY2_32_20, T>::threads;
}
Expand All @@ -63,6 +64,7 @@ struct generator_config_selector<ROCRAND_RNG_PSEUDO_THREEFRY2_32_20, T>
case target_arch::gfx90a: return 2048;
case target_arch::gfx908: return 1200;
case target_arch::gfx906: return 896;
case target_arch::gfx1201: return 2048;
default:
return generator_config_defaults<ROCRAND_RNG_PSEUDO_THREEFRY2_32_20, T>::blocks;
}
Expand Down
2 changes: 2 additions & 0 deletions library/src/rng/config/threefry2_64_20_config.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -46,6 +46,7 @@ struct generator_config_selector<ROCRAND_RNG_PSEUDO_THREEFRY2_64_20, T>
case target_arch::gfx90a: return 256;
case target_arch::gfx908: return 256;
case target_arch::gfx906: return 256;
case target_arch::gfx1201: return 512;
default:
return generator_config_defaults<ROCRAND_RNG_PSEUDO_THREEFRY2_64_20, T>::threads;
}
Expand All @@ -63,6 +64,7 @@ struct generator_config_selector<ROCRAND_RNG_PSEUDO_THREEFRY2_64_20, T>
case target_arch::gfx90a: return 2048;
case target_arch::gfx908: return 960;
case target_arch::gfx906: return 560;
case target_arch::gfx1201: return 1024;
default:
return generator_config_defaults<ROCRAND_RNG_PSEUDO_THREEFRY2_64_20, T>::blocks;
}
Expand Down
2 changes: 2 additions & 0 deletions library/src/rng/config/threefry4_32_20_config.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -46,6 +46,7 @@ struct generator_config_selector<ROCRAND_RNG_PSEUDO_THREEFRY4_32_20, T>
case target_arch::gfx90a: return 256;
case target_arch::gfx908: return 256;
case target_arch::gfx906: return 256;
case target_arch::gfx1201: return 1024;
default:
return generator_config_defaults<ROCRAND_RNG_PSEUDO_THREEFRY4_32_20, T>::threads;
}
Expand All @@ -63,6 +64,7 @@ struct generator_config_selector<ROCRAND_RNG_PSEUDO_THREEFRY4_32_20, T>
case target_arch::gfx90a: return 2048;
case target_arch::gfx908: return 1200;
case target_arch::gfx906: return 896;
case target_arch::gfx1201: return 1024;
default:
return generator_config_defaults<ROCRAND_RNG_PSEUDO_THREEFRY4_32_20, T>::blocks;
}
Expand Down
2 changes: 2 additions & 0 deletions library/src/rng/config/threefry4_64_20_config.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -46,6 +46,7 @@ struct generator_config_selector<ROCRAND_RNG_PSEUDO_THREEFRY4_64_20, T>
case target_arch::gfx90a: return 256;
case target_arch::gfx908: return 128;
case target_arch::gfx906: return 128;
case target_arch::gfx1201: return 128;
default:
return generator_config_defaults<ROCRAND_RNG_PSEUDO_THREEFRY4_64_20, T>::threads;
}
Expand All @@ -63,6 +64,7 @@ struct generator_config_selector<ROCRAND_RNG_PSEUDO_THREEFRY4_64_20, T>
case target_arch::gfx90a: return 2048;
case target_arch::gfx908: return 1920;
case target_arch::gfx906: return 1792;
case target_arch::gfx1201: return 2048;
default:
return generator_config_defaults<ROCRAND_RNG_PSEUDO_THREEFRY4_64_20, T>::blocks;
}
Expand Down
2 changes: 2 additions & 0 deletions library/src/rng/config/xorwow_config.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -46,6 +46,7 @@ struct generator_config_selector<ROCRAND_RNG_PSEUDO_XORWOW, T>
case target_arch::gfx90a: return 128;
case target_arch::gfx908: return 256;
case target_arch::gfx906: return 256;
case target_arch::gfx1201: return 1024;
default:
return generator_config_defaults<ROCRAND_RNG_PSEUDO_XORWOW, T>::threads;
}
Expand All @@ -63,6 +64,7 @@ struct generator_config_selector<ROCRAND_RNG_PSEUDO_XORWOW, T>
case target_arch::gfx90a: return 2048;
case target_arch::gfx908: return 600;
case target_arch::gfx906: return 560;
case target_arch::gfx1201: return 256;
default:
return generator_config_defaults<ROCRAND_RNG_PSEUDO_XORWOW, T>::blocks;
}
Expand Down
17 changes: 16 additions & 1 deletion library/src/rng/config_types.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -53,6 +53,7 @@ enum class target_arch : unsigned int
gfx1100 = 1100,
gfx1101 = 1101,
gfx1102 = 1102,
gfx1201 = 1201,
unknown = std::numeric_limits<unsigned int>::max(),
};

Expand Down Expand Up @@ -85,6 +86,8 @@ __host__ __device__ constexpr target_arch get_device_arch()
return target_arch::gfx1101;
#elif defined(__gfx1102__)
return target_arch::gfx1102;
#elif defined(__gfx1201__)
return target_arch::gfx1201;
#else
return target_arch::unknown;
#endif
Expand All @@ -106,7 +109,8 @@ inline target_arch parse_gcn_arch(const std::string& arch_name)
"gfx1030",
"gfx1100",
"gfx1101",
"gfx1102"};
"gfx1102",
"gfx1201"};
const target_arch target_architectures[] = {
target_arch::gfx900,
target_arch::gfx902,
Expand All @@ -120,6 +124,7 @@ inline target_arch parse_gcn_arch(const std::string& arch_name)
target_arch::gfx1100,
target_arch::gfx1101,
target_arch::gfx1102,
target_arch::gfx1201,
};
static_assert(sizeof(target_names) / sizeof(target_names[0])
== sizeof(target_architectures) / sizeof(target_architectures[0]),
Expand Down Expand Up @@ -320,6 +325,10 @@ hipError_t get_generator_config(const hipStream_t stream,
{
target_arch current_arch;
const hipError_t error = get_device_arch(stream, current_arch);
printf("%d\n", current_arch);
printf("%d\n", current_arch);
printf("%d\n", current_arch);
printf("%d\n", current_arch);
if(error != hipSuccess)
{
return error;
Expand All @@ -329,6 +338,12 @@ hipError_t get_generator_config(const hipStream_t stream,
}
else
{
printf("IN ELSE OF IS_ORDERING_DYNAMIC\n");
printf("IN ELSE OF IS_ORDERING_DYNAMIC\n");
printf("IN ELSE OF IS_ORDERING_DYNAMIC\n");
printf("IN ELSE OF IS_ORDERING_DYNAMIC\n");
printf("IN ELSE OF IS_ORDERING_DYNAMIC\n");

config.threads = generator_config_defaults<GeneratorType, T>::threads;
config.blocks = generator_config_defaults<GeneratorType, T>::blocks;
}
Expand Down
1 change: 1 addition & 0 deletions library/src/rng/lfsr113.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -392,6 +392,7 @@ class lfsr113_generator_template : public generator_impl_base
return ROCRAND_STATUS_SUCCESS;
}

printf("Blocks: %d ---- Threads: %d\n", config.blocks, config.threads);
status = dynamic_dispatch(
m_order,
[&, this](auto is_dynamic)
Expand Down

0 comments on commit eaca955

Please sign in to comment.