Merge pull request #1163 from gargrahul/fix_module_api_kernarg_sample

Fix sample to use kernelargs for launch

hipamd/samples/0_Intro/module_api/Makefile

@@ -5,7 +5,7 @@ endif
 HIPCC=$(HIP_PATH)/bin/hipcc
 HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --compiler)
 
-all: vcpy_kernel.code runKernel.hip.out launchKernelHcc.hip.out
+all: vcpy_kernel.code runKernel.hip.out launchKernelHcc.hip.out defaultDriver.hip.out
 
 runKernel.hip.out: runKernel.cpp
 	$(HIPCC) $(HIPCC_FLAGS) $< -o $@
@@ -13,8 +13,8 @@ runKernel.hip.out: runKernel.cpp
 launchKernelHcc.hip.out: launchKernelHcc.cpp
 	$(HIPCC) $(HIPCC_FLAGS) $< -o $@
 
-#defaultDriver.hip.out: defaultDriver.cpp
+defaultDriver.hip.out: defaultDriver.cpp
-#	$(HIPCC)  $(HIPCC_FLAGS) $< -o $@
+	$(HIPCC)  $(HIPCC_FLAGS) $< -o $@
 
 vcpy_kernel.code: vcpy_kernel.cpp
 	$(HIPCC)  --genco  $(GENCO_FLAGS) $^ -o $@

hipamd/samples/0_Intro/module_api/defaultDriver.cpp

@@ -21,7 +21,6 @@ THE SOFTWARE.
 */
 
 #include "hip/hip_runtime.h"
-#include "hip/hip_runtime_api.h"
 #include <iostream>
 #include <fstream>
 #include <vector>
@@ -29,20 +28,18 @@ THE SOFTWARE.
 #define LEN 64
 #define SIZE LEN << 2
 
-#define fileName "test.co"
+#define fileName "vcpy_kernel.code"
-#define kernel_name "vadd"
+#define kernel_name "hello_world"
 
 int main() {
-    float *A, *B, *C;
+    float *A, *B;
-    hipDeviceptr_t Ad, Bd, Cd;
+    hipDeviceptr_t Ad, Bd;
     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;
+        B[i] = 0.0f;
-        C[i] = 0.0f;
     }
 
     hipInit(0);
@@ -53,28 +50,25 @@ int main() {
 
     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[2] = {&Ad, &Bd};
-    void* args[4] = {&Ad, &Bd, &Cd, &n};
 
     hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, args, nullptr);
 
-    hipMemcpyDtoH(C, Cd, SIZE);
+    hipMemcpyDtoH(B, Bd, SIZE);
     int mismatchCount = 0;
     for (uint32_t i = 0; i < LEN; i++) {
-        if (A[i] + B[i] != C[i]) {
+        if (A[i] != B[i]) {
             mismatchCount++;
-            std::cout << "error: mismatch " << A[i] << " + " << B[i] << " != " << C[i] << std::endl;
+            std::cout << "error: mismatch " << A[i] << " != " << C[i] << std::endl;
         }
     }
 
@@ -84,6 +78,10 @@ int main() {
         std::cout << "FAILED!\n";
     };
 
+    hipFree(Ad);
+    hipFree(Bd);
+    delete A;
+    delete B;
     hipCtxDestroy(context);
     return 0;
 }