added sample to porting driver guide
Change-Id: Ia2b54311cec617e5e2a162fd31415c8caad4a374
이 커밋은 다음에 포함됨:
@@ -171,3 +171,83 @@ hipFunction k;
|
||||
hipModuleGetFunction(&k, module, "myKernel");
|
||||
```
|
||||
|
||||
The below sample shows how to use `hipModuleGetFunction`.
|
||||
|
||||
```
|
||||
#include<hip_runtime.h>
|
||||
#include<hip_runtime_api.h>
|
||||
#include<iostream>
|
||||
#include<fstream>
|
||||
#include<vector>
|
||||
|
||||
#define LEN 64
|
||||
#define SIZE LEN<<2
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
#define fileName "vcpy_isa.co"
|
||||
#endif
|
||||
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
#define fileName "vcpy_isa.ptx"
|
||||
#endif
|
||||
|
||||
#define kernel_name "hello_world"
|
||||
|
||||
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;
|
||||
std::cout<<A[i] << " "<<B[i]<<std::endl;
|
||||
}
|
||||
|
||||
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipInit(0);
|
||||
hipDevice_t device;
|
||||
hipCtx_t context;
|
||||
hipDeviceGet(&device, 0);
|
||||
hipCtxCreate(&context, 0, device);
|
||||
#endif
|
||||
|
||||
hipMalloc((void**)&Ad, SIZE);
|
||||
hipMalloc((void**)&Bd, SIZE);
|
||||
|
||||
hipMemcpyHtoD(Ad, A, SIZE);
|
||||
hipMemcpyHtoD(Bd, B, SIZE);
|
||||
hipModule_t Module;
|
||||
hipFunction_t Function;
|
||||
hipModuleLoad(&Module, fileName);
|
||||
hipModuleGetFunction(&Function, Module, kernel_name);
|
||||
|
||||
std::vector<void*>argBuffer(2);
|
||||
memcpy(&argBuffer[0], &Ad, sizeof(void*));
|
||||
memcpy(&argBuffer[1], &Bd, sizeof(void*));
|
||||
|
||||
size_t size = argBuffer.size()*sizeof(void*);
|
||||
|
||||
void *config[] = {
|
||||
HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0],
|
||||
HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
|
||||
HIP_LAUNCH_PARAM_END
|
||||
};
|
||||
|
||||
hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&con fig);
|
||||
|
||||
hipMemcpyDtoH(B, Bd, SIZE);
|
||||
for(uint32_t i=0;i<LEN;i++){
|
||||
std::cout<<A[i]<<" - "<<B[i]<<std::endl;
|
||||
}
|
||||
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtxDetach(context);
|
||||
#endif
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
```
|
||||
|
||||
새 이슈에서 참조
사용자 차단