int main(){
float *A, *B, *C;
hipDeviceptr_t Ad, Bd, Cd;
A = new float[LEN];
B = new float[LEN];
C = new float[LEN];
for(uint32_t i=0;i<LEN;i++){
A[i] = i*1.0f;
B[i] = 1.0f;
C[i] = 0.0f;
}
hipInit(0);
hipDevice_t device;
hipCtx_t context;
hipDeviceGet(&device, 0);
hipCtxCreate(&context, 0, device);
hipMalloc((void**)&Ad, SIZE);
hipMalloc((void**)&Bd, SIZE);
hipMalloc((void**)&Cd, SIZE);
hipMemcpyHtoD(Ad, A, SIZE);
hipMemcpyHtoD(Bd, B, SIZE);
hipMemcpyHtoD(Cd, C, SIZE);
hipModule_t Module;
hipFunction_t Function;
hipModuleLoad(&Module, fileName);
hipModuleGetFunction(&Function, Module, kernel_name);
int n = LEN;
void * args[4] = {&Ad, &Bd, &Cd, &n};
hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, args, nullptr);
hipMemcpyDtoH(C, Cd, SIZE);
int mismatchCount = 0;
for(uint32_t i=0;i<LEN;i++){
if (A[i] + B[i] != C[i]) {
mismatchCount++;
std::cout<<"error: mismatch " << A[i]<<" + "<<B[i]<<" != "<<C[i]<<std::endl;
}
}
if (mismatchCount == 0) {
std::cout << "PASSED!\n";
} else {
std::cout << "FAILED!\n";
};
hipCtxDestroy(context);
return 0;
}
hipModulePtr CreateModule(const std::string& program_name, std::string params, bool is_kernel_str)
{
std::string filename =
is_kernel_str ? "tinygemm.cl" : program_name; // jn : don't know what this is
tmp_dir dir{filename};
std::string bin_file = dir.path(filename) + ".bin";
std::string hsaco_file = dir.path(filename) + ".o";
std::string obj_file = dir.path(filename) + ".obj";
std::string src = is_kernel_str ? program_name : GetKernelSrc(program_name);
if(!is_kernel_str && miopen::EndsWith(program_name, ".so"))
{
WriteFile(src, hsaco_file);
}
else if(!is_kernel_str && miopen::EndsWith(program_name, ".s"))
{
AmdgcnAssemble(src, params);
WriteFile(src, hsaco_file);
}
else
{
WriteFile(src, dir.path(filename));
#if MIOPEN_BUILD_DEV
params += " -Werror" + KernelWarningsString();
// params += " -Werror -Weverything -Wno-shorten-64-to-32 -Wno-unused-macros -Wno-unused-function
// -Wno-sign-compare -Wno-reserved-id-macro -Wno-sign-conversion -Wno-missing-prototypes
// -Wno-cast-qual -Wno-cast-align -Wno-conversion -Wno-double-promotion";
#else
params += " -Wno-everything";
#endif
dir.execute(HIP_OC_COMPILER, params + " " + filename + " -o " + hsaco_file);
}
hipModule_t raw_m;
auto status = hipModuleLoad(&raw_m, hsaco_file.c_str());
hipModulePtr m{raw_m};
if(status != hipSuccess)
MIOPEN_THROW_HIP_STATUS(status, "Failed creating module");
return m;
}
int main() {
float *A, *B;
hipDeviceptr_t Ad, Bd;
A = new float[LEN];
B = new float[LEN];
for (uint32_t i = 0; i < LEN; i++) {
A[i] = i * 1.0f;
B[i] = 0.0f;
}
hipInit(0);
hipDevice_t device;
hipCtx_t context;
hipDeviceGet(&device, 0);
hipCtxCreate(&context, 0, device);
hipMalloc((void**)&Ad, SIZE);
hipMalloc((void**)&Bd, SIZE);
hipMemcpyHtoD(Ad, A, SIZE);
hipMemcpyHtoD(Bd, B, SIZE);
hipModule_t Module;
hipFunction_t Function;
HIP_CHECK(hipModuleLoad(&Module, fileName));
HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name));
uint32_t len = LEN;
uint32_t one = 1;
struct {
void* _Ad;
void* _Bd;
} args;
args._Ad = Ad;
args._Bd = Bd;
size_t size = sizeof(args);
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
HIP_LAUNCH_PARAM_END};
HIP_CHECK(
hipExtModuleLaunchKernel(Function, LEN, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config, 0));
hipMemcpyDtoH(B, Bd, SIZE);
int mismatchCount = 0;
for (uint32_t i = 0; i < LEN; i++) {
if (A[i] != B[i]) {
mismatchCount++;
std::cout << "error: mismatch " << A[i] << " != " << B[i] << std::endl;
}
}
if (mismatchCount == 0) {
std::cout << "PASSED!\n";
} else {
std::cout << "FAILED!\n";
};
hipCtxDestroy(context);
return 0;
}
