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