Use HIP atomicAdd as a workaround if compiler is hip-clang
@@ -458,6 +458,34 @@ class KernelWriterSource(KernelWriter):
kStr += " } while (assumed != old);%s" % (self.endLine)
kStr += "}%s" % (self.endLine)
"""
+ if globalParameters["CxxCompiler"] == "hipcc":
+ kStr += self.endLine
+ kStr += "__device__ inline int atomicAddType(int *fPtr, int operand)%s" % (self.endLine)
+ kStr += "{%s" % (self.endLine)
+ kStr += " return atomicAdd(fPtr,operand);%s" % (self.endLine)
+ kStr += "}%s" % (self.endLine)
+ kStr += self.endLine
+ kStr += "__device__ inline unsigned int atomicAddType(unsigned int *fPtr, unsigned int operand)%s" % (self.endLine)
+ kStr += "{%s" % (self.endLine)
+ kStr += " return atomicAdd(fPtr,operand);%s" % (self.endLine)
+ kStr += "}%s" % (self.endLine)
+ kStr += self.endLine
+ kStr += "__device__ inline unsigned long long int atomicAddType(unsigned long long int *fPtr, unsigned long long int operand)%s" % (self.endLine)
+ kStr += "{%s" % (self.endLine)
+ kStr += " return atomicAdd(fPtr,operand);%s" % (self.endLine)
+ kStr += "}%s" % (self.endLine)
+ kStr += self.endLine
+ kStr += "__device__ inline float atomicAddType(float *fPtr, float operand)%s" % (self.endLine)
+ kStr += "{%s" % (self.endLine)
+ kStr += " return atomicAdd(fPtr,operand);%s" % (self.endLine)
+ kStr += "}%s" % (self.endLine)
+ kStr += self.endLine
+ kStr += "__device__ inline double atomicAddType(double *fPtr, double operand)%s" % (self.endLine)
+ kStr += "{%s" % (self.endLine)
+ kStr += " return atomicAdd(fPtr,operand);%s" % (self.endLine)
+ kStr += "}%s" % (self.endLine)
+ kStr += self.endLine
+ else:
kStr += self.endLine
kStr += "template <typename T>%s" % (self.endLine)
kStr += "__device__ inline void atomicAddType(%s%sT *fPtr, T operand) {%s" \
