cuda kernel test

src/ngraph/CMakeLists.txt

@@ -203,7 +203,7 @@ if (NGRAPH_CPU_ENABLE AND LLVM_INCLUDE_DIR AND
             runtime/gpu/gpu_tensor_view.cpp
             runtime/gpu/gpu_tensor_view_wrapper.cpp
             runtime/gpu/gpu_util.cpp
-            runtime/gpu/gpu_cuda_kernel_emitter.cpp
+            runtime/gpu/gpu_cuda_kernel_emitters.cpp
             )
         set_property(SOURCE codegen/compiler.cpp APPEND_STRING PROPERTY COMPILE_DEFINITIONS
         "CUDA_HEADER_PATHS=\"${CUDA_INCLUDE_DIRS}\";")

src/ngraph/runtime/gpu/gpu_cuda_kernel_emitters.cpp

@@ -77,6 +77,71 @@ namespace ngraph
 
                     void emit_abs(void* in, void* out, size_t count)
                     {
+                            const char *op_abs = R"(  
+  extern "C" __global__  
+  void cuda_op_abs(float* in, float* out, size_t n)  
+  {  
+    size_t tid = blockIdx.x * blockDim.x + threadIdx.x;  
+    if(tid < n) 
+    {
+      out[tid] = fabsf(in[tid]);  
+    }
+  })";
+
+        size_t numBlocks = 4; size_t numThreads = 4; 
+    // Create an instance of nvrtcProgram with the code string. 
+
+    nvrtcProgram prog; 
+    NVRTC_SAFE_CALL(nvrtcCreateProgram(&prog, // prog i
+                op_abs, // buffer 
+                "op_abs.cu", // name 
+                0, // numHeaders 
+                NULL, // headers 
+                NULL)); // includeNames
+
+
+    const char *opts[] = {"--gpu-architecture=compute_35",
+        "--relocatable-device-code=true"};
+    nvrtcResult compileResult = nvrtcCompileProgram(prog, // prog 
+            2, // numOptions 
+            opts); // options
+    // Obtain compilation log from the program. 
+
+    size_t logSize; 
+
+    NVRTC_SAFE_CALL(nvrtcGetProgramLogSize(prog, &logSize)); 
+    char *log = new char[logSize]; 
+    NVRTC_SAFE_CALL(nvrtcGetProgramLog(prog, log)); 
+    std::cout << log << '\n'; 
+    delete[] log; 
+
+    if (compileResult != NVRTC_SUCCESS) {
+        exit(1);
+    }
+
+    size_t ptxSize; 
+    NVRTC_SAFE_CALL(nvrtcGetPTXSize(prog, &ptxSize)); 
+    char *ptx = new char[ptxSize]; 
+    NVRTC_SAFE_CALL(nvrtcGetPTX(prog, ptx)); // Destroy the program. 
+    NVRTC_SAFE_CALL(nvrtcDestroyProgram(&prog)); // Load the generated PTX and get a handle to the parent kernel. 
+
+    CUdevice cuDevice; 
+    CUcontext context; 
+    CUmodule module; 
+    CUfunction cuda_op_abs_kernel; 
+    CUDA_SAFE_CALL( cuInit(0)); 
+    CUDA_SAFE_CALL(cuDeviceGet(&cuDevice, 0)); 
+    CUDA_SAFE_CALL(cuCtxCreate(&context, 0, cuDevice)); 
+    // CUDA_SAFE_CALL(cuLinkCreate(0, 0 , 0, &linkState)); 
+    //CUDA_SAFE_CALL(cuLinkeAddFile(linkState, CU_JIT_INPUT_LIBRARY, ' ', 0, 0, 0));
+    //CUDA_SAFE_CALL(cuLinkAddData(linkState, CU_JIT_INPUT_PTX, (void *)ptx, ptxSize, "dynamic_parallelism.ptx", 0, 0, 0));
+
+    //size_t cubinSize;
+    //void *cubin;
+    //CUDA_SAFE_CALL(cuLinkComplete(linkState, &cubin, &cubinSize));
+    CUDA_SAFE_CALL(cuModuleLoadDataEx(&module, ptx, 0, 0, 0));
+    CUDA_SAFE_CALL(cuModuleGetFunction(&cuda_op_abs_kernel, module, "cuda_op_abs"));
+    
                         void *argsList[] = {In, Out, &count};
                         CUDA_SAFE_CALL(
                         cuLaunchKernel(cuda_op_abs_kernel, 

src/ngraph/runtime/gpu/gpu_emitter.cpp

@@ -85,78 +85,11 @@ void runtime::gpu::GPU_Emitter::EmitAbs(codegen::CodeWriter& writer,
         const vector<runtime::gpu::GPU_TensorViewWrapper>& args,
         const vector<runtime::gpu::GPU_TensorViewWrapper>& out)
 {
-    const char *op_abs = R"(  
-  extern "C" __global__  
-  void cuda_op_abs(float* in, float* out, size_t n)  
-  {  
-    size_t tid = blockIdx.x * blockDim.x + threadIdx.x;  
-    if(tid < n) 
-    {
-      out[tid] = fabsf(in[tid]);  
-    }
-  })";
-
-        size_t numBlocks = 4; size_t numThreads = 4; 
-    // Create an instance of nvrtcProgram with the code string. 
-
-    nvrtcProgram prog; 
-    NVRTC_SAFE_CALL(nvrtcCreateProgram(&prog, // prog i
-                op_abs, // buffer 
-                "op_abs.cu", // name 
-                0, // numHeaders 
-                NULL, // headers 
-                NULL)); // includeNames
-
-
-    const char *opts[] = {"--gpu-architecture=compute_35",
-        "--relocatable-device-code=true"};
-    nvrtcResult compileResult = nvrtcCompileProgram(prog, // prog 
-            2, // numOptions 
-            opts); // options
-    // Obtain compilation log from the program. 
-
-    size_t logSize; 
-
-    NVRTC_SAFE_CALL(nvrtcGetProgramLogSize(prog, &logSize)); 
-    char *log = new char[logSize]; 
-    NVRTC_SAFE_CALL(nvrtcGetProgramLog(prog, log)); 
-    std::cout << log << '\n'; 
-    delete[] log; 
-
-    if (compileResult != NVRTC_SUCCESS) {
-        exit(1);
-    }
-
-    size_t ptxSize; 
-    NVRTC_SAFE_CALL(nvrtcGetPTXSize(prog, &ptxSize)); 
-    char *ptx = new char[ptxSize]; 
-    NVRTC_SAFE_CALL(nvrtcGetPTX(prog, ptx)); // Destroy the program. 
-    NVRTC_SAFE_CALL(nvrtcDestroyProgram(&prog)); // Load the generated PTX and get a handle to the parent kernel. 
-
-    CUdevice cuDevice; 
-    CUcontext context; 
-    CUmodule module; 
-    CUfunction cuda_op_abs_kernel; 
-    CUDA_SAFE_CALL( cuInit(0)); 
-    CUDA_SAFE_CALL(cuDeviceGet(&cuDevice, 0)); 
-    CUDA_SAFE_CALL(cuCtxCreate(&context, 0, cuDevice)); 
-    // CUDA_SAFE_CALL(cuLinkCreate(0, 0 , 0, &linkState)); 
-    //CUDA_SAFE_CALL(cuLinkeAddFile(linkState, CU_JIT_INPUT_LIBRARY, ' ', 0, 0, 0));
-    //CUDA_SAFE_CALL(cuLinkAddData(linkState, CU_JIT_INPUT_PTX, (void *)ptx, ptxSize, "dynamic_parallelism.ptx", 0, 0, 0));
-
-    //size_t cubinSize;
-    //void *cubin;
-    //CUDA_SAFE_CALL(cuLinkComplete(linkState, &cubin, &cubinSize));
-    CUDA_SAFE_CALL(cuModuleLoadDataEx(&module, ptx, 0, 0, 0));
-    CUDA_SAFE_CALL(cuModuleGetFunction(&cuda_op_abs_kernel, module, "cuda_op_abs"));
-
-
     writer << "{  // " << n->get_name() << "\n";
     writer.indent++;
     writer << "int count = " << out[0].get_size() << ";\n";
     writer << "if(count == 0) return;\n";
-    writer << "void *argsList[] = {(void *)" << args[0].get_name() << ", (void *)" << out[0].get_name() << ", &count};\n";
-    writer << "//cuLaunchKernel(cuda_op_abs_kernel, count, 1, 1, 1, 1, 1, 0, NULL, argsList, 0);\n";
+    writer << "ngraph::runtime::gpu::cuda::kernel::emit_abs(" << args[0].get_name() <<  ", " << out[0].get_name() << ", count);\n";
     writer.indent--;
     writer << "}\n";