一个线程安全的问题

[zjhellofss]

src/ppl/cv/cuda/warp.hpp 第275行
线程调用transform.calculateCoordinates会改变一个transform实例内的x,y值。如果多个线程同时调用，他们操作的transoform实例属于同一个，我想这里是否需要考虑线程安全的问题？

template <typename Transform>
__global__
void warpLinearKernel(const uchar* src, int src_rows, int src_cols,
                      int channels, int src_stride, Transform transform,
                      uchar* dst, int dst_rows, int dst_cols, int dst_stride,
                      BorderType border_type, uchar border_value) {
  int element_x = (blockIdx.x << kBlockShiftX0) + threadIdx.x;
  int element_y = (blockIdx.y << kBlockShiftY0) + threadIdx.y;
  if (element_x >= dst_cols || element_y >= dst_rows) {
    return;
  }

  transform.calculateCoordinates(element_x, element_y);
  float src_x = transform.getX();
  float src_y = transform.getY();

[jimurk]

你是在多线程里调用warpaffine/warpperspective吗？

[zjhellofss]

我指的是一个block内多个threads的并发访问。

[zjhellofss]

block1 thread1: 使用 transform.calculateCoordinates(element_x, element_y)， 修改了transform中的x,y
block1 thread2: 同时也会访问calculateCoordinates(element_x, element_y)，来修改transform中的x,y
然后如果又切回到block1 thread1 ,它再调用 float src_x = transform.getX(); 就会得到被thread2修改过的x和y值

[jimurk]

block1 thread1: 使用 transform.calculateCoordinates(element_x, element_y)， 修改了transform中的x,y block1 thread2: 同时也会访问calculateCoordinates(element_x, element_y)，来修改transform中的x,y 然后如果又切回到block1 thread1 ,它再调用 float src_x = transform.getX(); 就会得到被thread2修改过的x和y值

嗯，是会有这个问题，我测和调一下。谢谢了。

[zjhellofss]

block1 thread1: 使用 transform.calculateCoordinates(element_x, element_y)， 修改了transform中的x,y block1 thread2: 同时也会访问calculateCoordinates(element_x, element_y)，来修改transform中的x,y 然后如果又切回到block1 thread1 ,它再调用 float src_x = transform.getX(); 就会得到被thread2修改过的x和y值

嗯，是会有这个问题，我测和调一下。谢谢了。

我觉得可以改成这样

  __device__ void calculateCoordinates(int element_x, int element_y,float&x ,float& y) {
    x = values[0] * element_x + values[1] * element_y + values[2];
    y = values[3] * element_x + values[4] * element_y + values[5];
  }

[jimurk]

block1 thread1: 使用 transform.calculateCoordinates(element_x, element_y)， 修改了transform中的x,y block1 thread2: 同时也会访问calculateCoordinates(element_x, element_y)，来修改transform中的x,y 然后如果又切回到block1 thread1 ,它再调用 float src_x = transform.getX(); 就会得到被thread2修改过的x和y值

嗯，是会有这个问题，我测和调一下。谢谢了。

我觉得可以改成这样

  __device__ void calculateCoordinates(int element_x, int element_y,float&x ,float& y) {
    x = values[0] * element_x + values[1] * element_y + values[2];
    y = values[3] * element_x + values[4] * element_y + values[5];
  }

可以这样改的，我再测一下，考虑一下。

[jimurk]

这个地方我debug了一下kernel, 无论x/y值，还是最终结果正确性没有问题。理论上传给kernel的参数应该放在constant memory，需要修改数值的变量需要放在global memory然后传指针；因为需要修改x/y数值，编译器应该把这个自定义struct放到了local memory或寄存器，每个线程保存一份，所以计算没有出错。但这么写不太规范，晚点我改一下，把x/y从struct剔除。

[zjhellofss]

这个地方我debug了一下kernel, 无论x/y值，还是最终结果正确性没有问题。理论上传给kernel的参数应该放在constant memory，需要修改数值的变量需要放在global memory然后传指针；因为需要修改x/y数值，编译器应该把这个自定义struct放到了local memory或寄存器，每个线程保存一份，所以计算没有出错。但这么写不太规范，晚点我改一下，把x/y从struct剔除。

明白了。我之前也实验过，确实不会影响到最后的结果。您耐心解释了那么多，我也比较有收获，非常感谢！

[jimurk]

这个地方我debug了一下kernel, 无论x/y值，还是最终结果正确性没有问题。理论上传给kernel的参数应该放在constant memory，需要修改数值的变量需要放在global memory然后传指针；因为需要修改x/y数值，编译器应该把这个自定义struct放到了local memory或寄存器，每个线程保存一份，所以计算没有出错。但这么写不太规范，晚点我改一下，把x/y从struct剔除。

明白了。我之前也实验过，确实不会影响到最后的结果。您耐心解释了那么多，我也比较有收获，非常感谢！

客气了，互相学习，多多交流，共同进步。