From e9532108bb1674e7262616f8e3401e499c1f89d1 Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Tue, 18 Dec 2018 15:48:21 -0500 Subject: [PATCH 1/2] Add default arguments for hipConfigureCall [ROCm/hip commit: ec03a8b3520235b085bfc915f1047846ad6a7f91] --- .../include/hip/hcc_detail/hip_runtime_api.h | 2 +- .../hip/tests/src/kernel/hipGridLaunch.cpp | 31 +++++++++++++++++++ 2 files changed, 32 insertions(+), 1 deletion(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index 97d6320f7f..d4b8c1faf4 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -2573,7 +2573,7 @@ hipError_t hipIpcCloseMemHandle(void* devPtr); * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue * */ -hipError_t hipConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem, hipStream_t stream); +hipError_t hipConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem = 0, hipStream_t stream = 0); /** diff --git a/projects/hip/tests/src/kernel/hipGridLaunch.cpp b/projects/hip/tests/src/kernel/hipGridLaunch.cpp index 8505a0758f..4f7883c656 100644 --- a/projects/hip/tests/src/kernel/hipGridLaunch.cpp +++ b/projects/hip/tests/src/kernel/hipGridLaunch.cpp @@ -74,12 +74,43 @@ int test_gl2(size_t N) { return 0; } +#if __HIP__ +int test_triple_chevron(size_t N) { + size_t Nbytes = N * sizeof(int); + + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N); + + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + + // Full vadd in one large chunk, to get things started: + HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + + vectorADD2<<>>(A_d, B_d, C_d, N); + + HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + + HIPCHECK(hipDeviceSynchronize()); + + HipTest::checkVectorADD(A_h, B_h, C_h, N); + + return 0; +} +#endif int main(int argc, char* argv[]) { HipTest::parseStandardArguments(argc, argv, true); test_gl2(N); +#if __HIP__ + test_triple_chevron(N); +#endif passed(); } From a2ab2b0370af9bac5381271a9440a30a8a2afcf9 Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Mon, 31 Dec 2018 13:07:19 -0500 Subject: [PATCH 2/2] Use __dparm for default parameter [ROCm/hip commit: 548f4dd4e51250a579d1b5b57154ea2368be1fca] --- projects/hip/include/hip/hcc_detail/hip_runtime_api.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index d4b8c1faf4..0d526ca25d 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -2573,7 +2573,7 @@ hipError_t hipIpcCloseMemHandle(void* devPtr); * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue * */ -hipError_t hipConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem = 0, hipStream_t stream = 0); +hipError_t hipConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem __dparm(0), hipStream_t stream __dparm(0)); /**