From 85914691a6fdb5fca9e453bd1eccca9223571ae5 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 2 Aug 2018 15:16:36 +0530 Subject: [PATCH 1/2] Remove adipose extension from genco output [ROCm/hip commit: 07f8f09aff211b2e67422ef70043813895359e1d] --- projects/hip/lpl_ca/lpl.hpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/hip/lpl_ca/lpl.hpp b/projects/hip/lpl_ca/lpl.hpp index cbd7fe8386..c3992e43c0 100644 --- a/projects/hip/lpl_ca/lpl.hpp +++ b/projects/hip/lpl_ca/lpl.hpp @@ -80,7 +80,7 @@ inline void copy_kernel_section_to_fat_binary(const std::string& tmp, const std: std::find_if(reader.sections.begin(), reader.sections.end(), [](const ELFIO::section* x) { return x->get_name() == kernel_section(); }); - std::ofstream out{output + fat_binary_extension()}; + std::ofstream out{output}; if (it == reader.sections.end()) { std::cerr << "Warning: no kernels were generated; fat binary shall " @@ -95,8 +95,8 @@ inline void generate_fat_binary(const std::vector& sources, const std::vector& targets, const std::string& flags, const std::string& output) { static const auto d = [](const std::string* f) { remove(f->c_str()); }; - - std::unique_ptr tmp{&output, d}; + std::string temp_str = output + ".tmp"; + std::unique_ptr tmp{&temp_str, d}; redi::ipstream hipcc{make_hipcc_call(sources, targets, flags, *tmp), redi::pstream::pstderr}; From 3bae902cae055e34f4416d303ab07f16b60fbee1 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 8 Aug 2018 22:28:13 +0530 Subject: [PATCH 2/2] Clean up module api samples [ROCm/hip commit: 8f0d6849884411f53581e7082257331f42426433] --- .../samples/0_Intro/module_api/runKernel.cpp | 25 +++-------------- .../0_Intro/module_api_global/runKernel.cpp | 25 +++-------------- .../11_texture_driver/texture2dDrv.cpp | 27 +++---------------- 3 files changed, 10 insertions(+), 67 deletions(-) diff --git a/projects/hip/samples/0_Intro/module_api/runKernel.cpp b/projects/hip/samples/0_Intro/module_api/runKernel.cpp index 129c458dd0..ed74a7e4eb 100644 --- a/projects/hip/samples/0_Intro/module_api/runKernel.cpp +++ b/projects/hip/samples/0_Intro/module_api/runKernel.cpp @@ -30,7 +30,7 @@ THE SOFTWARE. #define LEN 64 #define SIZE LEN << 2 -#define fileName "vcpy_kernel.code.adipose" +#define fileName "vcpy_kernel.code" #define kernel_name "hello_world" #define HIP_CHECK(status) \ @@ -66,32 +66,13 @@ int main() { HIP_CHECK(hipModuleLoad(&Module, fileName)); HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name)); -#ifdef __HIP_PLATFORM_HCC__ - uint32_t len = LEN; - uint32_t one = 1; - struct { void* _Ad; void* _Bd; } args; - args._Ad = Ad; - args._Bd = Bd; - -#endif - -#ifdef __HIP_PLATFORM_NVCC__ - struct { - uint32_t _hidden[1]; - void* _Ad; - void* _Bd; - } args; - - args._hidden[0] = 0; - args._Ad = Ad; - args._Bd = Bd; -#endif - + args._Ad = (void*) Ad; + args._Bd = (void*) Bd; size_t size = sizeof(args); diff --git a/projects/hip/samples/0_Intro/module_api_global/runKernel.cpp b/projects/hip/samples/0_Intro/module_api_global/runKernel.cpp index ebae029a16..0047f79b7f 100644 --- a/projects/hip/samples/0_Intro/module_api_global/runKernel.cpp +++ b/projects/hip/samples/0_Intro/module_api_global/runKernel.cpp @@ -30,7 +30,7 @@ THE SOFTWARE. #define LEN 64 #define SIZE LEN * sizeof(float) -#define fileName "vcpy_kernel.code.adipose" +#define fileName "vcpy_kernel.code" float myDeviceGlobal; float myDeviceGlobalArray[16]; #define HIP_CHECK(cmd) \ @@ -79,32 +79,13 @@ int main() { myDeviceGlobalArray[i] = i * 1000.0f; } -#ifdef __HIP_PLATFORM_HCC__ - uint32_t len = LEN; - uint32_t one = 1; - struct { void* _Ad; void* _Bd; } args; - args._Ad = Ad; - args._Bd = Bd; - -#endif - -#ifdef __HIP_PLATFORM_NVCC__ - struct { - uint32_t _hidden[1]; - void* _Ad; - void* _Bd; - } args; - - args._hidden[0] = 0; - args._Ad = Ad; - args._Bd = Bd; -#endif - + args._Ad = (void*) Ad; + args._Bd = (void*) Bd; size_t size = sizeof(args); diff --git a/projects/hip/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp b/projects/hip/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp index 2485d455ef..385ed4775a 100644 --- a/projects/hip/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp +++ b/projects/hip/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp @@ -21,13 +21,13 @@ THE SOFTWARE. */ #include "hip/hip_runtime.h" -#include "hip/hip_runtime_api.h" +//#include "hip/hip_runtime_api.h" #include #include #include -#include +//#include -#define fileName "tex2dKernel.code.adipose" +#define fileName "tex2dKernel.code" texture tex; bool testResult = false; @@ -87,34 +87,15 @@ bool runTest(int argc, char** argv) { float* dData = NULL; hipMalloc((void**)&dData, size); -#ifdef __HIP_PLATFORM_HCC__ - struct { void* _Ad; unsigned int _Bd; unsigned int _Cd; } args; - args._Ad = dData; + args._Ad = (void*) dData; args._Bd = width; args._Cd = height; -#endif - -#ifdef __HIP_PLATFORM_NVCC__ - struct { - uint32_t _hidden[1]; - void* _Ad; - unsigned int _Bd; - unsigned int _Cd; - } args; - - args._hidden[0] = 0; - args._Ad = dData; - args._Bd = width; - args._Cd = height; -#endif - - size_t sizeTemp = sizeof(args); void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE,