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; }