diff --git a/projects/hip/tests/src/runtimeApi/module/hipModule.cpp b/projects/hip/tests/src/runtimeApi/module/hipModule.cpp index 22b4e7ff49..019645ea32 100644 --- a/projects/hip/tests/src/runtimeApi/module/hipModule.cpp +++ b/projects/hip/tests/src/runtimeApi/module/hipModule.cpp @@ -17,6 +17,13 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ +/* HIT_START + * BUILD_CMD: vcpy_kernel.code %hc --genco %S/vcpy_kernel.cpp -o vcpy_kernel.code + * BUILD: %t %s ../../test_common.cpp + * TEST: %t + * HIT_END + */ + #include "hip/hip_runtime.h" #include "hip/hip_runtime_api.h" #include @@ -33,13 +40,15 @@ THE SOFTWARE. #define fileName "vcpy_kernel.code" #define kernel_name "hello_world" -__global__ void Cpy(hipLaunchParm lp, float* Ad, float* Bd) { - int tx = threadIdx.x; - Bd[tx] = Ad[tx]; -} +#define HIP_CHECK(status) \ + if (status != hipSuccess) { \ + std::cout << "Got Status: " << status << " at Line: " << __LINE__ << std::endl; \ + exit(0); \ + } int main() { - float *A, *B, *Ad, *Bd; + float *A, *B; + hipDeviceptr_t Ad, Bd; A = new float[LEN]; B = new float[LEN]; @@ -48,32 +57,38 @@ int main() { B[i] = 0.0f; } + HIPCHECK(hipInit(0)); + + hipDevice_t device; + hipCtx_t context; + HIPCHECK(hipDeviceGet(&device, 0)); + HIPCHECK(hipCtxCreate(&context, 0, device)); + HIPCHECK(hipMalloc((void**)&Ad, SIZE)); HIPCHECK(hipMalloc((void**)&Bd, SIZE)); - HIPCHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice)); HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice)); + hipModule_t Module; hipFunction_t Function; + hipFunction_t f; HIPCHECK(hipModuleLoad(&Module, fileName)); HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name)); - hipFunction_t f; - HIPCHECK(hipModuleGetFunction(&f, Module, kernel_name)); - assert(f == Function); + hipStream_t stream; HIPCHECK(hipStreamCreate(&stream)); - void* args[2] = {&Ad, &Bd}; - std::vector argBuffer(5); - memcpy(&argBuffer[3], &Ad, sizeof(void*)); - memcpy(&argBuffer[4], &Bd, sizeof(void*)); + struct { + void* _Ad; + void* _Bd; + } args; + args._Ad = (void*) Ad; + args._Bd = (void*) Bd; + size_t size = sizeof(args); - 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, stream, NULL, (void**)&config); + void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END}; + HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, stream, NULL, (void**)&config)); HIPCHECK(hipStreamDestroy(stream)); @@ -83,12 +98,6 @@ int main() { assert(A[i] == B[i]); } - std::vector vec(1024 * 1024 * 64); - for (unsigned i = 0; i < 1024 * 1024 * 64; i++) { - hipFunction_t func; - hipModuleGetFunction(&func, Module, kernel_name); - vec[i] = func; - } + HIPCHECK(hipCtxDestroy(context)); passed(); - return 0; } diff --git a/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.code b/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.code deleted file mode 100755 index 4fd6472496..0000000000 Binary files a/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.code and /dev/null differ diff --git a/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.cpp b/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.cpp index c99e3df985..d7c23a8810 100644 --- a/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.cpp +++ b/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -22,7 +22,7 @@ THE SOFTWARE. #include "hip/hip_runtime.h" -extern "C" __global__ void hello_world(hipLaunchParm lp, float* a, float* b) { - int tx = threadIdx.x; +extern "C" __global__ void hello_world(float* a, float* b) { + int tx = hipThreadIdx_x; b[tx] = a[tx]; }