GPU: Fix HIP memcpyToSymbol, which requires additional HIP_SYMBOL macro

davidrohr · ktf · commit 67cc2811cf24 · 2020-02-21T14:04:22.000+01:00
diff --git a/GPU/GPUTracking/Base/hip/GPUReconstructionHIP.hip.cxx b/GPU/GPUTracking/Base/hip/GPUReconstructionHIP.hip.cxx
@@ -412,9 +412,9 @@ size_t GPUReconstructionHIPBackend::WriteToConstantMemory(size_t offset, const v
 {
 #ifndef GPUCA_HIP_NO_CONSTANT_MEMORY
   if (stream == -1) {
-    GPUFailedMsg(hipMemcpyToSymbol(gGPUConstantMemBuffer, src, size, offset, hipMemcpyHostToDevice));
+    GPUFailedMsg(hipMemcpyToSymbol(HIP_SYMBOL(gGPUConstantMemBuffer), src, size, offset, hipMemcpyHostToDevice));
   } else {
-    GPUFailedMsg(hipMemcpyToSymbolAsync(gGPUConstantMemBuffer, src, size, offset, hipMemcpyHostToDevice, mInternals->HIPStreams[stream]));
+    GPUFailedMsg(hipMemcpyToSymbolAsync(HIP_SYMBOL(gGPUConstantMemBuffer), src, size, offset, hipMemcpyHostToDevice, mInternals->HIPStreams[stream]));
   }
   if (ev && stream != -1) {
     GPUFailedMsg(hipEventRecord(*(hipEvent_t*)ev, mInternals->HIPStreams[stream]));

Original file line number	Diff line number	Diff line change
`@@ -412,9 +412,9 @@ size_t GPUReconstructionHIPBackend::WriteToConstantMemory(size_t offset, const v`
`412`	`412`	`{`
`413`	`413`	`#ifndef GPUCA_HIP_NO_CONSTANT_MEMORY`
`414`	`414`	`if (stream == -1) {`
`415`		`- GPUFailedMsg(hipMemcpyToSymbol(gGPUConstantMemBuffer, src, size, offset, hipMemcpyHostToDevice));`
	`415`	`+ GPUFailedMsg(hipMemcpyToSymbol(HIP_SYMBOL(gGPUConstantMemBuffer), src, size, offset, hipMemcpyHostToDevice));`
`416`	`416`	`} else {`
`417`		`- GPUFailedMsg(hipMemcpyToSymbolAsync(gGPUConstantMemBuffer, src, size, offset, hipMemcpyHostToDevice, mInternals->HIPStreams[stream]));`
	`417`	`+ GPUFailedMsg(hipMemcpyToSymbolAsync(HIP_SYMBOL(gGPUConstantMemBuffer), src, size, offset, hipMemcpyHostToDevice, mInternals->HIPStreams[stream]));`
`418`	`418`	`}`
`419`	`419`	`if (ev && stream != -1) {`
`420`	`420`	`GPUFailedMsg(hipEventRecord((hipEvent_t)ev, mInternals->HIPStreams[stream]));`